Skip to content

Commit

Permalink
Fix CUDA kernel launch syntax formatting
Browse files Browse the repository at this point in the history
  • Loading branch information
9prady9 committed Jan 5, 2021
1 parent 7a69059 commit 74b5b4c
Show file tree
Hide file tree
Showing 8 changed files with 53 additions and 23 deletions.
14 changes: 9 additions & 5 deletions examples/cuda/bubblechart.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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
}
6 changes: 4 additions & 2 deletions examples/cuda/field.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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
}
4 changes: 3 additions & 1 deletion examples/cuda/fractal.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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
}
28 changes: 20 additions & 8 deletions examples/cuda/histogram.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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,
Expand Down Expand Up @@ -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,
Expand Down Expand Up @@ -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
}
4 changes: 3 additions & 1 deletion examples/cuda/plot3.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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
}
6 changes: 4 additions & 2 deletions examples/cuda/plotting.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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
}
10 changes: 7 additions & 3 deletions examples/cuda/stream.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down Expand Up @@ -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
}
4 changes: 3 additions & 1 deletion examples/cuda/surface.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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
}

0 comments on commit 74b5b4c

Please # to comment.