The CUDA Handbook ▸ Chapter 10. Texturing

10.13 Cubemap Textures

Cubemaps are another idea CUDA inherits from computer graphics, where they are the standard way to texture a direction rather than a surface: environment reflections, sky boxes, and omnidirectional lighting all read a cubemap.⁸ A cubemap is a set of six square images, one per face of an axis-aligned cube, addressed not by 2D coordinates but by a 3D direction vector radiating from the cube’s center. The largest-magnitude component of the direction selects the face; dividing the other two by it gives the coordinates on that face. CUDA stores the six faces in the standard graphics order—+X, −X, +Y, −Y, +Z, −Z—and texCubemap() applies exactly this addressing, so a cubemap sampled in a CUDA kernel matches one sampled in OpenGL or Direct3D.

A cubemap lives in a CUDA array allocated with cudaMalloc3DArray(), passing the cudaArrayCubemap flag and a depth extent of six—the faces occupy the array’s third dimension, so all six load with a single cudaMemcpy3D(). The faces must be square. Its texture object is built over a resource descriptor of type cudaResourceTypeArray, and because the fetch is by direction, a cubemap always uses normalized coordinates.

cudaArray_t cubeArray;
cudaExtent extent = make_cudaExtent( faceDim, faceDim, 6 );  // depth 6 = six faces
cuda(Malloc3DArray( &cubeArray, &channelDesc, extent, cudaArrayCubemap ));

// The six faces are consecutive in host memory, so they form the copy's depth.
cudaMemcpy3DParms p = {};
p.srcPtr   = make_cudaPitchedPtr( host, faceDim*sizeof(float), faceDim, faceDim );
p.dstArray = cubeArray;
p.extent   = extent;
p.kind     = cudaMemcpyHostToDevice;
cuda(Memcpy3D( &p ));

cudaResourceDesc resDesc = { .resType = cudaResourceTypeArray };
cudaTextureDesc  texDesc = {};
resDesc.res.array.array  = cubeArray;
texDesc.normalizedCoords = 1;
cuda(CreateTextureObject( &tex, &resDesc, &texDesc, NULL ));

The microdemo texcubemap.cu makes the face selection observable. It fills each face with a constant equal to that face’s index, then samples the cubemap along the six axis directions. Each axis points straight at one face center, so the sampled value comes back as that face’s index—confirming the direction-to-face mapping end to end:

  direction  -> face    sample
  +X         0          0.0
  -X         1          1.0
  +Y         2          2.0
  -Y         3          3.0
  +Z         4          4.0
  -Z         5          5.0

Cubemaps can themselves be arrayed: allocating with cudaArrayCubemap | cudaArrayLayered and a depth of six times the count creates a stack of cubemaps that texCubemapLayered() selects by integer index, just as layered textures (Section 10.11) do for the 2D case. The cubemap size limit is reported by cudaGetDeviceProperties() in cudaDeviceProp.maxTextureCubemap, with the layered form in maxTextureCubemapLayered.