9

When I bind an array to a texture in CUDA,

  1. is that array copy to a texture space? or,
  2. is that array reference as a texture?

If the answer is 1., then i can bind a texture and safety fetch data from the texture memory space while I write the result to the array, which is allocated in global memory.

If the answer is 2., then, is the texture memory a global memory space where the data is cached and spatially fetched?

I'd like to know about this topic, as I've seen some question related to this topic and I've not the answer clear right now.

Thanks in advance.

pQB
  • 3,077
  • 3
  • 23
  • 49

1 Answers1

16

The answer is the second option, but from there things get a little more complex. There is no such thing as "texture memory", just global memory which is accessed via dedicated hardware which include an on GPU read cache (6-8kb per MP depending on card, see table F-2 in Appending F of Cuda Programming Guide) and a number of hardware accelerated filtering/interpolation actions. There are two ways the texture hardware can be used in CUDA:

  1. Bind linear memory to a texture and read from it in a kernel using the 1D fetch API. In this case the texture hardware is really just acting as a read-through cache, and (IIRC) there are no filtering actions available.
  2. Create a CUDA array, copy the contents of linear memory to that array, and bind it to a texture. The resulting CUDA array contains a spatially ordered version of the linear source, stored in global memory in some sort of (undocumented) space filling curve. The texture hardware provides cached access to that array, include simultaneous memory reads with hardware accelerated filtering.

You might find the overview of the GT200 architecture written by David Kanter worth reading to get a better idea of how the actual architecture implements the memory hierarchy the APIs expose.

maxywb
  • 2,275
  • 1
  • 19
  • 25
talonmies
  • 70,661
  • 34
  • 192
  • 269
  • Amazing answer :) Undocumented and rumored is what i want to know ;) So, if i write to the original array and fetch data from the binded texture, the behavior is undefined, right?. Thanks you. – pQB Jul 11 '11 at 10:41
  • 3
    In CUDA, textures and arrays are opaque, read only objects. No writing is allowed (there isn't even a guarantee that the internal texture layout is the same between different hardware). If you have a Fermi card and are using CUDA 3.2 or later, you can use surfaces. The [surfaces API](http://developer.download.nvidia.com/compute/cuda/3_1/toolkit/docs/online/group__CUDART__SURFACE.html) provide both read and write to "textures", although there is no filtering support at present, AFIAK. – talonmies Jul 11 '11 at 10:51
  • 3
    Within a kernel invocation, the texture cache does no maintain coherency with the underlying global memory storage, see section 3.2.10.4 of the Programming Guide. 2D textures using filtering do not have to be bound to an opaque cudaArray, they can also be bound to pitch-linear memory allocated with cudaMallocPitch(). – njuffa Jul 11 '11 at 19:46