Day 4Module 1

in-technical-review

CUDA grid, block and thread indexing explained

Here is a launch that compiles, runs, returns zero, and writes three quarters of the array:

const int blocks = n / 256;   // n is 2000, so blocks is 7
writeOwner<<<blocks, 256>>>(d_blockOf, d_threadOf, n);

Seven blocks of 256 threads is 1792 threads for 2000 elements. The last 208 elements are never given to anybody. cudaGetLastError() returns cudaSuccess, the kernel returns, and if you spot-check element 1337 it is correct because 1337 sits inside the first 1792.

Only the tail is wrong, and nothing tells you.

Two questions are tangled together there, and both are today's subject: how a thread works out which element it owns, and how many threads you have to start so that every element has one. Work through this page and you can name the thread that touches element 1337 for any launch configuration, and write the kernel that proves you right.

Four numbers, and the one you build from them

A kernel launch does not create threads one at a time. The execution configuration, the <<<blocks, threads>>> between the kernel's name and its arguments, creates a grid of blocks, each holding the same number of threads. Every thread runs the same kernel body.

Values from the built-in index variables distinguish one thread from another.

There are four of them, and each thread sees its own values:

threadIdx.x is my index inside my block, from 0 to blockDim.x - 1.
blockIdx.x is my block's index in the grid, from 0 to gridDim.x - 1.
blockDim.x is how many threads are in a block. Every thread reads the same value.
gridDim.x is how many blocks are in the grid. Same for everyone.

All four are uint3 values with .x, .y and .z members. Today only .x matters; day 7 uses .y.

What is missing from that list is a number that is unique across the whole launch, and that is the one you build:

i = blockIdx.x * blockDim.x + threadIdx.x

Count the blocks in front of me, multiply by how many threads each one holds, add my position inside my own. With 256 threads per block, thread 57 of block 5 gets 5 * 256 + 57, which is 1337. That number is the global thread index, and on day 4 it is also the element index, because one thread owns one element.

Run it backwards and you have the question this lesson's exercise asks. Given an element i and a block size, the block is i / blockDim.x and the thread is i % blockDim.x.

One more number behaves differently from the others. Threads run in groups of 32 called warps, and a thread's lane is its position inside that group, threadIdx.x % 32. Because 32 divides every block size worth using, the lane is just i % 32.

Change the block size and the block number and the thread number both move; the lane does not. The block and thread numbers are a property of your launch. The lane is a property of the index.

Step the 1d-basic preset above one thread at a time and the formula stops being a line you copy. Click a cell to run it the other way and the widget names the thread that wrote it, which is the exercise at the bottom of this page done for you on a smaller array.

The grid is not a loop over your array

The intuition almost everyone arrives with is that <<<blocks, threads>>> is a loop header, and that CUDA will therefore cover the array the way for (i = 0; i < n; i++) covers it. It will not.

The launch starts exactly blocks * threads threads, no more and no fewer, and it has never seen your array. Whether those threads cover n elements is arithmetic you did, on the host, before the launch.

That has two consequences and they pull in opposite directions. You need ceiling division, (n + threads - 1) / threads, because truncating division leaves the last partial block out. And then you need the guard if (i < n) inside the kernel, because ceiling division starts more threads than you have elements.

The second habit is the argument order. <<<blocks, threads>>> reads left to right as "how many groups, of what size", and it is easy to write it the other way round:

// blocks is 8, so this starts 256 blocks of 8 threads. It covers every
// element, passes a correctness check, and wastes 24 lanes of every 32.
writeOwner<<<256, blocks>>>(d_blockOf, d_threadOf, n);

The swap does not always fail, which is exactly why it survives. Here it starts 2048 threads, covers all 2000 elements, and gives the right answer with 8-thread blocks that leave three quarters of every warp idle.

Day 10 measures what that costs. The swap only fails loudly when the block count climbs past 1024, and then it fails at the launch rather than in the results.

Asking the GPU which thread it was

The full program is code/day04-indexing/indexing.cu. It is small, and the three rules it follows matter more than its length.

The kernel computes nothing. Every thread writes its own blockIdx.x and threadIdx.x into the element it owns, and stops. There is no sum and no product for a wrong index to hide inside, so the output is the mapping itself rather than a result that happens to look plausible.

The buffers start at a value no thread can write. Both arrays are filled with -1 before every launch, so an element nobody owned stays visible. Fill them with zero and an untouched element is indistinguishable from one that block 0 thread 0 wrote, which is the number the first configuration exists to count.

The program checks the formula against the hardware. It rebuilds block * blockDim + thread from every recorded pair and exits non-zero on the first one that is not its own index. A page that prints a formula the hardware did not follow is worse than a page with no formula on it.

Here is the kernel:

// One thread writes its own two coordinates into the element it owns.
//
// Memory: consecutive threads in a warp have consecutive threadIdx.x, so one
// warp's 32 addresses cover 128 contiguous bytes of each output array. That is
// the coalesced pattern day 11 measures.
//
// Launch assumption: a 1D grid of 1D blocks. The guard is what makes an
// over-sized grid safe, so it is not optional.
__global__ void writeOwner(int* blockOf, int* threadOf, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        blockOf[i] = static_cast<int>(blockIdx.x);
        threadOf[i] = static_cast<int>(threadIdx.x);
    }
}

blockIdx.x and blockDim.x are both unsigned int, so the cast is not decoration: without it the multiply happens in 32 bits and wraps at 2^32 threads, which is 4 Gi floats and fits on hardware people rent.

And here are the two grid sizes the program runs, side by side:

    // Integer division truncates, so this grid is one block short whenever
    // kElems is not a multiple of the block size. Nothing warns you. The
    // launch succeeds, the kernel returns, and the tail of the array keeps
    // whatever was in it before.
    const int shortBlocks = static_cast<int>(kElems / kThreadsPerBlock);

    // Ceiling division. One more block than you strictly need, and the guard
    // inside the kernel switches off the threads that block over-provides.
    const int blocks =
        static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);

Note. The program records which thread wrote an element, not the order the threads ran in. Nothing here measures time, and thread numbers are addresses rather than a schedule. Day 21 records what actually ran when, with clock64() written to a device buffer.

Results

Re-verified without behavioral drift on a Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88) on 2026-09-02. The original CUDA 12.6 transcript remains beside the new run in the page's evidence array.

GPU: Tesla T4 (compute capability 7.5)
2000 elements, asking who owns element 1337

    7 x   256 =    1792 threads   unwritten  208   idle    0   element 1337 <- block 5, thread 57 (warp 1, lane 25)
    8 x   256 =    2048 threads   unwritten    0   idle   48   element 1337 <- block 5, thread 57 (warp 1, lane 25)
    2 x  1024 =    2048 threads   unwritten    0   idle   48   element 1337 <- block 1, thread 313 (warp 9, lane 25)

Every written element satisfied blockIdx.x * blockDim.x +
threadIdx.x == its own index. The only row that lost elements is the
one whose grid was sized with a truncating division.

The same element, 1337, gets a different owner when the block size changes. With 256-thread blocks, it belongs to block 5, thread 57. With 1024-thread blocks, it belongs to block 1, thread 313.

Both are correct because blockIdx.x * blockDim.x + threadIdx.x gives 1337. It stays in lane 25 because 57 and 313 are both 25 more than a multiple of 32.

The first row uses 7 x 256 = 1792 threads for 2000 elements, so it leaves 208 elements unwritten. The launch reports no error, and those 208 values keep the buffer's old contents.

The next two rows round the grid up and cover every element. Their extra 48 threads exit at the bounds check. Day 5 uses the same pattern.

The equation 1337 = 5 * 256 + 57 is arithmetic, not a measurement. The run checks that the hardware records the same owner.

It does so for every written element in all three launch configurations.

The transcript above has four parts: the device line, then one row per launch configuration, then the check. The row that matters is the first, because it is the one that quietly loses 208 elements.

Do not read the block and thread numbers as facts about CUDA. They are facts about one launch configuration. What carries to every other launch is the shape: one element has exactly one owner, that owner changes when the block size changes, the lane does not, and the number of elements with no owner at all is decided by a single division on the host.

Run it yourself

Compiler Explorer, because it costs nothing and needs no account. Two 8 KB arrays and three launches sits nowhere near Compiler Explorer's 20 second compile and 20 second run caps, and the whole program is one file carrying its own CUDA_CHECK, which is what an embed needs: it is a single editor buffer and cannot include a shared header.

nvcc -std=c++17 -O3 -arch=sm_75 -o indexing indexing.cu

Hardware. Target sm_75 or lower here: the Compiler Explorer runner is a Tesla T4 at compute capability 7.5, so a higher -arch compiles and then fails at run time. On Colab's free T4 or a card you own, the same line works unchanged.

Exercise

Part 1 is the whiteboard question a GPU kernel engineer interview opens with.

Part 1. Without running anything, name the block and the thread that touch element 1337 with 256 threads per block, with 512, and with 1024. Then say which number is the same in all three, and why.

Part 2. Write a kernel computing out[i] = in[i] + (float)i for every element, for an n you are not told in advance, and size the grid on the host.

Time: 25 to 40 minutes. Submit: your kernel and launch, plus the three answers.

Check: the harness runs N = 0, 1, 31, 32, 611, 1024 and 1025, the first case in integers stored as floats so a mismatch is an indexing bug and nothing else. On failure it prints the first ten wrong indices with the block and thread that owns each, and names the pattern: a tail bug is your bounds check, a constant shift is your formula. Look for the line saying you pass at 1024 and fail at 611.

Hint 1

Part 1 is the formula run backwards, from an element to a thread. Part 2 fixes the element count and the block size, so the only free number is the block count, and there are two candidates for it. Which elements does the smaller one never reach?

Hint 2

With 611 elements and 256 threads per block, what does 611 / 256 give, and how many elements does it leave with no thread? Now the other way: with the block count that covers 611, how many threads have no element, and what single line stops them writing anyway?

Solution

Part 1. 1337 = 5 * 256 + 57, and 2 * 512 + 313, and 1 * 1024 + 313: block 5 thread 57, block 2 thread 313, block 1 thread 313. The lane does not move.

1337 % 32 is 25 every time because each block size is a whole number of warps, which makes the lane i % 32 and nothing to do with the launch.

Part 2. The index line and the if (i < n) guard from indexing.cu, with out[i] = in[i] + static_cast<float>(i); inside, launched as (n + 255) / 256 blocks of 256. Day 4 is graded on correctness, not speed, and day 10 is where a block size first gets measured.

That formula does not change again for the rest of the course; what changes is how many elements one thread owns, several under the loop on day 8 and a whole tile from day 13.

Pitfalls

Most of the array is right and the tail is untouched. The grid was sized with n / threadsPerBlock, which drops the last partial block. Nothing fails and a spot check near the front passes.

Use (n + threadsPerBlock - 1) / threadsPerBlock, or the grid-stride loop on day 8, which sizes the grid to the machine instead of to n.

The program dies at a line that did not cause it. No bounds check, so the threads the ceiling grid over-provides read and write past the end of the array. The runtime string is an illegal memory access was encountered.

A launch is asynchronous, so it usually surfaces at the next cudaMemcpy or cudaDeviceSynchronize rather than at the launch.

The fix is if (i < n). Day 6 covers checking every call; day 61 covers finding it with compute-sanitizer.

The launch fails immediately with invalid configuration argument. That is cudaErrorInvalidConfiguration, code 9, and after a swapped <<<threads, blocks>>> it means the block count went past 1024, which is the maximum threads per block on every GPU this course targets.

The runtime reports it synchronously, so cudaGetLastError() on the line after the launch catches it. A zero in the grid, from a truncating division on a small n, produces the same string.

An int global index. blockIdx.x and blockDim.x are both unsigned int, so the product is computed in 32 bits and wraps past 2^32 threads. That is 4 Gi floats, 16 GiB, which fits on cards people rent by the hour.

The wrap is silent and the wrong answers land in the middle of the array. Cast to size_t before the multiply.

Reading execution order out of thread numbers. Thread 0 is not first and block 0 is not first. The numbers identify work; they do not define order.

Day 21 records what actually ran when.

Both quoted strings above are what cudaGetErrorString returns, captured on a Tesla T4 on 2026-08-29. The documentation describes those codes in different words than the runtime prints.

Go deeper

CUDA Programming Guide 2.3, "Writing SIMT Kernels": https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html (checked 2026-08-29)
CUDA Programming Guide, "Built-in Variables": https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html#built-in-variables (checked 2026-08-29)
Table 30, maximum threads per block and grid dimensions per compute capability: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-29)
The question this page answers, 178,647 views: https://stackoverflow.com/questions/2392250/understanding-cuda-grid-dimensions-block-dimensions-and-threads-organization-s (checked 2026-08-29)
PMPP 4th edition, chapter 3, multidimensional grids and data: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0 (checked 2026-08-29)

Day 5 puts this formula on 611 elements, where the last block runs 99 live threads and 157 idle ones, so the guard you just wrote decides whether the answer is right or merely plausible. Day 7 adds the second dimension. Day 11 is where the order of these addresses stops being a correctness question and becomes a bandwidth one, and the rest of module 1 is this formula under more pressure.