How to bind array of doubles to texture2D?

Question

I try to use texture memory/binding instead of global memory, but I cannot pass by binding texture. First thing I learned is CUDA does not support double for texture, so casting is needed, ok.

I declared global texture variable:

texture<int2, 2> texData;

then right after allocating device memory (cudaMalloc) with size (in bytes) width*height * sizeof(double) I try to bind it:

cudaChannelFormatDesc desc = cudaCreateChannelDesc<int2>();
cudaStatus = cudaBindTexture2D(nullptr, &texData, dev_data, &desc,  width, height, 0);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "Binding texture failed: %s\n", cudaGetErrorString(cudaStatus));
    goto Error;
}

And this binding fails with error "invalid argument". Width and height are 2048 which is much below texture 2d limits: 65536 x 65536 x 1048544.

So what did I do wrong here?

Sidenote: the signature of cudaBindTexture2D:

extern __host__ cudaError_t CUDARTAPI cudaBindTexture2D(size_t *offset, 
    const struct textureReference *texref, const void *devPtr, 
    const struct cudaChannelFormatDesc *desc, size_t width, size_t height, size_t pitch);

Answer 1

You should make a proper pitched allocation

size_t pitch;
cudaMallocPitch((void**)&dev_data, &pitch, width* sizeof(double),height);

cudaChannelFormatDesc desc = cudaCreateChannelDesc<int2>();

cudaStatus = cudaBindTexture2D(nullptr, texData, dev_data, desc,  width, height, pitch);

Note that while CUDA errors are generally not very informative, the one you got ""invalid argument"" is quite informative. The arguments you put in your functions are invalid.