Add __launch_bounds__ (CUDA/HIP) and reqd_work_group_size (SYCL) to kernels with fixed block sizes

README.md

@@ -1799,7 +1799,7 @@ Early results are shown [here](results/README.md)
 Authored and maintained by Zheming Jin (https://github.com/zjin-lcf) 
 
 ## Acknowledgement
-Abhishek Bagusetty, Andrew Barker, Andrey Alekseenko, Anton Gorshkov, Beau Johnston, Bernhard Esslinger, Bert de Jong, Chengjian Liu, Chris Knight, Daine McNiven, David Oro, David Sanchez, Douglas Franz, Edson Borin, Gabriell Araujo, Georgi Mirazchiyski, Henry Linjamäki, Henry Gabb, Hugh Delaney, Ian Karlin, Istvan Reguly, Jack Kirk, Jason Lau, Jeff Hammond, Jenny Chen, Jeongnim Kim, Jianxin Qiu, Jakub Chlanda, Jiya Su, John Tramm, Ju Zheng, Junchao Zhang, Kali Uday Balleda, Kinman Lei, Luke Drummond, Manish Motwani, Mark Stephenson, Martin Burtscher, Matthias Noack, Matthew Davis, Michael Kruse, Michel Migdal, Mike Franusich, Mike Giles, Mikhail Dvorskiy, Mohammed Alser, Mohammed ,Tarek Ibn Ziad, Muhammad Haseeb, Muaaz Awan, Nevin Liber, Nicholas Miller, Oscar Ludwig, Pavel Samolysov, Pedro Valero Lara, Piotr Różański, Rahulkumar Gayatri, Samyak Gangwal, Shaoyi Peng, Steffen Larsen, Rafal Bielski, Robert Harrison, Robin Kobus, Rod Burns, Rodrigo Vimieiro, Romanov Vlad, Tadej Ciglarič, Thomas Applencourt, Tiago Carneiro, Tiago Cogumbreiro, Timmie Smith, Tobias Baumann, Usman Roshan, Wayne Joubert, Ye Luo, Yongbin Gu, Zhe Chen
+Abhishek Bagusetty, Andrew Barker, Andrey Alekseenko, Anton Gorshkov, Beau Johnston, Bernhard Esslinger, Bert de Jong, Chengjian Liu, Chris Knight, Daine McNiven, David Oro, David Sanchez, Douglas Franz, Edson Borin, Gabriell Araujo, Georgi Mirazchiyski, Henry Linjamäki, Henry Gabb, Hugh Delaney, Ian Karlin, Istvan Reguly, Jack Kirk, Jason Lau, Jeff Hammond, Jenny Chen, Jeongnim Kim, Jianxin Qiu, Jakub Chlanda, Jiya Su, John Tramm, Ju Zheng, Junchao Zhang, Kali Uday Balleda, Kinman Lei, Luke Drummond, Manish Motwani, Mark Stephenson, Martin Burtscher, Matthias Noack, Matthew Davis, Michael Kruse, Michel Migdal, Mike Franusich, Mike Giles, Mikhail Dvorskiy, Mohammed Alser, Mohammed ,Tarek Ibn Ziad, Muhammad Haseeb, Muaaz Awan, Nevin Liber, Nicholas Miller, Oscar Ludwig, Pavel Samolysov, Pedro Valero Lara, Piotr Różański, Rahulkumar Gayatri, Samyak Gangwal, Shaoyi Peng, Steffen Larsen, Rafal Bielski, Robert Harrison, Robin Kobus, Rod Burns, Rodrigo Vimieiro, Romanov Vlad, Tadej Ciglarič, Thomas Applencourt, Tiago Carneiro, Tiago Cogumbreiro, Timmie Smith, Tobias Baumann, Usman Roshan, Wayne Joubert, Ye Luo, Yongbin Gu, Zhe Chen, Zhihui Du
 
 Codeplay<sup>®</sup> and Intel<sup>®</sup> for their contributions to the oneAPI ecosystem   
 

benchmarks.yaml

@@ -63,7 +63,7 @@ adv:
   models: [cuda, hip, omp, sycl]
   test:
     regex: '(?:elapsed time=)([0-9.+-e]+)(?: )'
-    args: ["16", "16", "16"]
+    args: ["7", "15", "16"]
     timeout: 300
 
 aes:

src/background-subtract-cuda/main.cu

@@ -8,6 +8,7 @@
 
 #define BLOCK_SIZE 256
 
+__launch_bounds__(BLOCK_SIZE)
 __global__ void findMovingPixels(
   const size_t imgSize,
   const unsigned char *__restrict__ Img,
@@ -26,6 +27,7 @@ __global__ void findMovingPixels(
 }
 
 // alpha = 0.92
+__launch_bounds__(BLOCK_SIZE)
 __global__ void updateBackground(
   const size_t imgSize,
   const unsigned char *__restrict__ Img,
@@ -38,6 +40,7 @@ __global__ void updateBackground(
 }
 
 // alpha = 0.92, c = 3
+__launch_bounds__(BLOCK_SIZE)
 __global__ void updateThreshold(
   const size_t imgSize,
   const unsigned char *__restrict__ Img,
@@ -56,6 +59,7 @@ __global__ void updateThreshold(
 //
 // merge three kernels into a single kernel
 //
+__launch_bounds__(BLOCK_SIZE)
 __global__ void merge(
   const size_t imgSize,
   const unsigned char *__restrict__ Img,

src/background-subtract-hip/main.cu

@@ -8,6 +8,7 @@
 
 #define BLOCK_SIZE 256
 
+__launch_bounds__(BLOCK_SIZE)
 __global__ void findMovingPixels(
   const size_t imgSize,
   const unsigned char *__restrict__ Img,
@@ -26,6 +27,7 @@ __global__ void findMovingPixels(
 }
 
 // alpha = 0.92
+__launch_bounds__(BLOCK_SIZE)
 __global__ void updateBackground(
   const size_t imgSize,
   const unsigned char *__restrict__ Img,
@@ -38,6 +40,7 @@ __global__ void updateBackground(
 }
 
 // alpha = 0.92, c = 3
+__launch_bounds__(BLOCK_SIZE)
 __global__ void updateThreshold(
   const size_t imgSize,
   const unsigned char *__restrict__ Img,
@@ -56,6 +59,7 @@ __global__ void updateThreshold(
 //
 // merge three kernels into a single kernel
 //
+__launch_bounds__(BLOCK_SIZE)
 __global__ void merge(
   const size_t imgSize,
   const unsigned char *__restrict__ Img,

src/background-subtract-sycl/main.cpp

@@ -158,7 +158,7 @@ int main(int argc, char* argv[]) {
         auto start = std::chrono::steady_clock::now();
         q.submit([&] (sycl::handler &cgh) {
           cgh.parallel_for<class merged_kernel>(
-            sycl::nd_range<1>(gws, lws), [=] (sycl::nd_item<1> item) {
+            sycl::nd_range<1>(gws, lws), [[sycl::reqd_work_group_size(BLOCK_SIZE)]] [=] (sycl::nd_item<1> item) {
             merge ( item, imgSize, d_Img, d_Img1, d_Img2, d_Tn, d_Bn );
           });
         }).wait();
@@ -169,19 +169,19 @@ int main(int argc, char* argv[]) {
         auto start = std::chrono::steady_clock::now();
         q.submit([&] (sycl::handler &cgh) {
           cgh.parallel_for<class k1>(
-            sycl::nd_range<1>(gws, lws), [=] (sycl::nd_item<1> item) {
+            sycl::nd_range<1>(gws, lws), [[sycl::reqd_work_group_size(BLOCK_SIZE)]] [=] (sycl::nd_item<1> item) {
             findMovingPixels ( item, imgSize, d_Img, d_Img1, d_Img2, d_Tn, d_Mp );
           });
         });
         q.submit([&] (sycl::handler &cgh) {
           cgh.parallel_for<class k2>(
-            sycl::nd_range<1>(gws, lws), [=] (sycl::nd_item<1> item) {
+            sycl::nd_range<1>(gws, lws), [[sycl::reqd_work_group_size(BLOCK_SIZE)]] [=] (sycl::nd_item<1> item) {
             updateBackground  ( item, imgSize, d_Img, d_Mp, d_Bn );
           });
         });
         q.submit([&] (sycl::handler &cgh) {
           cgh.parallel_for<class k3>(
-            sycl::nd_range<1>(gws, lws), [=] (sycl::nd_item<1> item) {
+            sycl::nd_range<1>(gws, lws), [[sycl::reqd_work_group_size(BLOCK_SIZE)]] [=] (sycl::nd_item<1> item) {
             updateThreshold  ( item, imgSize, d_Img, d_Mp, d_Bn, d_Tn );
           });
         });

src/bitonic-sort-cuda/main.cu

@@ -43,6 +43,7 @@
 
 #define BLOCK_SIZE 256
 
+__launch_bounds__(BLOCK_SIZE)
 __global__
 void bitonic_sort (const int seq_len, const int two_power, int *a)
 {

src/bitonic-sort-hip/main.cu

@@ -43,6 +43,7 @@
 
 #define BLOCK_SIZE 256
 
+__launch_bounds__(BLOCK_SIZE)
 __global__
 void bitonic_sort (const int seq_len, const int two_power, int *a)
 {

src/bitonic-sort-sycl/main.cpp

@@ -76,7 +76,7 @@ void ParallelBitonicSort(int input[], int n) {
       // Offload the work to kernel.
       q.submit([&](sycl::handler &h) {
         h.parallel_for(sycl::nd_range<1>(
-          sycl::range<1>(size), sycl::range<1>(BLOCK_SIZE)), [=](sycl::nd_item<1> item) {
+          sycl::range<1>(size), sycl::range<1>(BLOCK_SIZE)), [[sycl::reqd_work_group_size(BLOCK_SIZE)]] [=](sycl::nd_item<1> item) {
           int i = item.get_global_id(0);
 
           // Assign the bitonic sequence number.

src/channelSum-cuda/main.cu

@@ -49,6 +49,7 @@ __device__ void BlockReduce(T &input1, T &input2) {
 
 template <typename T, int kBlockDimX, int kBlockDimY>
 __global__
+__launch_bounds__(kBlockDimX * kBlockDimY)
 void ChannelSumNCHW(
     const int N,
     const int C,
@@ -79,6 +80,7 @@ void ChannelSumNCHW(
 
 template <typename T>
 __global__
+__launch_bounds__(NUM_THREADS)
 void ChannelSumNHWC(
     const int N,
     const int C,

src/channelSum-hip/main.cu

@@ -49,6 +49,7 @@ __device__ void BlockReduce(T &input1, T &input2) {
 
 template <typename T, int kBlockDimX, int kBlockDimY>
 __global__
+__launch_bounds__(kBlockDimX * kBlockDimY)
 void ChannelSumNCHW(
     const int N,
     const int C,
@@ -78,6 +79,7 @@ void ChannelSumNCHW(
 
 template <typename T>
 __global__
+__launch_bounds__(NUM_THREADS)
 void ChannelSumNHWC(
     const int N,
     const int C,

src/channelSum-sycl/main.cpp

@@ -19,7 +19,7 @@ typedef int scalar_t;
      sycl::range<2> lws (128, 1);                                      \
      q.submit([&](sycl::handler &cgh) {                                \
        cgh.parallel_for(                                               \
-         sycl::nd_range<2>(gws, lws), [=] (sycl::nd_item<2> item) {    \
+         sycl::nd_range<2>(gws, lws), [=] [[sycl::reqd_work_group_size(128, 1)]] (sycl::nd_item<2> item) {    \
          Func<T>(item, __VA_ARGS__);                                   \
        });                                                             \
      });                                                               \
@@ -28,7 +28,7 @@ typedef int scalar_t;
      sycl::range<2> lws (64, 2);                                       \
      q.submit([&](sycl::handler &cgh) {                                \
        cgh.parallel_for(                                               \
-         sycl::nd_range<2>(gws, lws), [=] (sycl::nd_item<2> item) {    \
+         sycl::nd_range<2>(gws, lws), [=] [[sycl::reqd_work_group_size(64, 2)]] (sycl::nd_item<2> item) {    \
          Func<T>(item, __VA_ARGS__);                                   \
        });                                                             \
      });                                                               \
@@ -37,7 +37,7 @@ typedef int scalar_t;
      sycl::range<2> lws (32, 4);                                       \
      q.submit([&](sycl::handler &cgh) {                                \
        cgh.parallel_for(                                               \
-         sycl::nd_range<2>(gws, lws), [=] (sycl::nd_item<2> item) {    \
+         sycl::nd_range<2>(gws, lws), [=] [[sycl::reqd_work_group_size(32, 4)]] (sycl::nd_item<2> item) {    \
          Func<T>(item, __VA_ARGS__);                                   \
        });                                                             \
      });                                                               \
@@ -46,7 +46,7 @@ typedef int scalar_t;
      sycl::range<2> lws (16, 8);                                       \
      q.submit([&](sycl::handler &cgh) {                                \
        cgh.parallel_for(                                               \
-         sycl::nd_range<2>(gws, lws), [=] (sycl::nd_item<2> item) {    \
+         sycl::nd_range<2>(gws, lws), [=] [[sycl::reqd_work_group_size(16, 8)]] (sycl::nd_item<2> item) {    \
          Func<T>(item, __VA_ARGS__);                                   \
        });                                                             \
      });                                                               \

src/dpid-cuda/Makefile

@@ -37,7 +37,7 @@ endif
 
 # Optimization Flags
 ifeq ($(OPTIMIZE),yes)
-  CFLAGS += -O3
+  CFLAGS += -O3 -use_fast_math
 endif
 #===============================================================================
 # Targets to Build
@@ -46,7 +46,10 @@ endif
 $(program): $(obj) Makefile
 	$(CC) $(CFLAGS) $(obj) -o $@ $(LDFLAGS)
 
-%.o: %.cu shared.h Makefile
+main.o: main.cu shared.h reference.h Makefile
+	$(CC) $(CFLAGS) -c $< -o $@
+
+kernels.o: kernels.cu shared.h Makefile
 	$(CC) $(CFLAGS) -c $< -o $@
 
 clean:

src/dpid-cuda/main.cu

@@ -8,6 +8,7 @@
 #include <cmath>
 #include <cstdint>
 #include "shared.h"
+#include "reference.h"
 
 double LCG_random_double(uint64_t * seed)
 {
@@ -70,14 +71,35 @@ int main(int argc, char** argv) {
   // run downsampling on a device
   run(p, hInput, hOutput);
 
-  int x = 0, y = 0, z = 0;
+  // per-pixel comparison for validation
+  uchar3 *hRef = (uchar3*) malloc(sizeof(uchar3) * p.oWidth * p.oHeight);
+  reference(p.oWidth, p.oHeight, p.iWidth, p.iHeight,
+            p.pWidth, p.pHeight, p.lambda,
+            hInput, hRef);
+  uint32_t mismatches = 0;
+  int maxDiff = 0;
   for (uint32_t i = 0; i < p.oWidth * p.oHeight; i++) {
-    x += hOutput[i].x;
-    y += hOutput[i].y;
-    z += hOutput[i].z;
+    int dx = std::abs((int)hOutput[i].x - (int)hRef[i].x);
+    int dy = std::abs((int)hOutput[i].y - (int)hRef[i].y);
+    int dz = std::abs((int)hOutput[i].z - (int)hRef[i].z);
+    if (dx > 0 || dy > 0 || dz > 0) mismatches++;
+    if (dx > maxDiff) maxDiff = dx;
+    if (dy > maxDiff) maxDiff = dy;
+    if (dz > maxDiff) maxDiff = dz;
+  }
+  uint32_t total = p.oWidth * p.oHeight;
+  if (mismatches == 0) {
+    printf("Verification PASS: GPU matches CPU reference exactly.\n");
+  } else {
+    printf("Verification: %u / %u pixels differ (%.2f%%), max channel diff = %d\n",
+           mismatches, total, 100.0 * mismatches / total, maxDiff);
+    if (maxDiff <= 1)
+      printf("  -> within rounding tolerance (maxDiff<=1, likely OK)\n");
+    else
+      printf("  -> WARNING: differences exceed rounding tolerance\n");
   }
-  printf("Checksums %d %d %d\n", x, y, z);
 
+  free(hRef);
   free(hInput);
   free(hOutput);
   return 0;