Cuda kernel - second block failed to launch?

Question

Here is my kernel call:

const size_t numBins = 1024;
const dim3 blockSize(512);
unsigned int * d_binHistogram, * d_cdf;
checkCudaErrors(cudaMalloc((void**) &d_binHistogram, sizeof(unsigned int)*numBins));
checkCudaErrors(cudaMemset(d_binHistogram, 0, sizeof(unsigned int)*numBins));
checkCudaErrors(cudaMalloc((void**) &d_cdf, sizeof(unsigned int)*numBins));
checkCudaErrors(cudaMemset(d_cdf, 0, sizeof(unsigned int)*numBins));
blellochScan<<<(numBins/512+1,1,1),blockSize,blockSize.x*sizeof(float)>>>(d_binHistogram, numBins, d_cdf);
cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());

Here is the kernel:

__global__ void blellochScan(const unsigned int* const d_bin, const size_t numBins,
                                  unsigned int* const d_cdf){
unsigned int i = blockIdx.x*blockDim.x+threadIdx.x;
unsigned int result = 0;
for(unsigned int j = 0; j<i; j++){
    result += d_bin[j];
}
d_cdf[i] = result; 
printf("%u %u\n", i, d_cdf[i]);
}

In the output, there are only i value from 0 to 511, there is no i value from 512 to 1023 of the second block. What did I do wrong?

Also, I am not sure of the third parameter in kernel launch blockSize.x*sizeof(float). How should I determine what to put in this third parameter?

I don't think this is your code. Did you just type this into the browser? you have blocksize with a lowercase s as your dim3 variable, and blockSize with an uppercase S as the variable you're using in your kernel call. This code seems to have a number of syntax errors, and is also incomplete. SO expects: "Questions concerning problems with code you've written must describe the specific problem — and include valid code to reproduce it — in the question itself. See SSCCE.org for guidance. " You haven't provided a SSCCE.org code, and the snippets you've provided don't make sense. — Robert Crovella
Thanks, I typed the first two lines so I mixed that syntax error. I have edited the code. Now it should be compilable. — Nam Thai
Read SSCCE.org. It's not just about compilable. It should be a complete code. Anyway, Roger's answer should sort out your code. Instead of (numBins/512+1,1,1) (which is not valid, the compiler should at least have indicated a warning on that line), you can pass dim3(numBins/512+1,1,1) (which is valid). However, that would end up launching 3 blocks, not 2, and since your kernel has no thread check, it would have other problems. Furthermore, printf may not be a reliable mechanism for large amounts of kernel output. — Robert Crovella
Thanks, I'll read SSCCE.org carefully. Currently I'm not entirely sure how I can put up the whole code. It has like 8 kernel launches to find the values of d_binHistogram — Nam Thai
It doesn't ask for your whole code either. It asks for a simplified version, reduced to the smallest example that will show the problem. Yes, it requires work. For these trivial examples, people like Roger can spot your errors for you. But for more complicated problems, we need to be able to compile and run the code to see what is happening for ourselves, in order to help you. — Robert Crovella

Roger Dahl Roger Dahl · Accepted Answer · 2013-12-08T16:08:31

In the future, please create a complete, compilable example that shows the issue you are asking about and copy/paste it into the question.

There is an error in your grid size calculation:

(numBins/512+1,1,1)

Will always evaluate to 1 (the value after the last comma) because you provided a comma separated list of expressions, not parameters for creating a dim3.

The third parameter in the kernel launch is number of bytes of shared memory to allocate. Since you're not using shared memory in your kernel, you can leave this parameter out.

Since your number of bins is divisible by the block size, you can further simplify your expression.

So, your kernel call should be:

const size_t numBins(1024);
const size_t blockSize(512);
dim3 blocks(numBins / blockSize);
blellochScan<<<blocks, blockSize>>>(...);

Cuda kernel - second block failed to launch?

1 Answers