Vector-add kernel

One element per thread

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

blockIdx.x * blockDim.x + threadIdx.x is the global index, and the only line in the kernel that knows a grid exists. Delete the guard and the kernel is correct for every n that divides the block size, which is exactly the size your first test uses.

Size the grid from n

const int threads = 256;
const int blocks = (n + threads - 1) / threads;
vadd<<<blocks, threads>>>(a, b, c, n);
cudaError_t err = cudaGetLastError();
cudaDeviceSynchronize();

For n = 1000, blocks is 4 and the grid launches 1024 threads. The last block runs 232 threads of work inside 256 lanes. Rounding up is what creates the guard; rounding down would silently drop 232 elements.

The index names and launch syntax here are stable across recent toolkits, but check the guide for your release if you use cluster dimensions or __grid_constant__ parameters.

The guard is the contract

Block 3 covers indices 768 to 1023, and threads 1000 to 1023 have no element to load. Without the check they write 24 floats past the end of c. Nothing about the launch is invalid, so cudaGetLastError is content; the damage lands in whatever allocation follows and surfaces later, often inside an unrelated kernel.

Things that break

A missing guard is not a crash, it is a write past the buffer, and compute-sanitizer --tool memcheck names the line. A block size that is not a multiple of 32 leaves lanes idle in the last warp of every block. An int index overflows above 2^31 elements, which reads as a wrong answer on a huge input rather than an error. A grid dimension beyond the limit is a launch error, which is one of the few failures you will actually see at cudaGetLastError.