GPU（CUDA）上的非线性优化，没有数据传输延迟

Question

我正在尝试完全在GPU上执行非线性优化问题。目标函数的计算以及从GPU到CPU的数据传输是瓶颈。为了解决这个问题，我想

1. 目标的高度并行计算和
2. 在GPU上执行整个优化。

更具体地说，问题在伪代码中如下：

x = x0  // initial guess of the vector of unknowns, typically of size ~10,000
for iteration = 1 : max_iter
      D = compute_search_direction(x)
      alpha = compute_step_along_direction(x)
      x = x   +   D * alpha  // update
end for loop

compute_search_direction(x)和compute_step_along_direction(x)函数每次迭代都调用目标函数f0(x)数十次。目标函数是一个复杂的CUDA内核，基本上是正向Bloch仿真（=描述磁场中核自旋动力学的方程组）。 f0(x)的输出为F（目标函数的值，标量）和DF（雅可比，或一阶导数的向量，大小与x相同，即〜10,000）。在GPU上，f0(x)确实非常快，但是将x从CPU传输到GPU，然后将F和DF从GPU传输回CPU需花费一段时间（总计约1秒）。由于该函数每次迭代被调用数十次，因此导致整体优化非常缓慢。

理想情况下，我希望在GPU上具有完整的伪代码。我现在能想到的唯一解决方案是递归内核。上面的伪代码将是“外部内核”，它以线程数= 1和块数= 1启动（即，该内核实际上不是并行的...）。每次需要评估目标函数和一阶导数的向量时，该内核便会调用目标函数（即“内部内核”，该内核是大规模并行的）。由于内核启动是异步的，因此我可以强制GPU等待，直到对f0内部内核进行完全评估，以移至外部内核的下一条指令（使用同步点）。

在某种意义上，这与常规的CUDA编程实际上相同，在常规的CUDA编程中，CPU控制内核启动以评估目标函数f0，不同之处在于CPU被未被并行化的外部内核代替（1个线程，1个块）。但是，由于所有内容都在GPU上，因此不再存在数据传输延迟。

我现在在一个简单的示例上测试这个想法，以测试可行性。但是，这似乎很麻烦……我的问题是：

1. 这对其他人有意义吗？
2. 是否有更直接的方法来获得相同的结果而又不增加嵌套内核的复杂性？

Answer 1

[您似乎混淆了“减少GPU和CPU之间的内存传输”，以及“使整个代码在设备上运行（也称为gpu）”。

为了减少内存传输，您不需要使整个代码在GPU上运行。

您可以一次将数据复制到GPU，然后在GPU代码和CPU代码之间来回切换。只要您不尝试从CPU代码访问任何GPU内存（反之亦然），就可以了。

这是您要执行的操作的正确方法的伪代码。

// CPU code
cudaMalloc(&x,...) //allocate memory for x on GPU
cudaMemCpy(x, x0, size, cudaMemCpyHostToDevice); //Copy x0 to the freshly allocated array 
cudaMalloc(&D, ....)    //allocate D and alpha before the loop
cudaMalloc(&alpha, ....)
for iteration = 1 : max_iter
      compute_search_direction<<<...>>>(x, D) //Call a kernel that does the computation and stores the result in D
      compute_step_along_direction<<<....>>>(x, alpha)
      combine_result<<<...>>>(x, D, alpha)  // x   +   D * alpha
end for loop
//Eventually copy x on CPU, if need be

希望有帮助！