Constant, Texture & Surface Memory

Cached and broadcast efficiently when all threads of a warp read the same address; serialized when addresses differ.

16.1 Constant memory

Cached and broadcast efficiently when all threads of a warp read the same address; serialized when addresses differ. Ideal for filter weights, parameters, small lookup tables. Kernel parameters themselves live in constant memory (up to 32 KB of params since CUDA 12.1 on Volta+).

__constant__ float coeffs[16];                         // 64 KB total per module
cudaMemcpyToSymbol(coeffs, h_coeffs, sizeof(h_coeffs));

16.2 Texture objects

Benefits: 2-D spatial locality caching, free interpolation, boundary modes, normalized coordinates, format conversion. Surface objects (surf2Dwrite) allow writing to CUDA arrays.

cudaResourceDesc res = {};
res.resType = cudaResourceTypeArray;
cudaArray_t arr;
cudaChannelFormatDesc ch = cudaCreateChannelDesc<float>();
cudaMallocArray(&arr, &ch, w, h);
cudaMemcpy2DToArray(arr, 0, 0, h_img, w * sizeof(float), w * sizeof(float), h, cudaMemcpyHostToDevice);
res.res.array.array = arr;

cudaTextureDesc td = {};
td.addressMode[0] = td.addressMode[1] = cudaAddressModeClamp;   // border handling
td.filterMode = cudaFilterModeLinear;                            // hardware bilinear interpolation
td.readMode = cudaReadModeElementType;
td.normalizedCoords = 0;

cudaTextureObject_t tex;
cudaCreateTextureObject(&tex, &res, &td, nullptr);

__global__ void sample(cudaTextureObject_t tex, float* out, int w, int h, float scale) {
    int x = blockIdx.x * blockDim.x + threadIdx.x, y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < w && y < h) out[y * w + x] = tex2D<float>(tex, x * scale + 0.5f, y * scale + 0.5f);
}
cudaDestroyTextureObject(tex); cudaFreeArray(arr);