The CUDA Handbook ▸ Chapter 10. Texturing

10.7 1D Surface Read/Write

Until the second generation of CUDA hardware, CUDA kernels could access the contents of CUDA arrays only via texturing. Other access to CUDA arrays, including all write access, could be performed only via memcpy functions such as cudaMemcpyToArray(). The only way for CUDA kernels to both read from and write to a given region of memory was to texture from linear device memory.

But with the surface read/write functions newly available in SM 2.x, developers can create surface objects over CUDA arrays and use the surf1Dread() and surf1Dwrite() intrinsics to read and write the CUDA arrays from a kernel. Unlike texture reads, which have dedicated cache hardware, these reads and writes go through the same L2 cache as global loads and stores.

The 1D surface read/write intrinsics are declared as follows:

template<class T> T surf1Dread(cudaSurfaceObject_t surfObj, int x, boundaryMode = cudaBoundaryModeTrap);
template<class T> void surf1Dwrite(T data, cudaSurfaceObject_t surfObj, int x, boundaryMode = cudaBoundaryModeTrap);

These intrinsics are not type-strong – the surface object carries no element type – and the size of the memory transaction depends on sizeof(T) for a given invocation of surf1Dread() or surf1Dwrite(). The x offset is in bytes and must be naturally aligned with respect to sizeof(T): for 4-byte operands such as int or float, offset must be evenly divisible by 4, for short it must be divisible by 2, and so on.

Support for surface read/write is far less rich than texturing functionality⁴. Only unformatted reads and writes are supported, with no conversion or interpolation functions; and the border handling is restricted to only two modes.

Boundary conditions are handled differently for surface read/write than for texture reads. For textures, this behavior is controlled by the addressing mode in the texture object’s descriptor. For surface read/write, the method of handling out-of-range offset values is specified as a parameter of surf1Dread() or surf1Dwrite(). Out-of-range indices can either cause a hardware exception (cudaBoundaryModeTrap), or read as 0 for surf1Dread() and are ignored for surf1Dwrite() (cudaBoundaryModeZero).

Because the surface object carries no element type, it is easy to write a templated 1D memset routine that works for all types:

template <typename T>
__global__ void
surf1Dmemset_kernel( cudaSurfaceObject_t surf1D, T value, int offset, size_t N )
{
    for ( size_t i = blockIdx.x*blockDim.x + threadIdx.x;
                 i < N;
                 i += blockDim.x*gridDim.x )
    {
        surf1Dwrite( value, surf1D, (offset+i)*sizeof(T) );
    }
}

This kernel is in the microdemo surf1Dmemset.cu, which creates a 1D CUDA array for illustrative purposes, initializes it with the above kernel, and prints the array in floating point and integer forms.

A generic template host function wraps this kernel: it builds a surface object over the CUDA array with cudaCreateSurfaceObject(), launches the kernel, and destroys the object.

template<typename T>

cudaError_t

surf1Dmemset( cudaArray *array, int offset, T value, size_t N )
{
    cudaError_t status;
    cudaSurfaceObject_t surfObj = 0;
    cudaResourceDesc resDesc = { .resType = cudaResourceTypeArray };
    resDesc.res.array.array = array;
    cuda(CreateSurfaceObject( &surfObj, &resDesc ));
    surf1Dmemset_kernel<<<2,384>>>( surfObj, value, offset, N );

Error:

    cudaDestroySurfaceObject( surfObj );
    return status;
}

Both surface objects and texture objects are ordinary handles that can be passed to a kernel as parameters, so a generic function like this one is straightforward to write for either. (With the file-scope texture and surface references used by older versions of CUDA, this was easy for surfaces – whose references were untyped – but awkward for textures, whose references were both type-strong and global.)

A one-line change from:

CUDART_CHECK(surf1Dmemset(array, 0, 3.141592654f, NUM_VALUES));

to:

CUDART_CHECK(surf1Dmemset(array, 0, (short) 0xbeef, 2*NUM_VALUES));

will change the output of this program from:

0x40490fdb 0x40490fdb ... (16 times)

3.141593E+00 3.141593E+00 ... (16 times)

to:

0xbeefbeef 0xbeefbeef ... (16 times)

-4.68253E-01 -4.68253E-01 ... (16 times)