texture with tex1Dfetch bind the texture, and try to get it out.

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

I just want to check how the texture is working and I tried the following code. My purpose is to bind an array of lines which contains 4 coordinates to a texture.

I can not get the some results. Maybe the texture is not bounded with the array, I am so sure. Can anyone help me? Thanks.

texture<float4, 1, cudaReadModeElementType> texLines;

__global__ void TestTexture(float4* test1)

{

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

	

	test1[idx] = tex1Dfetch(texLines, idx);

}

void MapGeometryToTexture(TLines Lines, int NumOfLines)

{

	float4* ScrCoords;

	cutilSafeCall( cudaMalloc((void **) &ScrCoords, NumOfLines * sizeof(float4)) );

	

	for (int i=0; i < NumOfLines; i++)

			ScrCoords[i] = make_float4(Lines[i].x1, Lines[i].y1, Lines[i].x2, Lines[i].y2);

	cutilSafeCall(cudaBindTexture(0, texLines, ScrCoords, NumOfLines * sizeof(float4)));

	

	float4* d_testing;

	cutilSafeCall(cudaMalloc((void **)&d_testing, 256 * sizeof(float4)) );

	

		TestTexture<<<1, 256>>>(d_testing);

	cudaThreadSynchronize();

	for (int i = 0; i < NumOfScreens; i++)

	{

		printf("\nOri: %f,%f,%f,%f", ScrCoords[i].x, ScrCoords[i].y, ScrCoords[i].z, ScrCoords[i].w);

		printf("\nTex: %f,%f,%f,%f", d_testing[i].x, d_testing[i].y, d_testing[i].z, d_testing[i].w);

	}

}
2

I see you’re allocating an array (ScrCoords) on the device first and afterwards you’re using the host to fill out the values for ScrCoords. This is not how CUDA works. The host (CPU) does not have direct access to device memory, so you need to create an array on the host side first, then fill out the values using host code and finally upload the host array to the device array(ScrCoords) using cudaMemCpy.

N.

3

Thank you!

So the ScrCoords = make_float4() is host code. But when I debug it, I saw the values were stored in the ScrCoords. Is it still in the host memory? And actually is there nothing in the device memory after that?

4

what is this function?
void TestTexture(float4* test1)
it is not a kernel function.

global void TestTexture(float4* test1)
you forgot to put “global” before this function.

5

Oh, sorry, I forget to type in this post, in my code it is there. thanks for pointing it out

6

Nico correctly,
you can not directly fill data to device memory.
for (int i=0; i < NumOfLines; i++)
ScrCoords[i] = make_float4(Lines[i].x1, Lines[i].y1, Lines[i].x2, Lines[i].y2);
you can not use printf() with device memory
for (int i = 0; i < NumOfScreens; i++)
{
printf(“\nOri: %f,%f,%f,%f”, ScrCoords[i].x, ScrCoords[i].y, ScrCoords[i].z, ScrCoords[i].w);
printf(“\nTex: %f,%f,%f,%f”, d_testing[i].x, d_testing[i].y, d_testing[i].z, d_testing[i].w);
}
you has to create data from host-> copy to device memory and bind to texture (in cuda2.2 you can bind host memory to texture directly, read guide for more detail).
if you want to see the result, you need to copy back data from device to host memory.