Application Case Studies

Decode (nvJPEG) → convert color space → Gaussian blur (separable convolution: horizontal then vertical pass with shared memory) → Sobel edge detection → histogram equalization (histogram + scan) → encode.

38.1 Image processing pipeline

Decode (nvJPEG) → convert color space → Gaussian blur (separable convolution: horizontal then vertical pass with shared memory) → Sobel edge detection → histogram equalization (histogram + scan) → encode. Use streams to pipeline batches and pinned memory for I/O.

__global__ void sobel(const unsigned char* in, unsigned char* out, int w, int h) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < 1 || y < 1 || x >= w - 1 || y >= h - 1) return;
    auto P = [&](int dx, int dy) { return (int)in[(y + dy) * w + (x + dx)]; };
    int gx = -P(-1,-1) - 2*P(-1,0) - P(-1,1) + P(1,-1) + 2*P(1,0) + P(1,1);
    int gy = -P(-1,-1) - 2*P(0,-1) - P(1,-1) + P(-1,1) + 2*P(0,1) + P(1,1);
    out[y * w + x] = (unsigned char)min(255, (int)sqrtf((float)(gx * gx + gy * gy)));
}

38.2 N-body simulation

All-pairs gravitational forces, O(N²) work: each block loads tiles of body positions into shared memory; each thread accumulates the force on its body over all tiles (rsqrtf for 1/r³); integrate with leapfrog. Classic GPU Gems 3 example reaching a large fraction of peak FLOP/s. Extensions: Barnes–Hut octrees (O(N log N)), multi-GPU.

38.3 Deep learning primitives

GEMM — the core of fully connected layers & attention (cuBLASLt, CUTLASS).

Convolutions — im2col + GEMM, implicit GEMM, Winograd, FFT (cuDNN).

Softmax / LayerNorm — warp/block reductions fused with element-wise ops.

Attention — FlashAttention: tiles Q, K, V in shared memory, computes softmax incrementally (online softmax), never materializes the N×N attention matrix → memory-bound to compute-bound.

Quantization — INT8/FP8 GEMMs with scaling factors.

// Row-wise softmax: one warp per row (rows of length ≤ a few thousand)
__global__ void softmaxRows(const float* in, float* out, int rows, int cols) {
    int row = blockIdx.x * (blockDim.x / 32) + threadIdx.x / 32;
    int lane = threadIdx.x & 31;
    if (row >= rows) return;
    const float* x = in + (size_t)row * cols;
    float m = -INFINITY;
    for (int c = lane; c < cols; c += 32) m = fmaxf(m, x[c]);
    for (int o = 16; o > 0; o >>= 1) m = fmaxf(m, __shfl_xor_sync(0xffffffff, m, o));
    float s = 0.f;
    for (int c = lane; c < cols; c += 32) s += __expf(x[c] - m);
    for (int o = 16; o > 0; o >>= 1) s += __shfl_xor_sync(0xffffffff, s, o);
    for (int c = lane; c < cols; c += 32) out[(size_t)row * cols + c] = __expf(x[c] - m) / s;
}

38.4 Scientific computing

Finite-difference PDE solvers (heat, wave, Navier–Stokes) with 2-D/3-D stencils.

Sparse iterative solvers (Conjugate Gradient with cuSPARSE SpMV + cuBLAS dot/axpy).

Molecular dynamics (GROMACS, AMBER, LAMMPS, NAMD use CUDA): neighbor lists, cell lists, PME with cuFFT.

Lattice Boltzmann, Monte Carlo (photon transport, option pricing), spectral methods (cuFFT).

38.5 Data analytics & databases

RAPIDS cuDF (GPU DataFrames), GPU-accelerated SQL (BlazingSQL successors, HeavyDB, Spark RAPIDS plugin), hash joins, group-by aggregation with hash tables in global memory (cuCollections), GPU-accelerated vector search (cuVS/CAGRA).

38.6 Cryptography & others

Hashing (SHA-256 brute force), password recovery (hashcat), zero-knowledge proof generation (MSM/NTT kernels), ray tracing (OptiX), audio/video processing, genomics (sequence alignment: Smith–Waterman), computational finance (Monte Carlo pricing, risk), weather/climate models (GPU ports), quantum simulation (cuQuantum).