通過CUDA紋理插值進行一維線性插值的精度

Question

我正在比較使用「標準」CUDA實現和基於「基於紋理」的CUDA實現的複數（float2）的一維線性插值。

「標準」 CUDA實現包括下列行：

/*************************************/ 
/* LINEAR INTERPOLATION KERNEL - GPU */ 
/*************************************/ 
__device__ float linear_kernel_GPU(float in) 
{ 
    float d_y; 
    return 1.-abs(in); 
} 

/**********************************************/ 
/* LINEAR INTERPOLATION KERNEL FUNCTION - GPU */ 
/**********************************************/ 
__global__ void linear_interpolation_kernel_function_GPU(float2* result_d, float2* data_d, float* x_in_d, float* x_out_d, int M, int N) 
{ 
    int j = threadIdx.x + blockDim.x * blockIdx.x; 

    if(j<N) 
    { 
     result_d[j].x = 0.; 
     result_d[j].y = 0.; 
     for(int k=0; k<M; k++) 
     { 
      if (fabs(x_out_d[j]-x_in_d[k])<1.) { 
       result_d[j].x = result_d[j].x + linear_kernel_GPU(x_out_d[j]-x_in_d[k])*data_d[k].x; 
       result_d[j].y = result_d[j].y + linear_kernel_GPU(x_out_d[j]-x_in_d[k])*data_d[k].y; } 
     } 
    } 
} 

extern "C" void linear_interpolation_function_GPU(cuComplex* result_d, cuComplex* data_d, float* x_in_d, float* x_out_d, int M, int N){ 

    dim3 dimBlock(BLOCK_SIZE,1); dim3 dimGrid(N/BLOCK_SIZE + (N%BLOCK_SIZE == 0 ? 0:1),1); 
    linear_interpolation_kernel_function_GPU<<<dimGrid,dimBlock>>>(result_d, data_d, x_in_d, x_out_d, M, N); 

} 

的「基於紋理的」 CUDA實現包括下列行：

texture<float2, 1, cudaReadModeElementType> data_d_texture; 

// ********************************************************/ 
// * LINEAR INTERPOLATION KERNEL FUNCTION - GPU - TEXTURE */ 
// ********************************************************/ 
__global__ void linear_interpolation_kernel_function_GPU_texture(cuComplex* result_d, float* x_out_d, int M, int N) 
{ 
    int j = threadIdx.x + blockDim.x * blockIdx.x; 

    if(j<N) result_d[j] = tex1D(data_d_texture,float(x_out_d[j]+M/2+0.5)); 

} 

// *************************************************/ 
// * LINEAR INTERPOLATION FUNCTION - GPU - TEXTURE */ 
// *************************************************/ 
extern "C" void linear_interpolation_function_GPU_texture(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.channelDesc, M, 1); 
    cudaMemcpyToArray(data_d, 0, 0, data, sizeof(float2)*M, cudaMemcpyHostToDevice); 
    cudaBindTextureToArray(data_d_texture, data_d); 
    data_d_texture.normalized = false; 
    data_d_texture.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<<<dimGrid,dimBlock>>>(result_d, x_out_d, M, N); 

} 

的「基於紋理的」內插是比「標準」快20倍以上。但是，我注意到結果中有一些不匹配，兩種實現之間的均方根誤差約爲0.07%。

CUDA C編程指南指出，插值係數以8位小數值的格式存儲在9位定點格式中，這可能是導致不匹配的原因。

我有那麼兩個問題：

1）是否有「貓膩」提升「基於紋理的」插值的準確性？

2）我認爲這個9位表示會限制這裏得到的精度，即使我移動到float4，對不對？換句話說，從float2到float4提高數字表示的準確性沒有意義？

在此先感謝。

Answer 1

你可以「預插值」你的紋理來提高你的分辨率，也就是說，如果你的初始紋理是100x100，那麼你可以預插值使它成爲200x200，那麼你已經加倍了內核插值的分辨率。