Simple texture problem Code will not compile.

Hi, I’m trying to allocate a pitched linear memory, and then binding to a texture.

Its a 2-dimensional array of floats (size NDETS, NDETSLICES);

The following code will not compile.

Any ideas?

float* d_proj;

size_t pitch, offset;

cudaMallocPitch((void**) &d_proj, &pitch, sizeof(float)*NDETS,NDETSLICES);

texture<float, 2, cudaReadModeElementType> texRef;

cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();

cudaBindTexture2D (&offset, texRef, d_proj, desc,  NDETS, NDETSLICES, pitch);

The compiler error i’m seeing is:
error C2018: unknown character ‘0x40’
corresponding to the line with the texture reference declaration.
that error makes absolutely no sense since ‘0x40’ is ASCII for ‘@’ and I certainly do not have that character in there.

Why is using textures so hard???

What encoding did you use in your text editor to create the program code? Is it unicode (utf-8) by any chance?

No clue…I’m using MS Visual Studio 2005.

Anyway, I was doing something illegal: declaring a texture reference inside a function.

When I moved it to file scope, the compiler errors went away.

Now that I know that texture declarations must be at file scope, I have a completely new problem:

If my host and device code are in separate files, how the heck do I make the texture visible to both files???

The simple answer is that you don’t. CUDA has no linker, so everything must be available in the same compilation unit. If you need to access the texture outside of file scope (which is only really going to happen in host code), then write an access function wrapper for it and call that from your other code (or pass it around as a function pointer).

The simpleTexture example in the SDK seems to work just fine without an “access function wrapper.”

At any rate, I will try to implement your suggestion if only I knew what you meant. I’m sure your suggestion can be explained with less than 5 lines of example code.

A picture is worth a thousand words, no?

Thanks again.

OK then, cut and pasted straight from some production code of mine:

texture<float4, 3, cudaReadModeElementType> microstructure;

static cudaChannelFormatDesc MSChannelDesc;

bool mybind2tex( cudaArray *MSArray_d )

{

	cudaError_t CUDAstatus;

	bool status = true;

	if ((CUDAstatus = cudaBindTextureToArray(microstructure, MSArray_d, MSChannelDesc)) != cudaSuccess) {

		fprintf(stderr,"Cuda Error : %s %s %d\n", cudaGetErrorString(CUDAstatus), __FILE__, __LINE__);

		status = false;

	}

	return status;

}

Put the texture, and the kernels that use it compiled in one file, with a wrapper function which takes a cudaArray from a caller defined wherever you like and binds it to the texture. You could add kernel calls inside the function as well if you wanted. As MrAnderson42 pointed out in another thread where you asked about this, binding to texture is a very low overhead operation that you won’t even notice. I am sure you can see how you might do it with something other than 3D textures or with linear memory instead of a cudaArray. If you have more levels of abstraction and more than a single texture, use function pointers to a generic function prototype and create wrapper for each texture definition you need for your kernels.

Thanks, I did as you suggested and there are no compile errors or runtime crashes. However my program output is incorrect. This is due to bad tex2D() read operations. Everytime my kernel calls tex2D() the returned value is 0 (or undefined.)

I am allocating a 1024 x 768 float array with cudaMallocPitch and then binding my texture to it. Curiously, the emuDebug and emuRelease configurations all work correctly. The GPU can’t seem to understand my texture. I couldn’t even begin to understand the reasons why the non-emulation builds do not work correctly.

This is how I declare and bind my texture. Pretty much verbatim as to how you did it.

texture<float,2,cudaReadModeElementType> texRef_proj;

bool texbind(float* d_proj, int pitch) {

	size_t offset;

	cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();

	// set texture parameters

	texRef_proj.addressMode[0] = cudaAddressModeClamp;

	texRef_proj.addressMode[1] = cudaAddressModeClamp;

	texRef_proj.filterMode = cudaFilterModeLinear;

	texRef_proj.normalized = false;	// access with normalized texture coordinates

	

	bool status = true;

	cudaError_t CUDAstatus;

	if ((CUDAstatus = cudaBindTexture2D (&offset, texRef_proj,(void*)d_proj, desc, NDETS, NDETSLICES, pitch)) != cudaSuccess) {

		fprintf(stderr,"Cuda Error : %s %s %d\n", cudaGetErrorString(CUDAstatus), __FILE__, __LINE__);

		status = false;

	}

	return status;

}

bool texunbind() {

	 bool status = true;

	cudaError_t CUDAstatus;

	if ((CUDAstatus = cudaUnbindTexture(texRef_proj)) != cudaSuccess) {

		fprintf(stderr,"Cuda Error : %s %s %d\n", cudaGetErrorString(CUDAstatus), __FILE__, __LINE__);

		status = false;

	}

	return status;

}

I used a texture bound to (1D) linear memory and it worked. I’m past the file scope issue now.
I wanted to use a texture bound to pitched linear memory so that I could use tex2D for access, but I gave up on that.
It must be some bug in how I’m allocating/binding my 2-dimensional data. At any rate the 1D texture is fine, I guess.

Is there a performance advantage in terms of cache hits if I use a 2D texture vs. a 1D texture??