Troubles creating a large 3D texture with a GTX 480

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi, I just upgraded from a GTX 260 to a GTX 480 in hopes that I could create larger 3D texture references. In the programming guide (version 3.1 beta) it states in appendix G that devices of compute capability 2.0 can have a 4096x 4096x 4096 3D texture. This is up from 2048 in each direction in the 1.3 devices.

Anyway, I can’t seem to be able to create any larger of a texture. It seems to choke on the allocation of the 3D cudaArray before I can even bind it to the texture. The following code works fine if mx < 2048, but I get an ‘invalid argument’ if mx is any larger.

Is there a different way I am supposed to be doing this? Any help would be appreciated, thanks!

-fna

[codebox]cudaArray *testArray = 0;

char *testPtr;

texture<char, 3, cudaReadModeElementType> testTexture;

void runTest( )

{

cudaError_t e1;

int mx = 1024*3;

int my = 128;

int mz = 128;

int numPixels = mx*my*mz;

cudaMalloc((void **)&testPtr, numPixels * sizeof(char)) ;

cudaMemset(testPtr,0,numPixels * sizeof(char)) ;

e1 = cudaGetLastError();

if(e1 != cudaSuccess)

{

    printf("\nError in cudaMalloc:  ");

    printf(cudaGetErrorString(e1));printf("\n");

    return;

}

cudaExtent volExtent = make_cudaExtent(mx, my, mz); 

cudaChannelFormatDesc charTex = cudaCreateChannelDesc<char>();      

cudaMalloc3DArray(&testArray, &charTex, volExtent);

e1 = cudaGetLastError();

if(e1 != cudaSuccess)

{

    printf("\nError in cudaMalloc3DArray:  ");

    printf(cudaGetErrorString(e1));printf("\n");

    return;

}

cudaMemcpy3DParms copyParams = {0};

copyParams.srcPtr   = make_cudaPitchedPtr((void*) (testPtr), volExtent.width*sizeof(char), volExtent.width, volExtent.height);

copyParams.dstArray = testArray;

copyParams.extent   = volExtent;

copyParams.kind     = cudaMemcpyDeviceToDevice;

cudaMemcpy3D(&copyParams) ;

cudaBindTextureToArray(testTexture, testArray) ;

}[/codebox]

I’m running windows XPx64, CUDA toolkit 3.1beta, driver 257.15 with a GTX 480.

2

So it states in the reference manual that a cudaArray can’t be larger than 2048 in any of the three directions.
So I guess my question is how do you get a 3D texture that is larger than 2048 if an array cannot be?
Is there something else I can bind to the texture? cudaBindTexture and cudaBindTexture2D seem to be only for 1 and 2D textures respectively…

3

Do you compile with “arch=sm_20” ?

4

Yes, I do.

From the reference manual (3.1 beta):

“• A3Darray is allocated if all three extents are non-zero. For 3D arrays valid extent ranges are {(1,2048),(1,2048),(1,2048)}.”

The only way I know how to create a 3D texture is using a cudaArray. But if textures can be up to 4096 with a GTX 480 whilst cudaArrays can only be up to 2048 there must be some other way?

5

Yes, although Fermi supports 4k 3D textures, apparently there is a software limitation in the driver that limits Cuda arrays to 2048 size currently. We’ll fix this in a future release.

6

Thanks for the reply Simon. Waiting for 3.2 beta ;)