cuda opengl interoperation problem

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I’m trying to use cudaPitchedPtr to write data to a 3d cudaArray, then map the cudaArray to openGL as GL_TEXTURE_2D_ARRAY.

In my shader I use imageLoad function to load data.

But the program doesn’t draw anything…

texture<float4, 3, cudaReadModeElementType> texState;

cudaArray *texArrayState;

cudaPitchedPtr writeablePitch;

void cloth::initTexture()

{	

	texArrayState = 0;

	// create the texture array

	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();

	cudaExtent volumeSize = make_cudaExtent(widthSize, heightSize, numberOfTexStateLayers);

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

	

	CUDA_SAFE_CALL(cudaBindTextureToArray(texState, texArrayState, channelDesc));

	// create memory using a cudaPitchedPtr

	cudaExtent pitchedVolSize = make_cudaExtent(widthSize*sizeof(float4), widthSize, heightSize);

	CUDA_SAFE_CALL(cudaMalloc3D(&writeablePitch, pitchedVolSize));

	

	glGenTextures(1, &texObjState);

	glActiveTexture(GL_TEXTURE0 + stage);

	glBindTexture(GL_TEXTURE_2D_ARRAY, texObjState);

	glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_MIN_FILTER, GL_NEAREST);

	glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_MAG_FILTER, GL_NEAREST);

	glTexImage3D(GL_TEXTURE_2D_ARRAY, 0, GL_RGBA32F, widthSize, heightSize, numberOfTexStateLayers, 0, GL_RGBA, GL_FLOAT, 0);

	glBindImageTextureEXT(stage, texObjState, 0, true, 0,  GL_READ_WRITE, GL_RGBA32F);

	glBindTexture(GL_TEXTURE_2D_ARRAY, 0);

	

	CUDA_SAFE_CALL(cudaGraphicsGLRegisterImage(&resourceState, texObjState, GL_TEXTURE_2D_ARRAY, cudaGraphicsMapFlagsNone));

	CUDA_SAFE_CALL(cudaGraphicsMapResources(1, &resourceState, NULL));

	CUDA_SAFE_CALL(cudaGraphicsSubResourceGetMappedArray(&texArrayState, resourceState, 0, 0));

}

void cloth::initClothState()

{

	float2 clothSize = make_float2(widthSize, heightSize);

	float3 pos = make_float3(position.x, position.y, position.z);

	dim3 blocks(widthSize/blockSize, heightSize/blockSize);

	dim3 threads(blockSize, blockSize);

	initState<<<blocks, threads>>>(writeablePitch, clothSize, length, pos);  // call the kernel to initialize the position of cloth

	

	// set parameters for copyParams

	cudaMemcpy3DParms copyParams = {0};	

	copyParams.srcPtr = writeablePitch;

	copyParams.dstArray = texArrayState;

	copyParams.kind = cudaMemcpyDeviceToDevice;

	copyParams.extent = make_cudaExtent( widthSize, heightSize, numberOfTexStateLayers);

	

	CUDA_SAFE_CALL(cudaMemcpy3D(&copyParams));  // copy data from cudaPitchedPtr to cudaArray

}

__global__ void initState(cudaPitchedPtr cpp, float2 clothSize, float clothWidthLength, float3 centerPos)

{

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

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

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

		return;

	

	int offset =x + y * clothSize.x;

	

	float cellSize = clothWidthLength / (clothSize.x - 1);

	float boundaryWidth_length = clothWidthLength;

	float boundaryHeight_length = cellSize * (clothSize.y - 1);

	

	float3 start_corner = make_float3(centerPos.x - boundaryWidth_length/2.0, centerPos.y, centerPos.z + boundaryHeight_length/2.0);

	float3 particlePos = make_float3(start_corner.x + x * cellSize, start_corner.y, start_corner.z - y * cellSize);

	

	// initialize positions

	((float4 *)cpp.ptr)[offset] = make_float4(particlePos.x, particlePos.y, particlePos.z, .0f);

}

I don’t know what wrong it is.

anyone can help me out?

million thanks!

2

I’m not sure if 2D texture arrays are supported for OpenGL interop (they use a different layout than 3D textures). Does it work with GL_TEXTURE_3D?

3

i’ve checked the nvidia doc about OpenGL Interoperability.

the list includes GL_TEXTURE_2D_ARRAY.

4

I solved the problem by renaming main.cpp.cu to main.cpp.
In my main file, the cloth object is created by “new” (clothObj = new cloth(…)).
I found that the “new” doesn’t work correctly in cu file…
dont know why…really strange