【发布时间】：2013-01-22 23:09:21
【问题描述】：
我正在比较使用“标准”CUDA 实现和复数 (float2) 上“基于纹理”的 CUDA 实现的一维线性插值。
“标准”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 编程指南说插值系数以 9 位定点格式存储，带有 8 位小数值，这可能是导致不匹配的原因。
我有两个问题：
1) 有没有什么“技巧”可以提高“基于纹理”插值的准确性？
2) 我认为即使我移动到 float4，这种 9 位表示也会将精度限制在此处获得的精度，对吧？也就是说，从float2提升到float4的数字表示精度没有意义吗？
提前致谢。
【问题讨论】：