gpu texture memory
**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**
| Memory | Scope | Cached? | Bandwidth |
|--------|-------|---------|----------|
| Global | All threads | L2 | ~900 GB/s |
| Shared | Block | On-chip | ~19 TB/s |
| Texture | All threads | L1 tex | ~900 GB/s + spatial cache |
| Constant | All threads | Const cache | Broadcast |
| Local | Thread | L2 | Slow |
**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**
```cuda
texture 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.