100 DAYS OF GPU Course preview
Day 10Module 1
in-technical-review

How to choose threads per block and blocks per grid

Someone on NVIDIA's forums, who had read the spec sheet for their card and done the arithmetic:

"I thought that the best and most optimum launch params for my kernel was <<<32, 768>>> (because it assigns 2 blocks to each SM, uses all SMs, and the maximum number of resident threads). However, when I run my kernel with <<<16, 1024>>>, it takes less time to execute. I can't explain this behavior."

https://forums.developer.nvidia.com/t/confusion-about-setting-kernel-block-and-grid-size-for-maximum-occupancy/287857 (checked 2026-08-30)

The arithmetic and the measurement are both right. The measurement decides which launch is faster.

This question has no closed-form answer. The Stack Overflow version had 163,053 views on 2026-08-30 (https://stackoverflow.com/questions/9985912/how-do-i-choose-grid-and-block-dimensions-for-cuda-kernels , view count from the Stack Exchange API).

A few constraints cut the candidates down to about six. One API gives you a sound default, and a twenty-minute sweep finds the best range for your kernel and GPU.

What the hardware does with the two numbers you pass

The launch configuration has two numbers with different roles. Blocks per grid should keep every multiprocessor busy: "The number of blocks in a grid should be larger than the number of multiprocessors so that all multiprocessors have at least one block to execute", and "To scale to future devices, the number of blocks per kernel launch should be in the thousands" (CUDA C++ Best Practices Guide 11.3, https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#thread-and-block-heuristics , checked 2026-08-30).

Use a grid-stride loop, as in day 8. Then the grid size can suit the GPU instead of matching the data size.

The hardware schedules threads per block in whole warps. A block runs on one SM, which assigns it ceil(T / 32) warp slots.

A 100-thread block takes four slots. Its fourth warp has 28 inactive lanes for the block's lifetime. The guide therefore says: "The number of threads per block should be a multiple of 32 threads, because this provides optimal computing efficiency and facilitates coalescing" (same section).

Two limits now apply. Each SM supports a fixed number of resident warps and resident blocks. Day 2 measured 32 warp slots and 16 block slots on a Tesla T4.

With 32 threads per block, you would need 32 blocks to fill the warp slots, but only 16 blocks fit. At 64 threads, the two limits meet. This is why the guide says: "A minimum of 64 threads per block should be used, and only if there are multiple concurrent blocks per multiprocessor."

Above 64 the ceiling stops being flat. Robert Crovella, answering someone who found this in a profiler and could not explain it:

"Now suppose I reduce my threadblock size to 992. That is 31 warps. I can still schedule at most 2 of these per SM (you cannot schedule 3- that would be 93 warps, or over 2900 threads). However, since I can schedule at most 2 of these, and each consists of 31 warps, the maximum warp load I can have is 62 warps - not 64. If I further reduce my threadblock size, the maximum achievable warp load will continue to decrease, until such point at which scheduling 3 blocks becomes feasible. Then the maximum achievable warp load will "jump up" again."

https://forums.developer.nvidia.com/t/question-about-threads-per-block-and-warps-per-sm/77491 (checked 2026-08-30)

Integer division causes this pattern. Block sizes that divide the warp cap evenly fill it; other sizes leave slots unused. On a 32-slot T4, this occurs at 96, 160, 192, 384, and 768 threads.

The t4-vs-5090-vs-h100 preset shows why one published block size may not suit another GPU. The widget models one SM, not run time. Read its "what this does not show" list first.

Occupancy helps hide latency; it does not predict speed

Occupancy is the fraction of those slots you fill. It gives the scheduler another warp to issue while one waits on memory. This is latency hiding.

The guide states the limit: "Higher occupancy does not always equate to higher performance-there is a point above which additional occupancy does not improve performance. However, low occupancy always interferes with the ability to hide memory latency, resulting in performance degradation" (Best Practices Guide 11.1, https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#occupancy , checked 2026-08-30).

Occupancy must be high enough to hide latency. Once it reaches that point, more warps may not help. They can also leave fewer registers per thread: "A lower occupancy kernel will have more registers available per thread than a higher occupancy kernel, which may result in less register spilling to local memory" (11.3).

Day 45 tunes one kernel to the same speed at two occupancies. That result explains the forum post at the top of this page.

You can compute the warp and block limits from the device specifications. Registers add another limit: threads times registers per thread must fit the SM's register file. Register pressure often sets a kernel's residency, and day 17 reads the count from ptxas.

Shared memory adds the fourth limit. More shared memory per block means fewer resident blocks. If either resource exceeds the device limit, the launch returns too many resources requested for launch.

Sweeping one variable at a time

The full program is in code/day10-block-size/block_size.cu. It sweeps thirteen block sizes over one 4K frame of grayscale conversion.

Unlike day 7's kernel, this one uses three planar float channels instead of interleaved unsigned char RGB and integer weights. That keeps the inner loop memory-bound. It uses the grid-stride loop from day 8, so correctness does not depend on the launch configuration.

The sweep follows four rules. The first depends on day 8.

The kernel is launch-configuration independent. The grid-stride loop covers every pixel for any grid and block size. A one-thread-per-element kernel changes its grid whenever the block size changes, so the benchmark could not isolate either value.

__global__ void grayscale(const float* r, const float* g, const float* b,
                          float* out, size_t n) {
    const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
    for (size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
         i < n; i += step) {
        out[i] = kRedWeight * r[i] + kGreenWeight * g[i] + kBlueWeight * b[i];
    }
}

The grid is one full wave, and the driver sizes it. Blocks per SM comes from cudaOccupancyMaxActiveBlocksPerMultiprocessor, which answers about the compiled kernel and therefore accounts for a register count your source never shows you.

        int blocksPerSM = 0;
        CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
            &blocksPerSM, grayscale, t, 0));

        size_t wave = static_cast<size_t>(blocksPerSM) *
                      static_cast<size_t>(prop.multiProcessorCount);
        const size_t needed =
            (kPixels + static_cast<size_t>(t) - 1) / static_cast<size_t>(t);
        if (wave > needed) {
            wave = needed;
        }
        if (wave < 1) {
            wave = 1;
        }
        const int gridBlocks = static_cast<int>(wave);

Correctness is checked at every block size. One check at 256 threads would not prove that every launch covers the image. The program clears the output buffer before each row, so an incomplete launch cannot reuse the last row's output.

Time with events, warm up every kernel you time. Three warm-ups discarded, the mean of ten timed runs, launches only inside the timed region. Day 9 covers why a host clock measures the wrong thing.

The program also asks the API for its choice, which gives the sweep a default for comparison:

    int suggestedGrid = 0;
    int suggestedBlock = 0;
    CUDA_CHECK(cudaOccupancyMaxPotentialBlockSize(&suggestedGrid,
                                                  &suggestedBlock, grayscale));

Note. Pinning the grid to one wave is a measurement choice and it contradicts the guide's "in the thousands" advice on purpose. That advice is about a grid sized to the data, which has to stay big enough for a future card with more SMs; a grid-stride kernel gets the same portability from the loop instead. The cost of the choice is that this program never sees the tail effect, where a grid of 1.1 waves spends its last tenth running on almost nothing.

Results

Re-verified 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 timing remains in the page's evidence array. The fastest block size and nine-of-thirteen plateau held; the block below is the CUDA 13.0 run.

GPU: Tesla T4 (compute capability 7.5)
  SMs 40, resident warps/SM 32, resident blocks/SM 16
Image: 3840 x 2160, 8294400 pixels, 126.6 MiB moved per pass
Timing: mean of 10 runs after 3 warm-ups, CUDA events

 thr/blk  warps  idle  blk/SM  warps/SM    occ    grid        ms     GB/s  of best
      16      1    16      16        16    50%     640     0.906    146.5     0.60
      32      1     0      16        16    50%     640     0.618    214.7     0.89
      64      2     0      16        32   100%     640     0.556    238.9     0.98
      96      3     0      10        30    94%     400     0.561    236.5     0.97
     100      4    28       8        32   100%     320     0.619    214.3     0.88
     128      4     0       8        32   100%     320     0.550    241.1     0.99
     160      5     0       6        30    94%     240     0.562    236.3     0.97
     192      6     0       5        30    94%     200     0.558    237.9     0.98
     256      8     0       4        32   100%     160     0.547    242.6     1.00
     384     12     0       2        24    75%      80     0.567    234.2     0.97
     512     16     0       2        32   100%      80     0.556    238.8     0.98
     768     24     0       1        24    75%      40     0.582    228.1     0.94
    1024     32     0       1        32   100%      40     0.574    231.1     0.95

cudaOccupancyMaxPotentialBlockSize suggests 1024 threads per block and a grid of 40 blocks.
Fastest measured: 256 threads per block at 0.547 ms.
9 of 13 configurations reach at least 95 percent of that row.

Nine of thirteen configurations land within 5 percent of the best one. Any tested size from 64 threads up performs well.

The two slow rows waste warp slots. Sixteen threads per block reaches 60 percent of the best result, and 32 reaches 89 percent. In both cases, the per-SM block limit applies before the warp limit.

The 100-thread row shows the cost of a partial warp. It is not a multiple of 32, so each block takes four warp slots and 28 lanes of the fourth sit idle for the block's whole life. It measures 0.619 ms against 0.550 for 128 threads, which is the same work done 12 percent slower because of the launch size.

cudaOccupancyMaxPotentialBlockSize suggested 1024, but 1024 is not the fastest. It measures 0.574 ms against 0.547 ms at 256 threads. The API optimises for occupancy, but this kernel is memory-bound, so maximum occupancy does not give maximum speed. Day 45 examines that trade.

Start with a multiple of 32 between 64 and 256. Then focus on memory access. Measure if the last few percent matter, because the best size depends on the kernel and GPU.

Run it yourself

The program allocates 126.6 MiB on the device and 158.2 MiB on the host and runs 169 timed launches, which is more than a shared Compiler Explorer slot should be asked for. Use a free Colab T4 or a card you own. The build line is the one in the repo's README:

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

Add -Xptxas -v to print the register count used by the occupancy calculation as ptxas info : Used N registers. A smaller version with three block sizes and a 512 by 512 image fits within Compiler Explorer's 20-second limits when you target sm_75 or lower.

Exercise

Run the sweep on your own GPU, then answer in two sentences: where does the curve go flat, and why does it stop improving there rather than at 100 percent occupancy?

Time: 25 to 40 minutes. Submit: your table, the two sentences, and the block size you would now ship.

Check: the program is its own check. It compares the output against a CPU reference at all thirteen block sizes and returns EXIT_FAILURE on the first row that disagrees, so a pass means the kernel really did cover the image at every launch configuration.

It then prints the fastest row, what cudaOccupancyMaxPotentialBlockSize suggested, and how many of the thirteen reach at least 95 percent of the fastest. That last count is what you are being graded on understanding.

Hint 1

Your two smallest block sizes report the same occupancy and should not report the same time. Ask what occupancy counts, and what those two rows do not share.

Hint 2

Compare the occupancy column against the of best column row by row. Where the two stop moving together, occupancy has stopped being the thing that limits you. What is left over on a kernel that reads three floats and writes one for every two multiply-adds?

Solution

Flat across most of the range, falling off only at the bottom, and the bottom fails for two reasons rather than one.

At 16 and 32 threads the per-SM block cap binds before the warp cap. Sixteen blocks is all you get, so 16 of 32 warp slots stay empty however large the grid is, and that shortage shows up as time. Sixteen threads is worse than 32 at the same occupancy, because occupancy counts warp slots and a 16-thread block runs its slot with half the lanes switched off.

From 64 up you have enough warps to cover the memory latency and more buy nothing. The sawtooth rows still lose slots to integer division, and they should lose much less time than that, because the deficit comes out of surplus: this kernel moves 16 bytes per pixel for two fused multiply-adds, so it is bandwidth-limited long before it is warp-limited.

Occupancy supplies warps that can run while others wait. Ask whether it is high enough to hide latency. Report ratios between rows, because the absolute values depend on the GPU.

Pitfalls

Picking a block size from a blog post. Every rule of thumb was measured on somebody else's card with somebody else's kernel. A T4 has 32 warp slots per SM, an RTX 5090 has 48, an H100 has 64 (FACT-SHEET.md section 3), so the size that fills one strands slots on another.

Reading "between 128 and 256" as an answer. The guide's exact words are "Between 128 and 256 threads per block is a good initial range for experimentation with different block sizes" (Best Practices Guide 11.3, checked 2026-08-30). That is the start of a search, and the same section says plainly that "inevitably some experimentation is required".

Chasing 100 percent occupancy. It is a means, and the guide says so: improving occupancy "from 66 percent to 100 percent generally does not translate to a similar increase in performance" (11.3). Day 45 measures two configurations that run at the same speed with different occupancy.

A block size that is not a multiple of 32. The launch succeeds without a warning, and every block runs its last warp with 32 - (T mod 32) idle lanes. At 100 threads, 28 of 128 lanes are idle.

A block bigger than the device allows. The product of the block dimensions has to stay at or below 1024 on every compute capability this course covers, and a dim3(32, 32, 2) block quietly asks for 2048. The launch returns invalid configuration argument, which day 5's cudaGetLastError() catches one line later and a program without that check never catches.

Sweeping a kernel that is not launch-configuration independent. If the grid is ceil(n / T), every row changes two things and the table means nothing. Rewrite it as a grid-stride loop first, which is day 8.

Go deeper

Next

That closes module 1. You can write a kernel, check it, size a launch, and time it.

Day 11 tests memory access, which stayed fixed throughout this sweep. Day 12 makes the access strided, and day 13 fixes it. Day 42 uses a profiler to measure achieved occupancy.