Cuda

Thread hierarchy
├── grid (1D/2D/3D)
│   └── block (1D/2D/3D, max 1024 threads)
│       └── thread (threadIdx, blockIdx, blockDim, gridDim)

// vector_add.cu
#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vector_add(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];
}

int main(void) {
    const int n = 1 << 20;
    size_t bytes = n * sizeof(float);
    float *h_a, *h_b, *h_c, *d_a, *d_b, *d_c;

    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    // ... host init and cudaMemcpy H2D ...

    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);
    cudaDeviceSynchronize();

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

#define TILE 16

__global__ void matmul_tiled(const float *A, const float *B, float *C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];

    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; t++) {
        As[threadIdx.y][threadIdx.x] = (row < N && t * TILE + threadIdx.x < N)
            ? A[row * N + t * TILE + threadIdx.x] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] = (col < N && t * TILE + threadIdx.y < N)
            ? B[(t * TILE + threadIdx.y) * N + col] : 0.0f;
        __syncthreads();

        for (int k = 0; k < TILE; k++)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();
    }
    if (row < N && col < N)
        C[row * N + col] = sum;
}

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, h_b, bytes, cudaMemcpyHostToDevice, stream2);
vector_add<<<blocks, threads, 0, stream1>>>(d_a, d_b, d_c, n);
cudaMemcpyAsync(h_c, d_c, bytes, cudaMemcpyDeviceToHost, stream1);
cudaStreamSynchronize(stream1);

# Single architecture (local GPU)
nvcc -O3 -arch=sm_80 -o prog vector_add.cu

# Fat binary for multiple GPUs
nvcc -O3 \
  -gencode arch=compute_80,code=sm_80 \
  -gencode arch=compute_90,code=sm_90 \
  -o prog vector_add.cu

# Debug symbols for cuda-gdb
nvcc -G -g -O0 -arch=sm_80 -o prog_debug vector_add.cu

# Show PTX/SASS
nvcc -arch=sm_80 -ptx vector_add.cu
nvcc -arch=sm_80 -cubin vector_add.cu
cuobjdump -sass prog

# CUDA Occupancy Calculator (spreadsheet) or programmatic:
./occupancy_tool --kernel vector_add --block-size 256 --regs 16 --smem 0

Kernel slow?
├── Low occupancy (< 25%) → reduce registers, shared mem, or block size
├── Memory-bound → check coalescing, use shared memory tiling
└── Compute-bound → increase arithmetic intensity, use tensor cores