c++ - 在 CUDA 中编写内核函数的正确方法？

Question

我正准备着手将我编写的程序转换为 CUDA，以期提高处理速度。

现在显然我的旧程序一个接一个地执行许多函数，我在主程序中将这些函数分开并按顺序调用每个函数。

void main ()
{
  *initialization of variables*
  function1()
  function2()
  function3()
  print result;
}

这些函数本质上是串行的，因为 funtion2 依赖于 funtion1 的结果。

好吧，现在我想把这些函数转换成内核，并并行运行函数中的任务。

是不是很简单，并行重写每个函数，然后在我的主程序中，一个接一个地调用每个内核？这比需要的慢吗？例如，我是否可以让我的 GPU 直接执行下一个并行操作，而无需返回 CPU 来初始化下一个内核？

显然，我会将所有运行时变量保留在 GPU 内存中以限制正在进行的数据传输量，所以我是否应该担心内核调用之间的时间间隔？

我希望这个问题很清楚，如果不清楚请让我详细说明。谢谢。

这是一个额外的问题，这样我就可以检查我的理智了。最终这个程序的输入是一个视频文件，通过不同的函数，每一帧都会导致一个结果。我的计划是一次抓取多个帧(比如 8 个独特的帧)，然后将我拥有的 block 总数划分为这 8 个帧，然后 block 中的多个线程将对图像数据进行更多的并行操作，如 vector 加法、傅里叶变换等。
这是解决问题的正确方法吗？

Answer 1

在某些情况下，您只需从普通 CPU 版本进行很少的移植工作，就可以让程序在 GPU 上以全速运行，这可能就是其中之一。

如果你有这样的功能是可能的:

void process_single_video_frame(void* part_of_frame)
{
  // initialize variables
  ...
  intermediate_result_1 = function1(part_of_frame);
  intermediate_result_2 = function2(intermediate_result_1);
  intermediate_result_3 = function3(intermediate_result_2);
  store_results(intermediate_result_3);
}

并且您可以同时处理多个part_of_frames。比如说，几千，

和 function1()、function2() 和function3() 执行几乎相同的代码路径(即程序流程不严重依赖于框架的内容)，

然后，本地内存可以为您完成所有工作。本地内存是一种存储在全局内存中的内存。它以一种微妙而深刻的方式与全局内存不同......内存以这样一种方式简单地交错，相邻线程将访问相邻的 32 位字，如果线程都从中读取，则可以完全合并内存访问他们本地内存阵列的相同位置。

您的程序流程是，您首先将 part_of_frame 复制到本地数组，然后为其他本地数组准备中间结果。然后，您可以在代码中的各个函数之间传递指向本地数组的指针。

一些伪代码:

const int size_of_one_frame_part = 1000;

__global__ void my_kernel(int* all_parts_of_frames) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int my_local_array[size_of_one_frame_part];
    memcpy(my_local_array, all_parts_of_frames + i * size_of_one_frame_part);
    int local_intermediate_1[100];
    function1(local_intermediate_1, my_local_array);
    ...
}

__device__ void function1(int* dst, int* src) {
   ...
}

总而言之，这种方法可以让您使用几乎没有变化的 CPU 函数，因为并行性不是来自创建函数的并行版本，而是并行运行整个函数链。硬件支持在本地数组中交错内存，这再次成为可能。

注意事项:

• part_of_frame 从全局内存到本地内存的初始拷贝没有合并，但希望您有足够的计算来隐藏它。

• 在计算能力 <= 1.3 的设备上，每个线程只有 16KiB 的本地内存可用，这可能不足以容纳您的 part_of_frame 和其他中间数据。但是在 >= 2.0 的计算能力上，这已经扩展到 512KiB，这应该足够了。