Counting Threads in a Launch

__global__ void scale(float* x, int n, float a) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) x[i] = a * x[i];
}

int n       = 1000000;
int threads = 256;
int blocks  = (n + threads - 1) / threads;   // 3907
scale<<<blocks, threads>>>(d_x, n, 2.0f);

How many threads does this launch start, how many warps is that, and what do the leftover threads do?

A

3,907 threads, one per block, each looping over its own 256 elements.

B

1,000,192 threads, and the extra 192 write past the end of x, since the guard runs after the index has already been used.

C

3,907 blocks x 256 threads = 1,000,192 threads, which is 31,256 warps of 32; the last 192 threads fail the i < n guard and exit without touching memory.

D

Exactly 1,000,000 threads, because the guard stops the extra threads from being created.

E

1,000,192 threads in 3,907 warps, since a warp is the hardware name for a block.