CUDA：并行化具有嵌套循环的函数调用的多个嵌套for循环

Question

问题

我有兴趣使用CUDA 并行化问题。有问题的C 代码遵循以下简化形式：

 int A, B, C; // 100 < A,B,C,D < 1,000

 float* v1, v2, v3;   
 //v1,v2, v3 will have respective size A,B,C
 //and will not be empty

 float*** t1, t2, t3; 
 //t1,t2,t3 will eventually have the size (ci,cj,ck)
 //and will not be empty

 int i, j , k, l;      

 float xi, xj, xk;

 for (i=0; i<A; ++i){   
   xi = ci - v1[i];
   for (j=0; j<B; ++j){
     xj = (j*cj)*cos(j*M_PI/180);   
     for (k=0;k<C; ++k){
       xk = xj - v3[k];
       if (xk < xi){
         call_1(t1[i], v1, t2[i], &t3[i][j][k]);
       }
       else t3[i][j][k] = some_number;
     }   
   } 
 }

这里是call_1

void call_1 (float **w, float *x, float **y, float *z){
 int k, max = some_value;
 float *v; //initialize to have size max
 for (k=0; k<max; ++k)
    call_2(x[k], y[k], max, &v[k]);
 call_2(y, v, max, z);
}

这里call_2是

void call_2 (float *w, float*x, int y, double *z)

在单个while 循环中仅包含位移、乘法、减法和加法等操作。

尝试的想法

到目前为止，我的想法是，函数call_1可以转化为内核代码__global__ void call_1；并且call_2 可以在不修改其内容的情况下转换为设备代码。特别是，我大概可以让__global__ void call_1 成为

double* v; //initialize to have size max

int index = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int k=index; k<max; k += stride)
    call_2 (x[k], y[k], max, &v[k]);

__syncthreads();

call_2 (y, v, max, z);
free (v);

我部分知道可以通过使用threadIdx, blockIdx, and gridDim 的组合来删除 for 循环，但我特别不确定如何特别是问题包含一个也使用函数调用的函数调用。

Answer 1

嗯，有两个可能的答案，虽然我没有勇气为你搜索所有这些，但我仍然会做一个答案，因为你似乎被公然忽视了。 ：/

首先。
最近的 CUDA API 和 nvidia 架构支持函数调用，甚至支持 CUDA 中的递归。我不确定它是如何工作的，因为我自己从未使用过它，但你可能想研究一下。 （或者做一些 Vulkan，因为它看起来很有趣并且也支持它。）

可能会帮助你：https://devtalk.nvidia.com/default/topic/493567/cuda-programming-and-performance/calling-external-kernel-from-cuda/ 以及其他具有相关关键字的东西。 : D 另一方面..
在解决简单问题时，特别是如果像我一样，你宁愿花时间编程而不是研究和学习一些随机的 API，你总是可以使用更原始的解决方案，只使用你正在使用的语言的基础。

在您的情况下，我将简单地内联对函数的调用以创建单个 CUDA 内核，因为它看起来很容易做到。

是的，对，如果对函数有多次调用，它可能包括一些复制粘贴...如果它能让你轻松高效地解决一个简单的问题并去做更多的事情，这并不重要富有成效。

int index = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int k=index; k<max; k += stride)
    call_2 (x[k], y[k], max, &v[k]); // Insert call_2 code here instead.

解决此问题的另一种方法是，当您确信自己的数据足够大，即使将代码和数据从 CPU 和 RAM 传递到 GPU 时，性能也会有很好的提升，那就是简单地有多个“波” cuda 内核调用。

您让第一波处理，同时准备第二波，然后在完成的第一波中启动。

它基本上等同于最近的 CUDA 实现提供的其他更智能的构造，因此您可能会通过一些研究找到更智能的事情，但话又说回来......取决于您的优先级。

但是，手动内联函数很棒。 :D
*几乎从不，但它可以很方便