使用CUDA纹理内存进行插值

Question

我想使用纹理内存进行数据插值。我有两个数组（即A[I]和B[I]），我想在它们之间插入Data。我想我可以将它们绑定到纹理内存并设置插值，但我不确定如何做到这一点。

CUDA附带的示例使用A[i-1]和A[i+1]进行插值。

有什么办法可以实现我的计划吗？我尝试这个是因为我认为我可以得到一个很好的加速。

Answer 1

是的，你可以用纹理内存做到这一点，而且速度很快。我个人使用ArrayFire来完成这些操作，因为它比我希望的手动编码更快。

如果你想在CUDA中手工编码，你需要这样的东西：

// outside kernel
texture<float,1>  A;
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
cudaArray *arr = NULL;
cudaError_t e = cudaMallocArray(&arr, &desc, 1, length);
A.filterMode = cudaFilterModePoint;
A.addressMode[0] = cudaAddressModeClamp;
cudaBindTextureToArray(A, arr, desc);
...
// inside kernel
    
valA = tex1D(A,1,idx)
valB = tex1D(B,1,idx)
float f = 0.5;
output  = (f)*valA + (1-f)*valB;

如果你只想插入ArrayFire（根据我的经验，它比我手动编码的更快，更不用说使用起来更简单了），那么你会想要：

// in arrayfire
array A = randu(10,1);
array B = randu(10,1);
float f = 0.5;
array C = (f)*A + (1-f)*B;

以上假设您希望在2个不同数组或矩阵的相应索引之间进行插值。还有其他插值函数可用。

Answer 2

如果您不习惯使用CUDA进行开发，那么使用纹理内存并不是最容易的事情。

我建议您尝试在CUDA中编写算法的第一个并行版本，而不进行优化。然后，在应用程序上使用NVIDIA Visual Profiler来确定是否需要设置纹理内存以优化内存访问。

请记住，优化得越早，调试就越困难。

最后但并非最不重要的是，最新的CUDA版本（CUDA 5，仍在候选版本中）能够自动将数据存储在纹理内存中，只要您将作为参数传递给内核的输入缓冲区声明为const restrict指针。