Write a CUDA vector add and explain the launch: grid math, the bounds check, and why the copies dominate.

TL;DR: Each thread computes one element: its global index is blockIdx.x * blockDim.x + threadIdx.x, it checks i < n because the grid is rounded up to whole blocks, and it writes c[i] = a[i] + b[i]. Launch with 256 threads per block and (n + 255) / 256 blocks; for 100M floats that is 390,625 blocks, which the hardware schedules onto the SMs in waves. The kernel reads 800 MB and writes 400 MB; on an H100 at 3.35 TB/s that is about 0.4 ms. The three copies (two arrays in, one out, 1.2 GB) over PCIe Gen5 at about 25 GB/s realized take about 50 ms, so the add is under 1% of the wall-clock; the point of the exercise is to say so, and to know that the fix is keeping data on the device across many kernels, not making the add faster.

How to approach it

Write the kernel and the host code, then explain each launch decision (block size, grid size, bounds check, error check). Then do the bandwidth arithmetic for the kernel and for the copies. Close with what the numbers mean for real workloads.

A strong answer

A typical situation: the interviewer gives 10 minutes for the kernel and 20 for questions about it; the candidate who cannot compute the grid or explain why the timing shows 50 ms for a 0.4 ms kernel does not pass.

The code:

#include <cuda_runtime.h>
#include <cstdio>

__global__ void vecAdd(const float* __restrict__ a, const float* __restrict__ b,
                       float* __restrict__ c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // global index: block offset + lane within block
    if (i < n) c[i] = a[i] + b[i];                    // last block may extend past n
}

#define CHECK(x) do { cudaError_t e = (x); if (e != cudaSuccess) { \
    fprintf(stderr, "%s:%d %s\n", __FILE__, __LINE__, cudaGetErrorString(e)); exit(1); } } while (0)

int main() {
    const int n = 100'000'000;
    size_t bytes = n * sizeof(float);
    float *ha, *hb, *hc;
    CHECK(cudaMallocHost(&ha, bytes)); CHECK(cudaMallocHost(&hb, bytes)); CHECK(cudaMallocHost(&hc, bytes)); // pinned
    for (int i = 0; i < n; ++i) { ha[i] = 1.0f; hb[i] = 2.0f; }

    float *da, *db, *dc;
    CHECK(cudaMalloc(&da, bytes)); CHECK(cudaMalloc(&db, bytes)); CHECK(cudaMalloc(&dc, bytes));
    CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));

    int block = 256;
    int grid = (n + block - 1) / block;              // ceiling division: 390,625 blocks
    vecAdd<<<grid, block>>>(da, db, dc, n);
    CHECK(cudaGetLastError());                       // launch errors (bad config) surface here
    CHECK(cudaDeviceSynchronize());                  // runtime errors (bad address) surface here

    CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
    printf("%f\n", hc[n - 1]);                       // 3.0
    return 0;
}

The launch, decision by decision:

block size 256: a multiple of the warp size (32) so no partial warps; small enough that an SM can hold several blocks
  (an H100 SM holds up to 2,048 threads → 8 blocks of 256) and large enough to amortize block scheduling; 128 to 512
  are all fine for this kernel; the answer is "a multiple of 32, tuned by measurement"
grid size: ceil(n / block) = (100,000,000 + 255) / 256 = 390,625 blocks; the grid's x dimension allows up to 2^31 − 1,
  so no need for a 2D grid; older code used a grid-stride loop when grids were limited
bounds check: 390,625 × 256 = 100,000,000 exactly here, but in general the last block has threads past n, and
  without the check they read and write past the arrays; the check costs one compare per thread
scheduling: the GPU has 132 SMs; each holds 8 blocks of 256 → 1,056 blocks resident at a time; 390,625 / 1,056 ≈ 370
  waves; the tail wave is nearly full, so no tail effect to speak of
__restrict__: tells the compiler the pointers do not alias, so it can schedule the loads freely; const on inputs
  lets it use the read-only path
error checking: cudaGetLastError after the launch catches configuration errors; the sync catches faults inside the
  kernel; without both, a failing kernel is silent and the output is stale memory

CUDA Programming Model has the hierarchy and the index arithmetic; GPU Execution Model has the SM's warp scheduling that the block size interacts with.

The arithmetic that the interviewer wants:

kernel traffic: read a and b (2 × 400 MB) + write c (400 MB) = 1.2 GB
H100 HBM at ~3.35 TB/s peak, ~3.0 TB/s achievable → 1.2 GB / 3.0 TB/s ≈ 0.4 ms
arithmetic intensity: 1 FLOP per 12 bytes ≈ 0.08 FLOP/B, against a ridge of ~300 FLOP/B → bandwidth-bound by 3,000×;
  the ALUs are idle 99.97% of the time and that is fine
copies: 3 × 400 MB over PCIe Gen5 x16 (64 GB/s theoretical, ~25 to 50 GB/s realized with pinned memory) → 1.2 GB /
  25 GB/s ≈ 48 ms; with pageable memory the runtime stages through a pinned buffer and it is slower still
total ≈ 48 ms, of which the kernel is 0.4 ms (under 1%)
sanity: an add on the CPU at ~20 GB/s of memory bandwidth takes ~60 ms; the GPU version is not faster end to end for
        a single add; it is faster when the data stays on the device across many operations, which is the lesson

Memory-Bound vs Compute-Bound Kernels and Roofline Model place the kernel on the roofline; GPU Memory Hierarchy explains why the HBM number is the one that matters here.

What follows from the numbers in practice: fuse elementwise operations so intermediate results never round-trip to HBM (Kernel Fusion); keep tensors on the device across the whole computation and copy once; use pinned host memory and asynchronous copies on streams to overlap transfer with compute when copies are unavoidable; and measure with events around the kernel and around the copies separately, because a single wall-clock number hides which one you are looking at.

The reversal condition: if the arrays are tiny (a thousand elements), the kernel launch overhead (a few microseconds) and the copies dominate even more, and the answer is to not use the GPU for that operation at all. ptxas -v on the build and compute-sanitizer on the first run are the two checks that make a launch trustworthy.

What interviewers probe next

• "What happens if you launch with 1,000 threads per block?" It launches and it is correct. 1,000 is under the 1,024 limit, so the only cost is efficiency: the block rounds up to 32 warps and the last warp runs 24 of its 32 lanes idle. Multiples of 32 are a performance rule, not a legality rule. 1,025 is the one that fails, with an invalid configuration error caught by cudaGetLastError.
• "Why 256 and not 1,024?" Both work; 1,024-thread blocks limit how many blocks fit per SM and reduce scheduling flexibility; measure. For this kernel the difference is noise.
• "How would you make the add faster?" You cannot, meaningfully; it is at the bandwidth ceiling. Vectorized loads (float4) help reach the ceiling on some architectures; fusing it with neighbours removes it.
• "Where does the time go if the arrays are already on the device?" 0.4 ms for the kernel; the interviewer wants the candidate to have separated the two numbers already.

Common mistakes

• Forgetting the bounds check.
• Grid computed as n / block (drops the tail).
• No error check after the launch, and a silent failure.
• Timing the whole program and concluding the GPU is slow.

Key takeaways

• Index = block offset + lane; bounds check because the grid is rounded up; block size a multiple of 32; grid = ceil(n / block).
• Check launch and runtime errors separately.
• The add is bandwidth-bound at ~0.4 ms for 100M floats; the copies take ~50 ms; the fix is to keep data on the device, not to tune the add.