使用多个 GPU OpenCL

Question

我有一个循环，在其中我将多个内核启动到 GPU 上。下面是片段：

for (int idx = start; idx <= end ;idx ++) {

            ret = clEnqueueNDRangeKernel(command_queue, memset_kernel, 1, NULL,
                                            &global_item_size_memset, &local_item_size, 0, NULL, NULL);
            ASSERT_CL(ret, "Error after launching 1st memset_kernel !");


            ret = clEnqueueNDRangeKernel(command_queue, cholesky_kernel, 1, NULL,
                                                    &global_item_size_cholesky, &local_item_size, 0, NULL, NULL);
            ASSERT_CL(ret, "Error after launching 1st cholesky_kernel !");


            ret = clEnqueueNDRangeKernel(command_queue, ckf_kernel1, 1, NULL,
                                            &global_item_size_kernel1, &local_item_size, 0, NULL,  NULL);
            ASSERT_CL(ret, "Error after launching ckf_kernel1[i] !");



            clFinish(command_queue);
            ret = clEnqueueNDRangeKernel(command_queue, memset_kernel, 1, NULL,
                                            &global_item_size_memset, &local_item_size, 0, NULL, NULL);
            ASSERT_CL(ret, "Error after launching 2nd memset_kernel !");


            ret = clEnqueueNDRangeKernel(command_queue, cholesky_kernel, 1, NULL,
                                                    &global_item_size_cholesky, &local_item_size, 0, NULL, NULL);
            ASSERT_CL(ret, "Error after launching 2nd cholesky_kernel !");


            ret = clSetKernelArg(ckf_kernel2, 4, sizeof(idx), (void *)&idx);

            ret = clEnqueueNDRangeKernel(command_queue, ckf_kernel2, 1, NULL,
                                            &global_item_size_kernel2, &local_item_size, 0, NULL, NULL);
            ASSERT_CL(ret, "Error after launching ckf_kernel2 !");

现在，我想将此代码用于具有多个 GPU 的系统。所以我已经完成了以下步骤：

• 为所有 GPU 创建了一个上下文。
• 每个设备创建一个命令队列。
• 为每个设备创建了单独的内核（假设有两个 gpu，下面的代码片段）

• 为每个设备分配单独的设备缓冲区

  cl_kernel ckf_kernel1[2];
  cl_kernel ckf_kernel2[2];
  cl_kernel cholesky_kernel[2];
  cl_kernel memset_kernel[2];
  
  // read get kernel.
  ckf_kernel1[0] = clCreateKernel(program, "ckf_kernel1", &ret);
  ASSERT_CL(ret, "Cannot load ckf_kernel1[i]!");
  ckf_kernel2[0] = clCreateKernel(program, "ckf_kernel2", &ret);
  ASSERT_CL(ret, "Cannot load ckf_kernel2!");
  memset_kernel[0] = clCreateKernel(program, "memset_zero", &ret);
  ASSERT_CL(ret, "Cannot load memset_kernel!");
  cholesky_kernel[0] = clCreateKernel(program, "cholesky_kernel", &ret);
  ASSERT_CL(ret, "Cannot load cholesky_kernel!");
  
  ckf_kernel1[1] = clCreateKernel(program, "ckf_kernel1", &ret);
  ASSERT_CL(ret, "Cannot load ckf_kernel1[i]!");
  ckf_kernel2[1] = clCreateKernel(program, "ckf_kernel2", &ret);
  ASSERT_CL(ret, "Cannot load ckf_kernel2!");
  memset_kernel[1] = clCreateKernel(program, "memset_zero", &ret);
  ASSERT_CL(ret, "Cannot load memset_kernel!");
  cholesky_kernel[1] = clCreateKernel(program, "cholesky_kernel", &ret);
  ASSERT_CL(ret, "Cannot load cholesky_kernel!");
  

现在，我不确定如何将内核启动到循环内的不同设备上。如何让它们并行执行？请注意，上面的循环中有一个 clFinish 命令。

另一个问题：在主机上使用多个线程/进程是否是标准做法，其中每个线程/进程负责在单个 GPU 上启动内核？

Answer 1

1. 您无需为所有设备创建单独的上下文。如果它们来自不同的平台，您只需要这样做。
2. 您也不需要创建单独的内核。您可以同时为多个设备编译内核（clBuildProgram 支持多设备编译），如果您在设备上启动内核，运行时将知道内核实体是否持有对给定设备有效的设备二进制文件。
3. 最简单的事情是：创建一个上下文，获取您需要的所有设备，然后将其放入一个数组中，然后使用该数组来构建您的内核，并为其中的每个设备创建一个 command_queue。
4. clEnqueueNDRange 内核是非阻塞的。您的 for 循环没有通过的唯一原因是 clFinish() 语句，并且很可能是因为您正在使用顺序队列，这意味着单个设备情况下也可以在没有 clFinish 的情况下正常工作。

在 OpenCL 中最佳使用多 GPU 的一般想法是按照我提到的方式创建上下文内核队列，并使队列无序。如果命令没有未满足的依赖项，则允许并行执行命令，例如。command2的输入不是command1的输出，那么就可以自由地开始与command1并行执行了。但是，如果您使用此方法，则必须使用 clEnqueueNDRangeKernels 的最后几个参数，因为您必须使用 cl_events 构建此依赖关系链。每个 clEnqueueWhatever 都可以等待一系列事件，这些事件源自其他命令。队列中的命令只有在满足所有依赖项后才会开始执行。

有一个您没有涉及的问题，那就是缓冲区的概念。如果您想运行多 GPU，您需要为您的设备分别显式地创建缓冲区，并对您的数据进行分区。在 2 个设备上将相同的缓冲区设置为参数是无效的，而它们都在尝试写入它。充其量，运行时将序列化您的工作，并且两个设备不会并行工作。这是因为缓冲区是内存句柄，运行时负责将缓冲区的内容移动到需要它的设备。（这可能会隐式发生（延迟内存移动），或者如果您调用 clEnqueueMigrateBuffer 则显式发生。）禁止运行时同时向 2 个设备提供带有 CL_MEM_READ_WRITE 或 CL_MEM_WRITE_ONLY 标志的相同缓冲区。即使你知道作为程序员，2 个设备可能不会写入缓冲区的同一部分，但运行时不会。你必须告诉它。优雅的方法是创建 2 个子缓冲区，它们是较大/原始缓冲区的一部分；不太优雅的方法是简单地创建 2 个缓冲区。第一种方法更好，因为它更容易从多个设备收集数据回主机，因为你只需要获取大缓冲区，运行时就会知道哪些子缓冲区在哪些设备上被修改了，它会占用小心收集数据。

如果我看到您的 clSetKernelArgument 调用和您正在使用的缓冲区，我可以看到您的内核的依赖关系并写出您需要做什么，但我认为这对您在多设备运行方面是一个相当好的开始. 归根结底，一切都与数据有关。（并开始使用乱序队列，因为它有可能更快，并且它迫使您开始使用事件，这使您和任何阅读代码的人都可以明确知道哪些内核可以并行运行。