Hi,
in my application I have a large linear segment in device memory interpreted as multiple equally-sized 2D-layers. For fast sparse access I want to bind each layer to a texture.
This works perfectly fine if the dimensions for these layers are somewhat “common”, like 512x512 or 800x640. But for more arbitrary sizes like 1000x700 or 800x641 cudaBindTexture2D() exits with the error “invalid argument” when attempting to bind the second texture. The first texture binding seems to complete successfully (however, I did not check this by actually accessing this first texture).
Here is the source code for a minimal program that reproduces the error:
#include <stdio.h>
#include <cuda.h>
#include <cutil_inline.h>
texture<float, 2> tex0, tex1;
int main(int argc, char** argv)
{
const size_t W = 1000;
const size_t H = 700;
const size_t SIZE = W * H;
const size_t MEMSIZE = W * H * sizeof(float);
const size_t PITCH = W * sizeof(float);
const size_t NUM_TEXS = 2;
/* Allocate and initialize device data. */
float* d_data;
cutilSafeCall(cudaMalloc((void**)&d_data, NUM_TEXS * MEMSIZE));
cutilSafeCall(cudaMemset(d_data, 0, NUM_TEXS * MEMSIZE));
/* Bind data to the two textures. */
cudaChannelFormatDesc chD = cudaCreateChannelDesc<float>();
float* ptr = d_data;
printf("Binding 1st %ux%u data segment to 1st texture...", W, H);
cutilSafeCall(cudaBindTexture2D(NULL, &tex0, ptr, &chD, W, H, PITCH));
printf(" done.\n");
/* Increase data pointer to the start of the next layer. */
ptr += SIZE;
printf("Binding 2nd %ux%u data segment to 2nd texture...", W, H);
cutilSafeCall(cudaBindTexture2D(NULL, &tex1, ptr, &chD, W, H, PITCH)); /* This line throws an error. */
printf(" done.\n");
cutilExit(argc, argv);
}
In case W and H are set to 800 and 640 the program terminates successfully:
Binding 1st 800x640 data segment to 1st texture... done.
Binding 2nd 800x640 data segment to 2nd texture... done.
Press ENTER to exit...
For WxH = 1000x700 it does not:
Binding 1st 1000x700 data segment to 1st texture... done.
Binding 2nd 1000x700 data segment to 2nd texture...cudaSafeCall() Runtime API error in file <d:/coding/CUDA/my projects/TexBi
ndTest//main.cu>, line 32 : invalid argument.
AFAIK the returned error for the second call of cudaBindTexture2D() is not an error actually returned by the first call, since cudaBindTexture2D() isn’t asynchronous. In particular inserting even multiple lines of cutilSafeCall(cudaThreadSynchronize()) between the two bindings does not change the behavior.
Some dimensions that work: 800x640, 800x700, 800x642, 1000x640
…and some that don’t: 1000x700, 800x641, 1000x641, 1000x642
I guess CUDA has problems with textures having a particular number of texels, but I’m not able to figure out what exactly and I couldn’t find anything about restrictions either.
Is this a bug or am I missing something?