linear filtering on float2 texture?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi, for texture linear filtering, I have read from cuda document that it only available on floating point texture. But does it apply on float2 or float4 texture?

2

You’re not exactly right - linear filtering can be used on any texture that returns floating point data - i.e. you can use it with integer formats, but only if the read mode is set to cudaReadModeNormalizedFloat.

3

So can I do linear interpolation with float4 using texture?

4

Yes.

5

You can do the following:

texture<float2, 1, cudaReadModeElementType> data_d_texture_filtering;

__global__ void linear_interpolation_kernel_function_GPU_texture_filtering(float2* __restrict__ result_d, const float* __restrict__ x_out_d, const int M, const int N)
{
   int j = threadIdx.x + blockDim.x * blockIdx.x; 
   if(j<N) result_d[j] = tex1D(data_d_texture_filtering,float(x_out_d[j]+M/2+0.5));
}

and

void linear_interpolation_function_GPU_texture_filtering(float2* result_d, float2* data, float* x_in_d, float* x_out_d, int M, int N){
   
	cudaArray* data_d = NULL; cudaMallocArray (&data_d, &data_d_texture_filtering.channelDesc, M, 1);
	cudaMemcpyToArray(data_d, 0, 0, data, sizeof(float2)*M, cudaMemcpyHostToDevice);
	cudaBindTextureToArray(data_d_texture_filtering, data_d);
	data_d_texture_filtering.normalized = false;
	data_d_texture_filtering.filterMode = cudaFilterModeLinear;
   
	dim3 dimBlock(BLOCK_SIZE,1); dim3 dimGrid(N/BLOCK_SIZE + (N%BLOCK_SIZE == 0 ? 0:1),1);
   	linear_interpolation_kernel_function_GPU_texture_filtering<<<dimGrid,dimBlock>>>(result_d, x_out_d, M, N);
}