OpenCL – 执行动态 for 循环

Posted 2023-02-16

【中文标题】OpenCL – 执行动态 for 循环

【英文标题】：OpenCL – execution of dynamic for loops

【发布时间】：2016-07-31 17:29:14

【问题描述】：

当 OpenCL 内核中的 for 循环边界是动态的时（即对于每个工作项，for 循环执行不同的次数时），它们如何在设备上执行？

---

AFAIK，内核是一组（或者更好地说是一个流）指令。 GPU 设备是一组独立的计算单元（流式多处理器 - SM），每个计算单元都包含多个计算单元（流式处理器 - SP）。

每个 SM 可以从内核（即指令流）加载一条指令（对于不同的 SM，这可能是不同的指令），并为当前 SM 中的 SP（每个 SP 运行相同的指令，但数据不同——SIMD）。

一个 SM 中的所有 SP 必须运行相同的指令，因此在执行 for 循环的条件后，必须根据每个工作项的条件结果做出动态决定，下一条指令将是什么可以在 SM 上运行和来运行它的工作项。

基于这个假设，我假设foobaz 内核（见下文）会执行得更快，因为当一个工作项完成执行时，另一个工作项可以取代它。

这个假设是错误的吗？

---

以下两个内核，foobar 和 foobaz，哪一个最终会执行得更快？ 性能取决于什么？ （一个元素的属性数量可能比其他元素多几个数量级）。

foobar;

__kernel void foobar(__global int* elements,            /* size N          */
                     __global int* element_properties,  /* size N*constant */
                     __global int* output)              /* size N          */

    size_t gid  = get_global_id(0);
    int reduced = 0;

    for (size_t i=N*gid; i<N+N*gid; i++)
      reduce += predict_future_events( reduce, element_properties[i] );


    output[gid] = reduced;

…和foobaz;

__kernel void foobaz(__global int*  elements,                   /* size N       */
                     __global int*  element_properties,         /* size upper-bounded */
                     __global int2* element_properties_ranges,  /* size N       */
                     __global int*  output)                     /* size N       */

    size_t gid  = get_global_id(0);
    int reduced = 0;

    // `range.x` = starting index in `element_properties`
    // `range.y` = ending   index in `element_properties`
    int2 range = element_properties_ranges[gid]; 

    for (size_t i=range.x; i<range.y; i++)
      reduce += predict_future_events( reduce, element_properties[i] );


    output[gid] = reduced;

【参考方案1】：

假设是opencl 1.2设备，

如果每个“predict_future_events”在性能方面都很混乱，您可以检查一些“硬件优化”更改。您可以同时抛出 2 个不同的内核（两个不同的完整内核（N），如果它们可以分开/独立），或者您可以将内核的一半（N/2）作为“恒定版本”推送，另一半推送为不同的内核（因为这在计算上与您的第一个示例没有什么不同），也许驱动程序可以处理一些情况，即一个内核永远延迟但至少另一半获得计算资源（如果驱动程序可以这样做）。因此，更多的管道将忙于做某事并最终为内核提供更好的时机。

除此之外，每个函数都有一个随机延迟，因此很难预测循环中的哪组函数给出了多少总延迟，因此给所有线程提供相同数量的步骤（如第一个示例/常量）更容易“假设" 将有更大的机会在线程之间实现负载平衡。

例如，光线追踪内核的 1000 深度折射 + 1000 深度反射会非常混乱，因此您可以只为每个线程提供 1 条光线来计算，因为您无法知道光线是否会被折射或反射在下一个表面上（如果有的话）。也许更近的分组可以更频繁地使用 L1-L2 缓存。

对于 opencl 2.0 设备，您可以在内核线程中生成更多线程/组，这将使这更加动态。

【参考方案2】：

您在这两种解决方案中所做的几乎相同。我敢打赌，他们必须几乎同时完成。

如果您想更快，请使用 SIMD 功能，使用 int4 作为参数并减少变量，然后调整您的 predict_future_events 函数以处理 int4 值。通过这种方式，您可以获得高达 4 倍的性能，因为每条指令并行处理 4 个元素。

根据您的硬件，您可以使用 int8 或最高 int16。

顺便说一句：我没有看到 N 变量在哪里分配，也没有看到任何元素数组的使用。

【讨论】：

我可能应该声明两个内核都只是伪代码（所以绝对加速不是我关心的问题）。然而，核心部分是，每个线程的 data-dependent 循环次数 与 每个线程的恒定（但大量）迭代如何影响性能，以及我对硬件（GPU）如何处理 for 循环的假设是否正确。

什么部分是动态的？ range.y?

@huseyintugrulbuyukisik 是的。好吧，实际上我事先不知道任何粒子的range.x 和range.y，所以这两个值都是动态的（我希望我们所说的“动态”是指同一件事）。假设element_properties_ranges 中的值是在执行foobar 和foobaz 之前由另一个内核生成的。

@sarasvati 所以你问迭代负载平衡是否比单个内核调用更快？

@huseyintugrulbuyukisik 我不明白other trying to change its range。上下文：首先我运行一个在element_properties_ranges 中设置值的内核。 （element_properties_ranges 中的每个int2 值代表element_properties 中elements 的适当元素的起始索引。）。然后我运行foobaz。