Early CUDA hardware gave texture a second role beyond graphics: it
was a way to read global memory through a separate cache and a separate
load path. Binding device memory to a texture and reading it with
tex1Dfetch() sidestepped the strict coalescing constraints that applied
to ordinary global loads, and the texture cache could reduce external
bandwidth on data with locality that the L1 and L2 caches did not
capture well.
On hardware from Maxwell onward, this motivation has gone away. The
texture cache was unified with the L1 cache, so reading through texture
no longer provides a cache or a load path that an ordinary memory access
does not already have. The modern way to get the same benefit – reading
through the read-only data cache – is to mark a pointer
const __restrict__, or to read through it explicitly with
the __ldg() intrinsic. The compiler can service both
without the setup cost of a texture, and without the addressing limits
that texture imposes – in particular the 27-bit element index that caps
how much memory a single texture read via tex1Dfetch() can cover.
Texturing remains worthwhile for its fixed-function features –
interpolation, the addressing modes, and integer-to-float promotion –
which the rest of this chapter covers. Two read-path microdemos are kept
in the source for historical interest: tex1dfetch_big.cu chains several
textures together to cover more than one texture’s worth of memory, and
tex1dfetch_int2float.cu exercises the integer-to-float promotion (Table
10-1) while texturing from mapped host memory.