使用 cuda 纹理内存进行一维插值

Question

我正在尝试使用纹理内存来解决插值问题，希望以比使用全局内存更快的方式。作为我第一次使用纹理内存，我将插值问题过度简化为线性插值问题。所以，我已经知道有比下面报道的更聪明、更快速的线性插值方法。 这是文件 Kernels_Interpolation.cuh。 __device__ 函数 linear_kernel_GPU 为简单起见被省略，但正确。

texture<cuFloatComplex,1> data_d_texture;

__global__ void linear_interpolation_kernel_function_GPU_texture(cuComplex* result_d, float* x_in_d, float* x_out_d, int M, int N)
{    
   int j = threadIdx.x + blockDim.x * blockIdx.x;

   cuComplex datum;

   if(j<N)
   {
       result_d[j] = make_cuComplex(0.,0.);
       for(int k=0; k<M; k++)
       {
           datum = tex1Dfetch(data_d_texture,k);
           if (fabs(x_out_d[j]-x_in_d[k])<1.) result_d[j] = cuCaddf(result_d[j],cuCmulf(make_cuComplex(linear_kernel_GPU(x_out_d[j]-x_in_d[k]),0.),datum));
       }  
   } 
}

这里是 Kernels_Interpolation.cu 函数

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

   cudaBindTexture(NULL, data_d_texture, data_d, M);

   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_in_d, x_out_d, M, N);

}

最后，在主程序中，data_d数组的分配和初始化如下

cuComplex* data_d;      cudaMalloc((void**)&data_d,sizeof(cuComplex)*M);
cudaMemcpy(data_d,data,sizeof(cuComplex)*M,cudaMemcpyHostToDevice);

result_d 数组的长度为 N。

奇怪的是，仅在前 16 个位置上正确计算了输出，尽管 N>16，其他为 0，例如

result.r[0] 0.563585 result.i[0] 0.001251 
result.r[1] 0.481203 result.i[1] 0.584259
result.r[2] 0.746924 result.i[2] 0.820994
result.r[3] 0.510477 result.i[3] 0.708008
result.r[4] 0.362980 result.i[4] 0.091818
result.r[5] 0.443626 result.i[5] 0.984452
result.r[6] 0.378992 result.i[6] 0.011919
result.r[7] 0.607517 result.i[7] 0.599023
result.r[8] 0.353575 result.i[8] 0.448551
result.r[9] 0.798026 result.i[9] 0.780909
result.r[10] 0.728561 result.i[10] 0.876729
result.r[11] 0.143276 result.i[11] 0.538575
result.r[12] 0.216170 result.i[12] 0.861384
result.r[13] 0.994566 result.i[13] 0.993541
result.r[14] 0.295192 result.i[14] 0.270596
result.r[15] 0.092388 result.i[15] 0.377816
result.r[16] 0.000000 result.i[16] 0.000000
result.r[17] 0.000000 result.i[17] 0.000000
result.r[18] 0.000000 result.i[18] 0.000000
result.r[19] 0.000000 result.i[19] 0.000000

其余的代码是正确的，也就是说，如果我用使用全局内存的函数替换 linear_interpolation_kernel_function_GPU_texture 和 linear_interpolation_function_GPU_texture 一切都很好。

我已经验证我可以正确访问纹理内存直到某个位置（取决于 M 和 N），例如 64，之后它返回 0。

如果我将 cuComplex 纹理替换为浮点纹理（强制数据为真实），我也会遇到同样的问题。

有什么想法吗？

Answer 1

我可以在您的程序的以下行中看到一个逻辑错误。

cudaBindTexture(NULL, data_d_texture, data_d, M);

cudaBindTexture 的最后一个参数采用字节为单位的数据大小，您指定了元素的数量。

您应该尝试以下方法：

cudaBindTexture(NULL, data_d_texture, data_d, M * sizeof(cuComplex));