使用多个 GPU OpenCL

问题描述 投票：0回答：2

我有一个循环，在其中将多个内核启动到 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 上启动内核）是标准做法吗？

opencl gpgpu
2个回答
5
投票
  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 个缓冲区。第一种方法更好，因为它更容易从多个设备收集数据回主机，因为您只需要获取大缓冲区，并且运行时会知道哪些子缓冲区在哪些设备上被修改，并且会花费小心收集数据。

如果我看到您的 clSetKernelArgument 调用以及您正在使用的缓冲区，我可以看到您的内核的依赖关系并写出您需要执行的操作，但我认为这对于您获得多设备来说是一个相当好的开始跑步。最终，一切都与数据有关。 （并开始使用无序队列，因为它有可能更快，并且它迫使您开始使用事件，这使您和任何阅读代码的人都清楚地知道哪些内核可以并行运行。

0
投票

@流星头 为什么要在writebuffer后面加上clfinish，这样会导致内核以PCIE数据传输串行运行。

最新问题
© www.soinside.com 2019 - 2024. All rights reserved.