CUDA 3.1beta: writes to texture?

I was reading the programming guide of the new CUDA 3.1 beta when I came across the following in section 3.2.8.1 (page 40):

It’s clearly stated that an OpenGL texture or renderbuffer can be read and written by kernels. Is it a mistake or is it really possible to write to texture now? I haven’t used CUDA for a while so it’s possible that I just missed it, but I always believed that writes to texture wasn’t possible with CUDA. That’d be great if it was the case now :)

Edit: I checked and it’s already stated in the programming guide 3.0. But I didn’t find any reference to writes to texture on the forums. That’s why I’m wondering if it has actually happened…

It says:

In CUDA, it appears as a CUDA array and be read and written by kernels or via cudaMemcpy2D() calls.

So you can read from an write to it because it is a CUDA array. You can also bind a texture to it (as you can bind textures to CUDA arrays)

Read the new programming guide more carefully. Specifically sections 3.2.5 and B.9. They detail a new surface type (looks like texture<>) that can be read with surf?Dread and written with surf?Dwrite. Surfaces are bound to cudaArrays, so this allows direct writing to a cudaArray from a kernel.

It is odd that the mention of this new type is not included in the bullet points for CUDA 3.1beta. I haven’t had time to see how well these new methods work, yet.

I had a quick can of the programming guide about surfaces. From my quick parse, surfaces look to have a subset of the functionality of textures, but with the important added ability of being able to write to them. I could find no mention of being able to define read mode types for surfaces. What I would like is the ability to write to a texture in the normalized float type, i.e., take a float between -1…1, and write this to a texture stored as a signed integer. Anyone know if the various texture reading modes are supported on surfaces, and as writing modes?

I assume those ?D’s are arrows, so these are probably C++ operators. Does this also work on 1.x hardware?

No, I think that’s standing in for 1D, 2D, and 3D. :)

Surfaces should be in the 3.1 beta.

Hum. It’s still a bit confusing for me.

I always believed that kernels could not write to CUDA arrays. In my understanding, using cudaMemcpy2D() was the only solution to update a CUDA array. And it’s indeed confirmed section 5.3.2.5 (about texture) of the programming guide (3.1):

In the same guide, the section 3.2.8.1 that I quoted in my original message states:

So section 3.2.8.1 states clearly that because the openGL texture (or renderbuffer) is registered in CUDA as a CUDA array, it’s therefore possible to write into it using kernels (and cudaMemcpy2D). So that would mean that kernels can write to CUDA arrays.

At last, as MisterAnderson mentionned, 3.1beta seems to have introduced a new type with surfaces. So writing to CUDA arrays is indeed possible via surf?Dwrite.

Question 1

Is is possible that a CUDA array is writable in some cases (a surface reference is bound to it or the CUDA array is an openGL buffer) but not writable if a texture reference is bound to it??? Sounds weird but… To me it seems that section 3.2.8.1 and 5.3.2.5 are contradictory.

Question 2

Can someone (probably from NVIDIA) explain what this new surface type is and what the differences are with a texture? Because section B.9 (that presents surface functions) refers to the maximum texture sizes as if texture features were relevant to surfaces as well. So they must have some relation with textures, while being writable…

Aaarggghhh, been looking too much at garbled URL’s lately ;)

Well, it is a beta version. A few random sentences that have yet to be updated should not deter you from trying out the new surface functions that are detailed in an entire new section of the guide.

Judging from 3.2.5, all one needs to do is allocate the cudaArray with the cudaArraySurfaceLoadStore flag to make read and written with the surface functions. It does not explicitly state that a texture reference cannot be bound to the same cudaArray. Presumably if you wanted to, you could bind a texture reference and a surface reference to the same array (with the normal cache coherency limitation). Someone will just have to write a code to try it and see if it works :)

That is my guess. It seems that surfaces are just exposing some of the previously “graphics only” hardware to CUDA that allows write-to-texture support.

Too bad it’s still only 2D max… We want 3D :)