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

The zero pitch was the illegal argument. – talonmies Oct 12 '18 at 12:28 — talonmies, Oct 12 '18 at 12:28

Ander Biguri · Accepted Answer · 2018-10-12T12:37:52.827

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.

edited Oct 12 '18 at 12:37

answered Oct 12 '18 at 11:25

Ander Biguri

35,140
11
74
120

Many thanks, so the `cudaMalloc` was the culprit, after using `cudaMallocPitch` (and passing pitch) it is OK. I though the problem lies with dimensions or the description. Btw. I preserved the pointers passing because this is what header shows me -- this functions expects pointer to texture and description. – astrowalker Oct 12 '18 at 11:43
@astrowalker uhm.... I don't think so. It may not error, but does the code actually work? – Ander Biguri Oct 12 '18 at 12:20
1

I added the signature to my post, there are 2 more overloads, all of them use pointers. And the code works, and produces correct results. – astrowalker Oct 12 '18 at 12:37

How to bind array of doubles to texture2D?

1 Answers1