TL;DR: Give each thread a global index, guard
i < n, and launch enough blocks to cover the tail. For 1,000,003 values with 256 threads per block, 3,907 blocks cover the array; 189 threads in the last block skip the access. Separate kernel time from transfers and setup. A memory-bandwidth estimate is a lower bound, not proof that this kernel already reaches peak performance.
How to approach it
First state which input elements each thread reads and which output it writes, then write the kernel. Compute the rounded-up grid, including empty input. Then explain launch errors, asynchronous execution and resource limits. Finish with byte arithmetic for resident data versus an end-to-end host round trip.
A strong answer
A vector add works for a round array length but leaves the final values unchanged when the length changes. The likely error is floor division in the grid size. This is a useful first kernel because the output ownership is simple enough to prove: every valid index must have exactly one writer.
One thread owns one output
The input pointers refer to readable device arrays of at least n floats, and the output refers to a writable, non-overlapping array of that size. This deliberately strict alias contract matches __restrict__. The host validates allocation sizes before launching.
#include <cstddef>
__global__ void vecAdd(const float* __restrict__ a,
const float* __restrict__ b,
float* __restrict__ c, std::size_t n) {
const std::size_t i = static_cast<std::size_t>(blockIdx.x)
* blockDim.x + threadIdx.x;
if (i < n) c[i] = a[i] + b[i];
}
Widening before multiplication avoids doing the index product in a narrower type. On the host, skip the launch for n == 0. For a block size of 256, calculate the number of blocks as n / 256 + (n % 256 != 0); this avoids adding 255 to a near-maximum size value. Check the selected device's grid limit before converting that count to the launch dimension. The launch expression is vecAdd<<<blocks, 256>>>(a, b, c, n) after successful allocations and input copies.
NVIDIA's programming model defines the grid/block hierarchy. Blocks can run in any order; this kernel needs no barrier because each output depends only on its own two inputs. CUDA Programming Model and GPU Execution Model explain how the logical blocks map onto the device.
1,000,003 / 256 = 3,906 full blocks, remainder 67
ceiling grid = 3,907 blocks; launched threads = 1,000,192
final block starts at 3,906 × 256 = 999,936
valid threads = 67; skipped threads = 256 − 67 = 189
sanity: last valid index is n − 1 = 1,000,002, not n.
The same bounds rule applies even when today's array happens to divide evenly. For 100,000,000 values, 390,625 blocks cover the input exactly. A grid-stride loop is another useful design: it lets a bounded grid process more elements and tune work per thread, not merely work around old grid limits.
A legal launch is not a performance result
A multiple of 32 is a sensible starting block size because current CUDA warps contain 32 threads. A 1,000-thread block has 32 warps with 24 inactive lanes in its final warp; it may be legal if the device and this kernel's resource limits permit it. A supported maximum thread count does not guarantee that any kernel can launch at that size: registers, shared memory and other limits also matter.
Do not assume every block reaches theoretical residency or that a particular number of SMs implies a fixed schedule of waves. Check the device, compiled resource use and achieved behavior. For this kernel, compare reasonable block sizes while holding the data and timing method constant.
Check every allocation and copy result. Check the launch status promptly with cudaGetLastError(), then check completion at the synchronization or event boundary where the result is needed. Launch success is not evidence that asynchronous device execution completed without error. Validate the entire output and release both host and device allocations in a complete program.
Count transfers on the path you measure
For 100M FP32 elements, one array is 400 MB decimal. Reading two arrays and writing one output gives 1.2 GB logical device traffic and 1/12 ≈ 0.0833 FLOP/byte. With a hypothetical sustained device rate of 3 TB/s, the traffic-only estimate is 0.4 ms. It does not imply the ALUs are idle for a precisely calculable percentage of wall time.
If the two inputs start on the host and the result returns there, another 1.2 GB crosses the host/device path. At an assumed effective 25 GB/s with sequential transfers, that is 48 ms, before setup and other overhead. A 0.4 ms kernel would then be about 0.83% of the 48.4 ms transfer-plus-kernel subtotal. These rates are illustrative inputs, not measurements on this machine or promises for PCIe Gen5.
Memory-Bound vs Compute-Bound Kernels explains the intensity calculation. GPU Memory Hierarchy helps separate logical loads from actual cache and HBM traffic. Time allocation/setup, copies and the kernel separately, using device events for the device work and an appropriate end-to-end timer for the user-visible operation.
Keeping data resident across a larger computation can remove repeated host round trips. Pinned buffers and asynchronous operations can enable overlap when hardware, streams and dependencies permit; they do not make every copy overlap automatically. For a small isolated add, a CPU implementation may be the better choice. Reverse that choice when the surrounding GPU workload already owns the arrays or measured end-to-end performance justifies moving them.
Verification scope: the kernel's index/write logic was executed in a CPU harness with empty, exact-block and tail cases plus output guards. Native CUDA compilation, device execution and timing require a working CUDA environment; no GPU benchmark is claimed here.
What interviewers probe next
- "Why not
n / block?" It omits the partial block and leaves the tail unwritten. - "Does launch success mean the output is ready?" No. Respect stream ordering and a checked completion boundary.
- "Can you optimize the add?" Possibly, but first identify whether device traffic, launches, transfers or setup dominate the measured path.
Common mistakes
- Accessing index
n, launching an empty grid, or overflowing size arithmetic. - Treating peak bandwidth and maximum residency as achieved values.
- Timing only the kernel when the product pays for copies too.
Key takeaways
- 1,000,003 values need 3,907 blocks of 256, with 189 guarded tail threads.
- The 100M-element operation moves 1.2 GB logically on the device.
- The toy 48 ms transfer estimate changes the optimization priority; actual rates must be measured.