Is there a way to write to a cudaMipmappedArray_t?

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;
}

Did you managed to achieve to write to cudaMipmappedArray ?