Understanding this CUDA kernels launch parameters

Question

I am trying to analyze some code I have found online and I keep thinking myself into a corner. I am looking at a histogram kernel launched with the following parameters

histogram<<<2500, numBins, numBins * sizeof(unsigned int)>>>(...); 

I know that the parameters are grid, block, shared memory sizes.

So does that mean that there are 2500 blocks of numBins threads each, each block also having a numBins * sizeof(unsigned int) chunk of shared memory available to its threads?

Also, within the kernel itself there are calls to __syncthreads(), are there then 2500 sets of numBins calls to __syncthreads() over the course of the kernel call?

Answer 1

So does that mean that there are 2500 blocks of numBins threads each, each block also having a numBins * sizeof(unsigned int) chunk of shared memory available to its threads?

From the CUDA Toolkit documentation:

The execution configuration (of a global function call) is specified by inserting an expression of the form <<<Dg,Db,Ns,S>>>, where:

• Dg (dim3) specifies the dimension and size of the grid.
• Db (dim3) specifies the dimension and size of each block
• Ns (size_t) specifies the number of bytes in shared memory that is dynamically allocated per block for this call in addition to the statically allocated memory.
• S (cudaStream_t) specifies the associated stream, is an optional parameter which defaults to 0.

So, as @Fazar pointed out, the answer is yes. This memory is allocated per block.

Also, within the kernel itself there are calls to __syncthreads(), are there then 2500 sets of numBins calls to __syncthreads() over the course of the kernel call?

__syncthreads() waits until all threads in the thread block have reached this point. Is used to coordinate the communication between threads in the same block.

So, there is a __syncthread() call per block.