如何优化我的 OpenCl 内核

Question

我已经分离出一些彼此不同但需要在不同工作项中并行运行的功能。因此，当内核被调用时，它需要决定必须执行哪个函数..

void call_calc0() {
    // code
}

void call_calc1() {
    // code
}

void call_calc2() {
    // code
}

void call_calc3() {
    // code 
}

__kernel void perform (__global double* A, __global double* B) {
    int idx = get_global_id(0);
    if (idx == 0) {
        call_calc0();
    } else if (idx == 1) {
        call_calc1();
    } else if (idx == 2) {
        call_calc2();
    } else if (idx == 3) {
        call_calc3();
    }
}

如果有 256/512 个工作项，此代码示例将不是正确的方法。我该如何优化呢？

Answer 1

如果可能的话，最好的优化是使用四个不同的内核。如果您调用这个内核的组大小超过 1，那么在并行执行时就会出现问题。

如果可能，请尝试分离您的全局内存或以非常小心、不冲突的方式使用它。这应该允许您创建四个单独的内核，并摆脱条件代码执行。

当遇到第一个 if/case 时，该组的一些工作项将运行代码，但其他 75% 的工作项将等待。大多数 opencl 设备，尤其是 GPU，都以这种方式运行。当前 25% 的工作项完成后，它们将等待下一个 if/case 代码执行。

这适用于 opencl 中的所有分支，例如 if/else、switch、for 和 while/do。每当组中的某些工作项不满足条件时，它们就会等待其他满足条件的工作项。然后在 'if' 组等待时执行 'else' 组工作项。

查看它的另一种方法是比较 CPU 和 GPU 硬件。 CPU 有很多专门用于分支预测和高速缓存的晶体管。 GPU 在本质上更加基于矢量，并且直到最近才开始支持 CPU 的一些更高级的流控制功能。

Answer 2

由于 OpenCL 不支持函数指针，因此您只能使用 if/else 或 switch。这两者的性能应该是相同的，只是编码偏好不同。

您可以使用预处理器宏让事情变得更简单/更简洁。例如，您可以这样做：

#define CALL_CASE(i)    \
    case i:             \
        call_calc##i(); \
        break;          \

__kernel void perform (__global double* A, __global double* B) {
    int idx = get_global_id(0);
    switch (idx) {
        CALL_CASE(0);
        CALL_CASE(1);
        CALL_CASE(2);
        CALL_CASE(3);
        ... // etc
    }
}

如果您要自动生成call_calcX() 函数，那么同时生成这个switch 块也很容易。如果您手动编写这些call_calcX() 函数，那么只需多行代码即可将每个函数添加到块中。不理想，但也不可怕。

根据上面的 cmets，这个问题似乎根本不是数据并行的，这将限制利用大多数 OpenCL 设备中可用的 SIMD 执行的能力。