Optimization Techniques

Launch exactly enough blocks to fill the GPU; each block loops pulling work items from a global queue (atomic counter).

22.1 Checklist (in rough priority order)

Maximize parallelism and keep the GPU busy (enough threads, async work, concurrency).

Minimize host↔device transfers.

Coalesce global memory accesses; use SoA, vectorized loads.

Reuse data via shared memory / registers (tiling); avoid bank conflicts.

Avoid warp divergence in hot paths.

Kernel fusion: combine element-wise kernels to avoid round trips through global memory.

Use the right precision: FP16/BF16/TF32/FP8 with Tensor Cores; avoid accidental FP64 (1.0 vs 1.0f).

Instruction optimizations: fast math intrinsics, __fmaf_rn, rsqrtf, avoid integer division/modulo (use shifts/masks for powers of 2).

Loop unrolling (#pragma unroll), increase ILP (multiple independent accumulators per thread).

Tune launch configuration and occupancy; control register usage.

Use libraries (cuBLAS, cuDNN, CUB, CUTLASS) — they encode years of tuning.

CUDA graphs to cut launch overhead.

22.2 Fast math

__sinf(x), __cosf(x), __expf(x), __logf(x), __powf(x, y), __fdividef(a, b)   // fast, less accurate intrinsics
rsqrtf(x), fmaf(a, b, c), sincosf(x, &s, &c), __saturatef(x)
// nvcc --use_fast_math: maps standard functions to intrinsics, flushes denormals, approximate division/sqrt

22.3 Instruction-level parallelism example

// One accumulator: each add waits for the previous one
float s = 0; for (int i = tid; i < n; i += stride) s += x[i];
// Four independent accumulators → more in-flight operations per thread
float s0 = 0, s1 = 0, s2 = 0, s3 = 0;
for (int i = tid; i + 3 * stride < n; i += 4 * stride) {
    s0 += x[i]; s1 += x[i + stride]; s2 += x[i + 2 * stride]; s3 += x[i + 3 * stride];
}
float s = s0 + s1 + s2 + s3;   // (plus tail handling)

22.4 Kernel fusion example

// Unfused: three kernels, 3 reads + 3 writes of the array
scale<<<...>>>(x, a); addBias<<<...>>>(x, b); relu<<<...>>>(x);
// Fused: one read + one write
__global__ void fused(float* x, float a, float b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) x[i] = fmaxf(fmaf(x[i], a, b), 0.0f);
}

22.5 Persistent kernels & work queues

Launch exactly enough blocks to fill the GPU; each block loops pulling work items from a global queue (atomic counter). Reduces launch overhead and balances irregular workloads.

22.6 L2 cache persistence (Ampere+)

Keeps frequently reused data resident in L2.

cudaStreamAttrValue attr = {};
attr.accessPolicyWindow.base_ptr = d_lookup;
attr.accessPolicyWindow.num_bytes = bytes;
attr.accessPolicyWindow.hitRatio = 1.0f;
attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);