help needed: texture read and bind problem


I’m tring to read value from a 3d cudaArray which is bound to a texture reference.

but it seems tex3d doesn’t work at all.

I dont know the problem is whether i use tex3d incorrectly or the cudaArray is bound to texref wrongly…

all in one cu file

texture<float4, 3, cudaReadModeElementType> stateTex; // 3D texture

// at first i initialize the array, pitch, texref and bind the array to the texref.

extern "C" void cudoInitTexture(cudaPitchedPtr *pitch, cudaArray **array, float3 arraySize, GLuint *texObj, int glStage)


        *array = 0;

        // create the texture array

        cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();

        cudaExtent volumeSize = make_cudaExtent(arraySize.x, arraySize.y, arraySize.z);

        CUDA_SAFE_CALL(cudaMalloc3DArray(array, &channelDesc, volumeSize) );

// create memory using a cudaPitchedPtr

        cudaExtent pitchVolSize = make_cudaExtent(arraySize.x*sizeof(float4), arraySize.y, arraySize.z);

        CUDA_SAFE_CALL(cudaMalloc3D(pitch, pitchVolSize));

glGenTextures(1, texObj);

        glActiveTexture(GL_TEXTURE0 + glStage);

        glBindTexture(GL_TEXTURE_2D_ARRAY, *texObj);



        glTexImage3D(GL_TEXTURE_2D_ARRAY, 0, GL_RGBA32F, arraySize.x, arraySize.y, arraySize.z, 0, GL_RGBA, GL_FLOAT, 0);

        glBindImageTextureEXT(glStage, *texObj, 0, true, 0,  GL_READ_WRITE, GL_RGBA32F);

        glBindTexture(GL_TEXTURE_2D_ARRAY, 0);

stateTex.normalized = false;                   

        stateTex.filterMode = cudaFilterModePoint;    

        stateTex.addressMode[0] = cudaAddressModeClamp; 

        stateTex.addressMode[1] = cudaAddressModeClamp;

        stateTex.addressMode[2] = cudaAddressModeClamp;

        CUDA_SAFE_CALL(cudaBindTextureToArray(stateTex, *array, channelDesc));

        cutilCheckMsg( "cudaBindTextureToArray" );


// then I do the initialize for position of each particles of the object, and it works.

// the initial shape is rendered as expectable by openGL, which means array, pitch and gl interoperability work.

// then i use tex3d to read from the array via texref

__global__ void simulateExplicit(cudaPitchedPtr cpp, float2 clothSize, float clothWidthLength, float3 windForce, float3 displacement, float deltaTime)


        int x = threadIdx.x + blockIdx.x * blockDim.x;

        int y = threadIdx.y + blockIdx.y * blockDim.y;

        if(x >= clothSize.x || y >= clothSize.y)


int index = x + y * clothSize.x;

// in order to find out the reason, i commented all other simulation operations.

        // just read the value and store it again without any computation, so the object should remain the same shape,

        // but it just disappear immediately in like very very short time

        float4 parPosTmp = tex3D(stateTex, x, y, 0);

        ((float4 *)cpp.ptr)[index] = parPosTmp;


The object’s shape should remain the same, but the whole object disappear, hence i guess the problem may be because I use tex3d or bind function incorrectly.

sadly, i have no idea why…

any can one help?

thanks alot!