Fix CUDA kernel launch syntax formatting

arrayfire · Jan 5, 2021 · 74b5b4c · 74b5b4c
1 parent 7a69059
commit 74b5b4c
Show file tree

Hide file tree

Showing 8 changed files with 53 additions and 23 deletions.
diff --git a/examples/cuda/bubblechart.cu b/examples/cuda/bubblechart.cu
@@ -40,7 +40,9 @@ __global__ void setupRandomKernel(curandState* states,
 int main(void) {
     FORGE_CUDA_CHECK(
         cudaMalloc((void**)&state, DATA_SIZE * sizeof(curandState_t)));
-    setupRandomKernel<<<divup(DATA_SIZE, 32), 32> > >(state, 314567);
+    // clang-format off
+    setupRandomKernel<<<divup(DATA_SIZE, 32), 32>>>(state, 314567);
+    // clang-format on
 
     float* cos_out;
     float* tan_out;
@@ -173,11 +175,13 @@ void kernel(float* dev_out, int functionCode, float* colors, float* alphas,
     static const dim3 threads(32);
     dim3 blocks(divup(DATA_SIZE, 32));
 
-    mapKernel<<<blocks, threads> > >(dev_out, functionCode, FRANGE_START, DX);
+    // clang-format off
+    mapKernel<<<blocks, threads>>>(dev_out, functionCode, FRANGE_START, DX);
 
-    if (colors) colorsKernel<<<blocks, threads> > >(colors, state);
+    if (colors) colorsKernel<<<blocks, threads>>>(colors, state);
 
-    if (alphas) randKernel<<<blocks, threads> > >(alphas, state, 0, 1);
+    if (alphas) randKernel<<<blocks, threads>>>(alphas, state, 0, 1);
 
-    if (radii) randKernel<<<blocks, threads> > >(radii, state, 20, 60);
+    if (radii) randKernel<<<blocks, threads>>>(radii, state, 20, 60);
+    // clang-format on
 }
diff --git a/examples/cuda/field.cu b/examples/cuda/field.cu
@@ -114,6 +114,8 @@ void generatePoints(float* points, float* dirs) {
     dim3 blocks(divup((int)(NELEMS), threads.x),
                 divup((int)(NELEMS), threads.y));
 
-    pointGenKernel<<<blocks, threads> > >(points, dirs, (int)(NELEMS), MINIMUM,
-                                          STEP);
+    // clang-format off
+    pointGenKernel<<<blocks, threads>>>(points, dirs, (int)(NELEMS), MINIMUM,
+                                        STEP);
+    // clang-format on
 }
diff --git a/examples/cuda/fractal.cu b/examples/cuda/fractal.cu
@@ -99,5 +99,7 @@ void kernel(unsigned char* dev_out) {
     static const dim3 threads(8, 8);
     dim3 blocks(divup(DIMX, threads.x), divup(DIMY, threads.y));
 
-    julia<<<blocks, threads> > >(dev_out);
+    // clang-format off
+    julia<<<blocks, threads>>>(dev_out);
+    // clang-format on
 }
diff --git a/examples/cuda/histogram.cu b/examples/cuda/histogram.cu
@@ -59,7 +59,9 @@ int main(void) {
     Bitmap bmp = createBitmap(IMGW, IMGH);
 
     FORGE_CUDA_CHECK(cudaMalloc((void**)&state, NBINS * sizeof(curandState_t)));
-    setupRandomKernel<<<1, NBINS> > >(state, 314567);
+    // clang-format off
+    setupRandomKernel<<<1, NBINS>>>(state, 314567);
+    // clang-format on
 
     /*
      * First Forge call should be a window creation call
@@ -227,18 +229,24 @@ void PerlinNoise::generateNoise() {
     float amp         = 1.0f;
     float tamp        = 0.0f;
 
-    perlinInitKernel<<<blocks, threads> > >(base, perlin, state);
+    // clang-format off
+    perlinInitKernel<<<blocks, threads>>>(base, perlin, state);
+    // clang-format on
 
     for (int octave = 6; octave >= 0; --octave) {
         int period = 1 << octave;
 
-        perlinComputeKernel<<<blocks, threads> > >(perlin, base, amp, period);
+        // clang-format off
+        perlinComputeKernel<<<blocks, threads>>>(perlin, base, amp, period);
+        // clang-format on
 
         tamp += amp;
         amp *= persistence;
     }
 
-    perlinNormalize<<<blocks, threads> > >(perlin, tamp);
+    // clang-format off
+    perlinNormalize<<<blocks, threads>>>(perlin, tamp);
+    // clang-format on
 }
 
 __global__ void fillImageKernel(unsigned char* ptr, unsigned width,
@@ -268,8 +276,10 @@ void kernel(Bitmap& bmp, PerlinNoise& pn) {
 
     dim3 blocks(divup(bmp.width, threads.x), divup(bmp.height, threads.y));
 
-    fillImageKernel<<<blocks, threads> > >(bmp.ptr, bmp.width, bmp.height,
-                                           pn.perlin);
+    // clang-format off
+    fillImageKernel<<<blocks, threads>>>(bmp.ptr, bmp.width, bmp.height,
+                                         pn.perlin);
+    // clang-format on
 }
 
 __global__ void histogramKernel(const unsigned char* perlinNoise, int* histOut,
@@ -300,7 +310,9 @@ void populateBins(Bitmap& bmp, int* histOut, const unsigned nbins,
 
     cudaMemset(histOut, 0, nbins * sizeof(int));
 
-    histogramKernel<<<blocks, threads> > >(bmp.ptr, histOut, nbins);
+    // clang-format off
+    histogramKernel<<<blocks, threads>>>(bmp.ptr, histOut, nbins);
 
-    histColorsKernel<<<1, nbins> > >(histColors, state);
+    histColorsKernel<<<1, nbins>>>(histColors, state);
+    // clang-format on
 }
diff --git a/examples/cuda/plot3.cu b/examples/cuda/plot3.cu
@@ -94,5 +94,7 @@ void kernel(float t, float dx, float* dev_out) {
     static const dim3 threads(1024);
     dim3 blocks(divup(ZSIZE, 1024));
 
-    generateCurve<<<blocks, threads> > >(t, dx, dev_out, ZMIN, ZSIZE);
+    // clang-format off
+    generateCurve<<<blocks, threads>>>(t, dx, dev_out, ZMIN, ZSIZE);
+    // clang-format on
 }
diff --git a/examples/cuda/plotting.cu b/examples/cuda/plotting.cu
@@ -144,6 +144,8 @@ void kernel(float* dev_out, int functionCode) {
     static const dim3 threads(1024);
     dim3 blocks(divup(DATA_SIZE, 1024));
 
-    simple_sinf<<<blocks, threads> > >(dev_out, DATA_SIZE, functionCode, dx,
-                                       FRANGE_START);
+    // clang-format off
+    simple_sinf<<<blocks, threads>>>(dev_out, DATA_SIZE, functionCode, dx,
+                                     FRANGE_START);
+    // clang-format on
 }
diff --git a/examples/cuda/stream.cu b/examples/cuda/stream.cu
@@ -103,7 +103,9 @@ void generateColors(float* colors) {
     static const dim3 threads(512);
     dim3 blocks(divup(numElems, threads.x));
 
-    genColorsKernel<<<blocks, threads> > >(colors, numElems);
+    // clang-format off
+    genColorsKernel<<<blocks, threads>>>(colors, numElems);
+    // clang-format on
 }
 
 __global__ void pointGenKernel(float* points, float* dirs, int nBBS0,
@@ -137,6 +139,8 @@ void generatePoints(float* points, float* dirs) {
 
     dim3 blocks(blk_x * NELEMS, blk_y);
 
-    pointGenKernel<<<blocks, threads> > >(points, dirs, blk_x, NELEMS, MINIMUM,
-                                          STEP);
+    // clang-format off
+    pointGenKernel<<<blocks, threads>>>(points, dirs, blk_x, NELEMS, MINIMUM,
+                                        STEP);
+    // clang-format on
 }
diff --git a/examples/cuda/surface.cu b/examples/cuda/surface.cu
@@ -88,5 +88,7 @@ void kernel(float dx, float* dev_out) {
     static const dim3 threads(8, 8);
     dim3 blocks(divup(XSIZE, threads.x), divup(YSIZE, threads.y));
 
-    sincos_surf<<<blocks, threads> > >(dx, dev_out, XMIN, YMIN, XSIZE, YSIZE);
+    // clang-format off
+    sincos_surf<<<blocks, threads>>>(dx, dev_out, XMIN, YMIN, XSIZE, YSIZE);
+    // clang-format on
 }