Texture Bilinear filtering

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

I use textures to get the interpolated value of a pixel of this texture.

My problem is that some parts of my image, the interpolated values are negative.

I can not explain the phenomenon.

texture’s configuration:

myTexture.filterMode = cudaFilterModeLinear;

myTexture.normalized = false;

Do I forget something?

Thx for your help.

Etienne

2

The interpolated value should only be negative if you have negative values in the array. What texture type are you using?

You’ll have to post more code than that if you want anyone to help.

3

declaration and fonction to access texture in kernel and main :

texture<float, 2, cudaReadModeElementType> dataTest_tex;

texture<float, 2, cudaReadModeElementType> &getTexture() { return dataTest_tex; }

static __inline__ __device__ texture<float, 2, cudaReadModeElementType> &getDeviceTexture() { return dataTest_tex; }

Kernel using the texture :

__global__ void bilinear(float *d_out, int dimIn, int dimOut, float ax, float ay, float bx, float by, float cx, float cy) {

	int ix		= IMUL(blockDim.x, blockIdx.x) + threadIdx.x;

	int iy		= IMUL(blockDim.y, blockIdx.y) + threadIdx.y;

	float xp	= ax*ix + bx*iy + cx;

	float yp	= ay*ix + by*iy + cy;

	int w1		= int(xp);

	int h1		= int(yp);

	d_out[iy*dimOut + ix] = tex2D(getDeviceTexture(), w1, h1);

}

Kernel call :

dim3	blockSizeEXT(16,16);

dim3	gridSizeEXT(512/16,512/16);

bilinear<<< gridSizeEXT, blockSizeEXT >>>(d_DataTemp, SIZE_TEST, SIZE_REF, tabParam[i].ax, tabParam[i].ay, tabParam[i].bx, tabParam[i].by, tabParam[i].cx, tabParam[i].cy);

I use a cudaArray to bind texture with datas :

float *h_DataTest;	// <= datas are read from raw file

cudaArray	*d_dataTest;

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();

cudaMallocArray( &d_dataTest, &channelDesc, SIZE_TEST, SIZE_TEST );

cudaMemcpyToArray( d_dataTest, 0, 0, (void*)h_DataTest, test_size, cudaMemcpyHostToDevice);

texture<float, 2, cudaReadModeElementType> &myTexture = getTexture();

myTexture.filterMode = cudaFilterModeLinear;

myTexture.normalized = false;

cudaBindTextureToArray( myTexture, d_dataTest, channelDesc);

I am sure there is no negative value in *h_DataTest but I do not know if the datas can be modified when I copy them.

ps : sorry for my bad english.