优化 CUDA 插值答案

【问题标题】：Optimizing CUDA interpolation优化 CUDA 插值
【发布时间】：2013-02-22 14:29:16
【问题描述】：

我已经使用 CUDA 开发了以下插值，我正在寻找一种改进这种插值的方法。由于某些原因，我不想使用 CUDA 纹理。

由于某些未知原因，我注意到的另一点是，在我的情况下，如果向量的大小优于线程数（例如，向量大小为 1000，线程数等于 512。线程执行其第一项工作，仅此而已。我想优化 singleInterp 函数。

这是我的代码：

__device__ float singleInterp(float* data, float x, int lx_data) { 

float res = 0;    
int i1=0;
int j=lx_data; 
int imid;

while (j>i1+1)  
{  
    imid = (int)(i1+j+1)/2;
    if (data[imid]<x) 
         i1=imid;
    else 
        j=imid;
} 
if (i1==j)
    res = data[i1+lx_data];
else
    res =__fmaf_rn( __fdividef(data[j+lx_data]-data[i1+lx_data],(data[j]-data[i1])),x-data[i1], data[i1+lx_data]);

return res;

}

内核：

__global__ void linearInterpolation(float* data, float* x_in, int lx_data) {

int i = threadIdx.x + blockDim.x * blockIdx.x; 
int index = i; 
if (index < lx_data) 
    x_in[index] = singleInterp(data, x_in[index], lx_data);
}

【问题讨论】：

你用多少块启动？你能列出你的内核设置和调用代码吗？考虑到要处理的项目数以及每个块有多少线程，您需要正确计算所需的块数。在上面的示例中，您需要 2 个块。

标签： performance optimization cuda interpolation

【解决方案1】：

您似乎对一维线性插值感兴趣。我已经遇到了优化这种插值的问题，我最终得到了以下代码

__global__ void linear_interpolation_kernel_function_GPU(double* __restrict__ result_d, const double* __restrict__ data_d, const double* __restrict__ x_out_d, const int M, const int N)
{
    int j = threadIdx.x + blockDim.x * blockIdx.x;

    if(j<N)
    {
        double reg_x_out = x_out_d[j/2]+M/2;
        int k = floor(reg_x_out);
        double a = (reg_x_out)-floor(reg_x_out);
        double dk = data_d[2*k+(j&1)];
        double dkp1 = data_d[2*k+2+(j&1)];
        result_d[j] = a * dkp1 + (-dk * a + dk);
    } 
 }

假设数据在-M/2 和M/2 之间的整数节点处采样。该代码“等效于”一维纹理插值，如下面的web-page 所述。对于一维线性纹理插值，请参见CUDA-Programming-Guide 的图 13。不同解决方案的比较，请看以下thread。

【讨论】：