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

3

我希望使用纹理内存来解决插值问题,希望比使用全局内存更快。由于这是我第一次使用纹理内存,因此我将我的插值问题简化为线性插值问题。所以,我已经意识到有比下面报告的更聪明和更快的方法来进行线性插值。

这里是文件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。

奇怪的是,尽管N>16,但仅在前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纹理替换为一个浮点数(强制数据为实数),我也遇到了同样的问题。
有什么想法吗?

BLOCK_SIZE是什么?M是什么?N是什么?如果您只是发布包含API调用错误检查的小型自包含示例代码,那将会更简单,这样其他人就可以研究、编译和运行它。 - talonmies
1个回答

3
我可以看到你的程序中以下行存在逻辑错误。
cudaBindTexture(NULL, data_d_texture, data_d, M);
cudaBindTexture函数的最后一个参数需要以字节为单位指定数据的大小,而您正在指定元素数量。
您应该尝试以下操作:
cudaBindTexture(NULL, data_d_texture, data_d, M * sizeof(cuComplex));

那绝对解决了问题。非常感谢。现在我可以将这个可行的简单示例作为更复杂的插值问题的起点。 - Vitality
3
作为一种健壮的编程实践,我建议永远不要将NULL作为cudaBindTexture的第一个参数传递。相反,传递一个指向size_t对象的指针,这样代码可以处理纹理偏移量的非零值,以防发生错误。传递NULL会引发日后潜在的静默失败。 - njuffa
@njuffa,我唯一要提到的例外是,如果您正在传递cudaMalloc()或cudaMallocPitch()的返回值,则始终应返回符合纹理基地址对齐要求的指针。 - ArchaeaSoftware
cudaMalloc() 可能与消耗指针的 cudaBindTexture() 不在紧密的文本接近度内,或者这两个调用在源代码重构期间可能被分开。然后下一个程序员可能会引入由 cudaMalloc() 返回的指针的偏移量,突然之间对 cudaBindTexture() 的调用可能会默默失败。 - njuffa

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接