Your First Kernel

Variable qualifiers: __device__ (global memory variable), __constant__, __shared__, __managed__ (unified memory), __grid_constant__.

4.1 Hello from the GPU

#include <cstdio>

__global__ void hello() {
    printf("Hello from block %d, thread %d\n", blockIdx.x, threadIdx.x);
}

int main() {
    hello<<<2, 4>>>();          // launch 2 blocks × 4 threads = 8 threads
    cudaDeviceSynchronize();    // wait for the GPU (kernel launches are asynchronous)
    return 0;
}

nvcc hello.cu -o hello && ./hello

4.2 Function qualifiers

host (and device with dynamic parallelism) — a kernel, returns void host (default) __host__ __device__ compiled for both

Variable qualifiers: __device__ (global memory variable), __constant__, __shared__, __managed__ (unified memory), __grid_constant__.

4.3 Vector addition — the canonical example

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

__global__ void vecAdd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // global thread index
    if (i < n) {                                     // bounds check
        c[i] = a[i] + b[i];
    }
}

int main() {
    const int n = 1 << 20;                           // 1M elements
    size_t bytes = n * sizeof(float);

    // 1. Allocate host memory
    float *h_a = (float*)malloc(bytes), *h_b = (float*)malloc(bytes), *h_c = (float*)malloc(bytes);
    for (int i = 0; i < n; ++i) { h_a[i] = 1.0f; h_b[i] = 2.0f; }

    // 2. Allocate device memory
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 3. Copy inputs host → device
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 4. Launch the kernel
    int threads = 256;
    int blocks = (n + threads - 1) / threads;        // ceiling division
    vecAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);

    // 5. Copy result device → host (cudaMemcpy synchronizes with the kernel)
    cuda
…

4.4 The CUDA workflow

Allocate host & device memory.

Copy input data host → device.

Launch kernel(s).

Copy results device → host.

4.5 Grid-stride loops (robust kernels for any size)

Decouples the grid size from the problem size and lets one launch handle arbitrarily large arrays.

__global__ void vecAddStride(const float* a, const float* b, float* c, int n) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += blockDim.x * gridDim.x) {
        c[i] = a[i] + b[i];
    }
}
// Launch with a fixed number of blocks (e.g., 4 × number of SMs)