Pitched Texture delivers wrong results?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I believe I am doing something fundamentally wrong when binding pitched memory to 2D textures. I have a kernel that works fine, when the pitch equals the width. Somehow the texture seems to deliver wrong values when the pitch is not equal to the width.

//Texture reference to work represent image data

        texture<unsigned char, 2, cudaReadModeElementType> bayTex;

//Bind Texture to reference

	const textureReference* bayTexPtr;

	cudaGetTextureReference(&bayTexPtr, "bayTex");

	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned char>();

	//Set acces parameters

	bayTex.addressMode[0] = cudaAddressModeClamp;

	bayTex.addressMode[1] = cudaAddressModeClamp;

	//Bind pitched data to texture

	unsigned int width_px = source->d_width / (source->d_padding * source->element_size);

	cudaBindTexture2D(NULL, bayTexPtr, source->d_pointer, &channelDesc, width_px, source->d_height, source->d_pitch);

I am fetching my values with tex2D(bayTex, x, y). When I put in the same image with different crop (resulting in different pitch), the same texture position returns different results.

Any ideas?

Kwyjibo

P.S.: What about that offset parameter? Do I need to use it?

2

Hi,

I believe I am doing something fundamentally wrong when binding pitched memory to 2D textures. I have a kernel that works fine, when the pitch equals the width. Somehow the texture seems to deliver wrong values when the pitch is not equal to the width.

//Texture reference to work represent image data

        texture<unsigned char, 2, cudaReadModeElementType> bayTex;

//Bind Texture to reference

	const textureReference* bayTexPtr;

	cudaGetTextureReference(&bayTexPtr, "bayTex");

	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned char>();

	//Set acces parameters

	bayTex.addressMode[0] = cudaAddressModeClamp;

	bayTex.addressMode[1] = cudaAddressModeClamp;

	//Bind pitched data to texture

	unsigned int width_px = source->d_width / (source->d_padding * source->element_size);

	cudaBindTexture2D(NULL, bayTexPtr, source->d_pointer, &channelDesc, width_px, source->d_height, source->d_pitch);

I am fetching my values with tex2D(bayTex, x, y). When I put in the same image with different crop (resulting in different pitch), the same texture position returns different results.

Any ideas?

Kwyjibo

P.S.: What about that offset parameter? Do I need to use it?

3

Just tried the minimal example and it works. I am relly confused right now. The error must be somewhere else.

[/code]

#include <stdio.h>

//Texture reference to work represent image data

texture<unsigned char, 2, cudaReadModeElementType> arr1Tex;

texture<unsigned char, 2, cudaReadModeElementType> arr2Tex;

//Test texture fetch

global void textureTestKernel()

{

printf("arr1Tex (3,3): %i\n", tex2D(arr1Tex, 3, 3));

printf("arr2Tex (3,3): %i\n", tex2D(arr2Tex, 3, 3));

}

int main()

{

//Allocate arrays

unsigned char arr1[64*64] = {0};

arr1[64*3+3] = 1;

unsigned char* d_pointer_arr1;

size_t pitch_arr1;



unsigned char arr2[50*50] = {0};

arr2[50*3+3] = 1;

unsigned char* d_pointer_arr2;

size_t pitch_arr2;



//Upload arrays

cudaMallocPitch((void**) &d_pointer_arr1, &pitch_arr1, 64, 64);

cudaMallocPitch((void**) &d_pointer_arr2, &pitch_arr2, 50, 50);

cudaMemcpy2D(d_pointer_arr1, pitch_arr1, arr1, 64, 64, 64, cudaMemcpyHostToDevice);

cudaMemcpy2D(d_pointer_arr2, pitch_arr2, arr2, 50, 50, 50, cudaMemcpyHostToDevice);



//Print pitch values

printf("pitch1: %i\n", pitch_arr1);

printf("pitch2: %i\n", pitch_arr2);

//Bind arr1 to reference

const textureReference* arr1TexPtr;

cudaGetTextureReference(&arr1TexPtr, "arr1Tex");

cudaChannelFormatDesc channelDesc1 = cudaCreateChannelDesc<unsigned char>();

//Set acces parameters

arr1Tex.addressMode[0] = cudaAddressModeClamp;

arr1Tex.addressMode[1] = cudaAddressModeClamp;

//Bind pitched data to texture

cudaBindTexture2D(NULL, arr1TexPtr, d_pointer_arr1, &channelDesc1, 64, 64, pitch_arr1);



//Bind arr2 to reference

const textureReference* arr2TexPtr;

cudaGetTextureReference(&arr2TexPtr, "arr2Tex");

cudaChannelFormatDesc channelDesc2 = cudaCreateChannelDesc<unsigned char>();

//Set acces parameters

arr2Tex.addressMode[0] = cudaAddressModeClamp;

arr2Tex.addressMode[1] = cudaAddressModeClamp;

//Bind pitched data to texture

cudaBindTexture2D(NULL, arr2TexPtr, d_pointer_arr2, &channelDesc2, 50, 50, pitch_arr1);

//Start Kernel

textureTestKernel<<<1, 1>>>();

}[/code]

4

Just tried the minimal example and it works. I am relly confused right now. The error must be somewhere else.

[/code]

#include <stdio.h>

//Texture reference to work represent image data

texture<unsigned char, 2, cudaReadModeElementType> arr1Tex;

texture<unsigned char, 2, cudaReadModeElementType> arr2Tex;

//Test texture fetch

global void textureTestKernel()

{

printf("arr1Tex (3,3): %i\n", tex2D(arr1Tex, 3, 3));

printf("arr2Tex (3,3): %i\n", tex2D(arr2Tex, 3, 3));

}

int main()

{

//Allocate arrays

unsigned char arr1[64*64] = {0};

arr1[64*3+3] = 1;

unsigned char* d_pointer_arr1;

size_t pitch_arr1;



unsigned char arr2[50*50] = {0};

arr2[50*3+3] = 1;

unsigned char* d_pointer_arr2;

size_t pitch_arr2;



//Upload arrays

cudaMallocPitch((void**) &d_pointer_arr1, &pitch_arr1, 64, 64);

cudaMallocPitch((void**) &d_pointer_arr2, &pitch_arr2, 50, 50);

cudaMemcpy2D(d_pointer_arr1, pitch_arr1, arr1, 64, 64, 64, cudaMemcpyHostToDevice);

cudaMemcpy2D(d_pointer_arr2, pitch_arr2, arr2, 50, 50, 50, cudaMemcpyHostToDevice);



//Print pitch values

printf("pitch1: %i\n", pitch_arr1);

printf("pitch2: %i\n", pitch_arr2);

//Bind arr1 to reference

const textureReference* arr1TexPtr;

cudaGetTextureReference(&arr1TexPtr, "arr1Tex");

cudaChannelFormatDesc channelDesc1 = cudaCreateChannelDesc<unsigned char>();

//Set acces parameters

arr1Tex.addressMode[0] = cudaAddressModeClamp;

arr1Tex.addressMode[1] = cudaAddressModeClamp;

//Bind pitched data to texture

cudaBindTexture2D(NULL, arr1TexPtr, d_pointer_arr1, &channelDesc1, 64, 64, pitch_arr1);



//Bind arr2 to reference

const textureReference* arr2TexPtr;

cudaGetTextureReference(&arr2TexPtr, "arr2Tex");

cudaChannelFormatDesc channelDesc2 = cudaCreateChannelDesc<unsigned char>();

//Set acces parameters

arr2Tex.addressMode[0] = cudaAddressModeClamp;

arr2Tex.addressMode[1] = cudaAddressModeClamp;

//Bind pitched data to texture

cudaBindTexture2D(NULL, arr2TexPtr, d_pointer_arr2, &channelDesc2, 50, 50, pitch_arr1);

//Start Kernel

textureTestKernel<<<1, 1>>>();

}[/code]

5

Okay,

I’m trying to hunt this bug down.

The values of the first row are fetched correctly, but as soon as the kernel fetches values from the second or third row I get wrong values.

Has anyone ever experienced auch a behaviour?

Regards,

Kwyjibo

6

Okay,

I’m trying to hunt this bug down.

The values of the first row are fetched correctly, but as soon as the kernel fetches values from the second or third row I get wrong values.

Has anyone ever experienced auch a behaviour?

Regards,

Kwyjibo