GPU Texture and Constant Memory are specialized GPU memory spaces with dedicated caches — optimized for specific access patterns that offer higher effective bandwidth than global memory for appropriate workloads.

CUDA Memory Hierarchy Summary

Texture Memory

• Cached in a separate L1 texture cache (separate from L1 data cache).
• Spatial locality caching: Optimized for 2D access patterns — if a thread accesses (x,y), neighbors (x+1,y), (x,y+1) likely cached.
• Hardware interpolation: GPU hardware performs bilinear/trilinear interpolation for free.
• Address modes: Wrap, clamp, mirror — hardware boundary handling.
• Usage: Image processing (sampling at non-integer coordinates), simulation stencils.

Texture Example

texture<float, 2, cudaReadModeElementType> tex;
cudaBindTexture2D(0, tex, d_data, channelDesc, width, height, pitch);

__global__ void sample_kernel() {
    float val = tex2D(tex, u, v);  // Bilinear interpolation included
}

Constant Memory

• 64KB total, cached in dedicated constant cache per SM.
• Broadcast: If all threads in a warp access the same address → single cache transaction (vs. 32 separate loads from global memory).
• Best use: Read-only data accessed uniformly by all threads (filter coefficients, LUT, camera parameters).
• Performance: Matching all-uniform access → as fast as registers. Divergent access → serialized (slow).

Read-Only Cache (__ldg)

• Modern alternative to texture for read-only global data.
• __ldg(&ptr[i]): Use L1 read-only cache (separate from L1 data cache).
• No setup required — simpler than texture objects.
• Good for gather patterns with spatial locality.

Texture and constant memory are specialized caches that provide free speedups for specific access patterns — image processing kernels using texture memory can achieve 2-4x better cache hit rates than equivalent global memory accesses on spatially correlated data.