1

我正在使用“标准”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 倍以上。但是，我注意到结果中存在一些不匹配，在 about 的两个实现之间存在均方根误差0.07%。

CUDA C 编程指南说插值系数以 9 位定点格式存储，带有 8 位小数值，这可能是导致不匹配的原因。

然后我有两个问题：

1）是否有任何“技巧”来提高“基于纹理”插值的准确性？

2）我认为即使我移动到 float4，这种 9 位表示也会将精度限制在此处获得的精度，对吗？换句话说，将数字表示精度从 float2 提高到 float4 没有意义吗？

提前致谢。

4

1 回答 1

2

你可以“预插值”你的纹理来增加你的分辨率，即如果你的初始纹理是 100x100，那么你可以预插值使其成为 200x200，那么你已经将内核插值的分辨率加倍。

于 2013-01-22T15:09:21.403 回答