Invalid texture reference

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I have this piece of code:

// Declarations

unsigned int m_set_size(1232);

float* m_p_host_set(new float[m_set_size]);

float* m_p_device_set;

texture<float, 1, cudaReadModeElementType> m_set_texture;

// Set texture parameters

m_set_texture.addressMode[0] = cudaAddressModeWrap;

m_set_texture.addressMode[1] = cudaAddressModeWrap;

m_set_texture.filterMode = cudaFilterModePoint;

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

// Allocate memory on GPU

cudaMalloc((void**) &m_p_device_set,	  m_set_size * sizeof(float));

CheckCudaError(__FILE__, __FUNCTION__, __LINE__);

// Copy data from RAM to GPU

cudaMemcpy(m_p_device_set, m_p_host_set, m_set_size * sizeof(float), cudaMemcpyHostToDevice);

CheckCudaError(__FILE__, __FUNCTION__, __LINE__);

// Bind CUDA texture

cudaChannelFormatDesc set_channel_descriptor = cudaCreateChannelDesc<float>();

cudaBindTexture(0, &m_set_texture, m_p_device_set, &set_channel_descriptor, m_set_size);

CheckCudaError(__FILE__, __FUNCTION__, __LINE__);

The call to cudaBindTexture(…) will always generate the error as follows:

“invalid texture reference”

Any idea why it does not work?

How to create a 1D texture in CUDA?

Cheers,

2

Your texture reference is invalid because you are using the wrap address mode with non-normalized coordinates. See the programming guide section 3.2.4.3 where it stats that this is not supported. I also think that normalized coordinates are only supported with cudaBindTextureToArray, but don’t see any mention of that in the guide - maybe it was a limitation in older versions of CUDA?

3

Your texture reference is invalid because you are using the wrap address mode with non-normalized coordinates. See the programming guide section 3.2.4.3 where it stats that this is not supported. I also think that normalized coordinates are only supported with cudaBindTextureToArray, but don’t see any mention of that in the guide - maybe it was a limitation in older versions of CUDA?

4

Good point, thanks. However, changing the address mode to cudaAddressModeClamp did not help. I get the same error.

I will rather prefer not using a cudaArray, nor normalized coordinates.

5

Good point, thanks. However, changing the address mode to cudaAddressModeClamp did not help. I get the same error.

I will rather prefer not using a cudaArray, nor normalized coordinates.

6

Odd, I thought I responded to this on Friday…

On a second look, I noticed two additional problems in your bind statement: Try

cudaBindTexture(0, m_set_texture, m_p_device_set, &set_channel_descriptor, m_set_size * sizeof(float));

Attempting to bind to the address of the reference is likely what is causing your error message. Missing the *sizeof(float) would only clamp your addresses sooner than you intended :)

7

Odd, I thought I responded to this on Friday…

On a second look, I noticed two additional problems in your bind statement: Try

cudaBindTexture(0, m_set_texture, m_p_device_set, &set_channel_descriptor, m_set_size * sizeof(float));

Attempting to bind to the address of the reference is likely what is causing your error message. Missing the *sizeof(float) would only clamp your addresses sooner than you intended :)

8

Thanks!

Weird, my reply was missing too.

I made the changes that you suggested. But I still get the invalid texture reference error.

9

Thanks!

Weird, my reply was missing too.

I made the changes that you suggested. But I still get the invalid texture reference error.

10

In fact, I declare the texture variable as a local variable or class attribute.

In that case I always get the error.

I tried as a global variable, and it works.

I am not very keen on global variables, particularly in header files…

11

In fact, I declare the texture variable as a local variable or class attribute.

In that case I always get the error.

I tried as a global variable, and it works.

I am not very keen on global variables, particularly in header files…

12

Unfortunately that is how it is, according to section 3.2.4.1 of the Programming Guide:

“A texture reference is declared at file scope as a variable of type texture: […] A texture reference can only be declared as a static global variable and cannot be passed as an argument to a function.”

13

Unfortunately that is how it is, according to section 3.2.4.1 of the Programming Guide:

“A texture reference is declared at file scope as a variable of type texture: […] A texture reference can only be declared as a static global variable and cannot be passed as an argument to a function.”

14

If you don’t want to use your texture as a global variable in header files, there is a little trick I use, don’t know if it’s what you are looking for.

I declare the texture as global variable in my .cu, write a function getref to get the ref on that texture, declare the function in my header file and include this header and I can use it anywhere in my program :

*** .cu ***

texture<unsigned short, 2, cudaReadModeElementType> ref;

texture<unsigned short, 2, cudaReadModeElementType> &get_ref(void)

{

	return ref;

}

*** .h ***

texture<unsigned short, 2, cudaReadModeElementType> &get_ref(void);

*** .c ***

cudaBindTextureToArray(get_ref(), ref_ary);
15

If you don’t want to use your texture as a global variable in header files, there is a little trick I use, don’t know if it’s what you are looking for.

I declare the texture as global variable in my .cu, write a function getref to get the ref on that texture, declare the function in my header file and include this header and I can use it anywhere in my program :

*** .cu ***

texture<unsigned short, 2, cudaReadModeElementType> ref;

texture<unsigned short, 2, cudaReadModeElementType> &get_ref(void)

{

	return ref;

}

*** .h ***

texture<unsigned short, 2, cudaReadModeElementType> &get_ref(void);

*** .c ***

cudaBindTextureToArray(get_ref(), ref_ary);
16

Nice trick, I didn’t know that would work given that textures are implicit static and all that. I guess references to static variables still translate across compilation units.

My solution is to only declare texture references in the .cu file where they are used. Then I simply call cudaBindTexture prior to every single kernel call as these functions map calls from classes and it is not known what instances will call the functions in what order. Bind is very cheap: the overhead of calling bind before every single kernel launch is barely measurable.

17

Nice trick, I didn’t know that would work given that textures are implicit static and all that. I guess references to static variables still translate across compilation units.

My solution is to only declare texture references in the .cu file where they are used. Then I simply call cudaBindTexture prior to every single kernel call as these functions map calls from classes and it is not known what instances will call the functions in what order. Bind is very cheap: the overhead of calling bind before every single kernel launch is barely measurable.

18

Good to know, thanks!

19

Good to know, thanks!

20

I wonder if it’s also possible to use references to other objects in a different module, such as areas of constant storage - and whether these will actually be cached on access.