About cudaBindTexture2D

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi, I use CUDA 2.2Beta now and in the documentation, it provide a new funciton to fetch 2d global memroy by texture.(cudaBindTexture2D)

the way to use it is

cudaError_t cudaBindTexture2D (size_t *offset, const struct textureReference texref, const void devPtr, const struct cudaChannelFormatDesc * desc, size_t width, size_t height, size_t pitch)

And it describe the parameters as:

offset - Offset in bytes

texref - Texture reference to bind

devPtr - 2D memory area on device

desc - Channel format

width - Width in texel units

height - Height in texel units

pitch - Pitch in bytes

But I still don’t know how to use offset, what does it used for? And so do Pitch.

Any one know how to use it?

2.About prarameter x,y,z,w of cudaChannelFormatDesc means how many data be fetched when we read it by tex2d(ptr,x,y)?

If I want to do convolution whose filter size is 100100 and the image is 48723248 16bit unsigned , does it correct to set x=10016(bits), and y=10016(bits)

Below is my code, but it can not work and will cause memory confliciton.I don’t know why, When I use cudaBindTextureToArray, it works…

#include <conio.h>

#include <stdio.h>

#include <stdlib.h>

#include <cuda_runtime.h>

int main(void) {

	cudaChannelFormatDesc chDescUSht = cudaCreateChannelDesc(16*100,16*100, 0, 0, cudaChannelFormatKindUnsigned);

	texture<unsigned short,2> texUShtConvSource;

	unsigned short *ptrS_d;

	char  *ptrD1_d, *ptrD2_d;

	const int SizeX=4872;

	const int SizeY=3248;

	int cudaError;

	size_t TexOffset;

	size_t PitchinByte=2;	

	cudaMalloc((void**) &ptrS_d, sizeof(unsigned short) * SizeX * SizeY);

	cudaMalloc((void**) &ptrD1_d, sizeof(char) * SizeX * SizeY);

	cudaMalloc((void**) &ptrD2_d, sizeof(char) * SizeX * SizeY);

	dim3 dimBlockNorm(BLOCK_SIZEX, BLOCK_SIZEY);

	dim3 dimGridNorm(SizeX/dimBlockNorm.x+1, SizeY/dimBlockNorm.y+1);

	cudaError=cudaBindTexture2D(&TexOffset, texUShtConvSource,ptrS_d,chDescUSht,SizeX,SizeY,PitchinByte);

	if (cudaError) printf ("Failed to Bind Data \n");

	Mykernel<<<dimGridNorm,dimBlockNorm>>>(ptrD1_d,ptrD2_d);

	cudaUnbindTexture(texUShtConvSource);

	printf ("Sucess \n");

	getch();

	return 0;

}
2

I am using Cuda 2.0,
in cuda 2.0, Can not bind texture with 2D global memory array.
I don’t know that if in cuda 2.2 Beta we can do.

when I read you code. I saw pitchinByte value = 2.
Why PitchinByte = 2?
And you are used cuda Malloc to allocate global memory.
It is not correct.
In this case, if you want to bind 2D texture with 2D global memory array.

  • Allocate 2D global memory Array with cudaMallocPitch() function. remember that Pitch is very important variable.
  • bind this area to your array.
    :)
3

Hi, to bind 2d global memory to 2d texture is a new funciton in CUDA 2.2Beta.

So that cuda array is not critcally needed.

Its why I use cudaMalloc to allocate Global memory.

And about pitch, do you mean the pitch should be equal to SizeX*Number of Byte per unsigned short?

4

when you use cudaMallocPitch to allocate 2D array on global memory.

the value of Pitch variable will be setted automaticly .

in your case, if you use cudaMaloc to allocate an Array in global memory,

and try to bind this area with 2D texture, pitch is the SizeX*sizeof(your type).

:)