Mipmaps come to CUDA straight from computer graphics, where Lance Williams introduced them in his 1983 SIGGRAPH paper Pyramidal Parametrics.7 A mipmap is a precomputed pyramid of successively halved images: the base level holds the texture at full resolution, the next level holds it at half the width and height (one-quarter the texels), and so on down to a 1×1 image. The name abbreviates the Latin multum in parvo—“much in a small space”—which is what the pyramid delivers: every scale of the image, stored together for only one-third more memory than the base level alone.
The pyramid exists to combat aliasing under minification, when a texture is drawn so small that a single output pixel covers many texels. Sampling only the base level in that situation misses most of the texels that ought to contribute, producing the shimmering and moiré patterns characteristic of undersampling. A mipmapped fetch instead selects the level whose texels most nearly match the size of the region being sampled, so that one filtered read already reflects the whole region. The level it selects—generally a fractional one—is the level of detail, or LOD; blending the two nearest levels, on top of bilinear filtering within each level, is what is meant by trilinear filtering.
A graphics pipeline computes the LOD automatically from the
screen-space derivatives of the texture coordinates. A CUDA kernel has
no such derivatives, so it selects the level itself:
tex2DLod() takes the LOD as an explicit parameter, while
tex2DGrad() takes the coordinate gradients and lets the
hardware derive the LOD as the graphics pipeline would.
A mipmapped array is allocated with cudaMallocMipmappedArray(), which
takes the channel format, the base-level extent, and the number of
levels. Each level is itself a CUDA array: cudaGetMipmappedArrayLevel()
hands back the array for a given level so its texels can be filled with
a memcpy. Applications either downsample the levels on the host—a box
filter that averages each 2×2 block is the usual starting point—or
generate them on the device. The texture object is created over a
resource descriptor of type cudaResourceTypeMipmappedArray;
its filterMode selects filtering within a level and its
mipmapFilterMode selects filtering between levels, so
setting both to cudaFilterModeLinear requests trilinear
sampling, and
minMipmapLevelClamp/maxMipmapLevelClamp bound
the range of levels the sampler may reach. A mipmapped texture always
uses normalized coordinates.
cudaMipmappedArray_t mipArray;
cudaExtent extent = make_cudaExtent( Width, Height, 0 );
cuda(MallocMipmappedArray( &mipArray, &channelDesc, extent, numLevels, 0 ));
// Each level is a CUDA array in its own right; fill it with a memcpy.
for ( unsigned int level = 0; level < numLevels; level++ ) {
cudaArray_t levelArray;
cuda(GetMipmappedArrayLevel( &levelArray, mipArray, level ));
cuda(Memcpy2DToArray( levelArray, 0, 0, /* this level's texels */ ));
}
cudaResourceDesc resDesc = { .resType = cudaResourceTypeMipmappedArray };
cudaTextureDesc texDesc = {};
resDesc.res.mipmap.mipmap = mipArray;
texDesc.normalizedCoords = 1;
texDesc.filterMode = cudaFilterModeLinear; // within a level
texDesc.mipmapFilterMode = cudaFilterModeLinear; // between levels
texDesc.maxMipmapLevelClamp = (float) (numLevels - 1);
cuda(CreateTextureObject( &tex, &resDesc, &texDesc, NULL ));The microdemo tex2d_mipmap.cu exercises this path end to
end. It fills each level of a small texture with a constant equal to
that level’s index, then samples the center of the image with
tex2DLod() across a range of LODs. Because every level is a
constant, linear mipmap filtering makes the sampled value equal to the
LOD—a request for LOD 1.5 returns the even blend of levels 1 and 2—so
both the level selection and the inter-level blend are visible in the
output:
LOD sampled
0.00 0.000
0.50 0.500
1.00 1.000
1.50 1.500
2.00 2.000
2.50 2.500
3.00 3.000That demo fills its levels with constants; real mipmap levels come
from downsampling the base image, and here CUDA differs from the
graphics APIs: it ships no mipmap generator—there is no equivalent of
OpenGL’s glGenerateMipmap. The pyramid must be built by the
application, on the host or on the device. One convenient, high-quality
option is NPP’s image-resize primitive. The microdemo
mipmapGenNPP.cu wraps it in a
GenerateMipmapsNPP() helper, written to copy verbatim, that
takes a base image in pitched device memory and, for each level in turn,
area-resamples the base down to that level’s dimensions with
nppiResize_32f_C1R() and writes the result straight into
the mipmapped array. Because the source is always the base image, it
never aliases the destination, so a single scratch buffer—allocated
once, ahead of the loop—serves every level. The resampling filter is a
parameter, because the right one is application-dependent:
NPPI_INTER_SUPER area-averages—the box filter for a 2:1
reduction—and preserves the image mean at any ratio, so the 1×1 apex
equals the average of the base image, a property the demo checks—while
NPPI_INTER_CUBIC and NPPI_INTER_LANCZOS trade
ringing for sharpness. Generating the levels by hand is also
straightforward: bind the finer level as a texture and the coarser level
as a surface (allocating the array with
cudaArraySurfaceLoadStore), then have a kernel read each
2×2 footprint through the texture unit—one bilinear fetch at the block
center is the box average—and surf2Dwrite() the
result.
The size limits for mipmapped textures are reported by
cudaGetDeviceProperties() in the
cudaDeviceProp.maxTexture2DMipmap field, along with the
corresponding 1D and layered fields.