Is there a way to write to a cudaMipmappedArray_t, similar to how you can use a surface to write to a cudaArray?
There are API calls to allocate cudaMipmappedArrays and bind textures to cudaMipmappedArrays, but there doesn’t seem to be an equivalent call to bind a surface to a mipmap. Attempting to call cudaCreateSurfaceObject() with cudaResDesc that points to a mipmapped array, for example, will throw "“invalid argument”
Do cudaMipmappedArray_t’s exist purely for OpenGL interop?
Attached is code to bind a texture object to a mipmap which also throws an error when you attempt to bind a surface to a mipmap. This is only CUDA-6.5.
#include <stdio.h>
#include <iostream>
#define checkCudaError(err) if(err != cudaSuccess){std::cout << "Cuda Error at " << __FILE__ << " , "<< __LINE__ << " calling (" << #err << ") of type " << cudaGetErrorString(err) << ". Code " << (int)err << std::endl; cudaThreadExit(); exit(0); }
int main(){
int NX = 1024;
int NY = 1024;
int LEVELS = 3;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned short>();
cudaExtent extent = make_cudaExtent(NX,NY,0);
cudaMipmappedArray_t dma_mipmap;
unsigned int mipmapFlags = cudaArraySurfaceLoadStore;
checkCudaError( cudaMallocMipmappedArray(&dma_mipmap, &channelDesc, extent, LEVELS, mipmapFlags) );
cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(cudaResourceDesc) );
resDesc.resType = cudaResourceTypeMipmappedArray;
resDesc.res.mipmap.mipmap = dma_mipmap;
cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(cudaTextureDesc) );
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
texDesc.addressMode[2] = cudaAddressModeClamp;
texDesc.filterMode = cudaFilterModeLinear;
texDesc.readMode = cudaReadModeNormalizedFloat;
texDesc.sRGB = 0;
texDesc.normalizedCoords = 0;
texDesc.maxAnisotropy = 1;
texDesc.mipmapFilterMode = cudaFilterModeLinear;
texDesc.mipmapLevelBias = 0;
texDesc.minMipmapLevelClamp = 0;
texDesc.maxMipmapLevelClamp = LEVELS;
cudaResourceViewDesc viewDesc;
viewDesc.format = cudaResViewFormatUnsignedShort1;
viewDesc.width = NX;
viewDesc.height = NY;
viewDesc.depth = 0;
viewDesc.firstMipmapLevel = 0;
viewDesc.lastMipmapLevel = LEVELS;
viewDesc.firstLayer = 0; // for non-layered resources this value has to be zero
viewDesc.lastLayer = 0;
cudaTextureObject_t to_mipmap;
checkCudaError( cudaCreateTextureObject( &to_mipmap, &resDesc, &texDesc, &viewDesc) );
// this doesn't work, because you can only bind a surface to cudaArray, not to a cudaMipmappedArray
// cudaCreateSurfaceObject will throw "invalid argument' error
cudaSurfaceObject_t so_mipmap;
checkCudaError( cudaCreateSurfaceObject( &so_mipmap, &resDesc) );
checkCudaError( cudaFreeMipmappedArray(dma_mipmap) );
return 0;
}