OpenCL

Question

当边界是动态的时，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;
}

Answer 1

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

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

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

顺便说一句：我没有看到 N 变量赋值的位置，也没有看到元素数组的任何使用。

Answer 2

假设是opencl 1.2设备，

如果每个 "predict_future_events" 在性能方面都很混乱，您可以检查一些 "hardware optimized" 更改。你可以同时抛出 2 个不同的内核（两个不同的全内核（N），如果它们可以是 separated/independent）或者你可以将一半的内核（N/2）推为 "constant version"和另一半在不同的内核中（因为这在计算上与您的第一个示例没有什么不同），也许驱动程序可以处理一些情况，其中一个内核永远延迟但至少另一半获得计算资源（如果驱动程序可以做到这一点）。因此，更多的管道将忙于做一些事情，并最终为内核提供更好的时序。

除此之外，每个函数都有一个随机延迟使得很难预见循环中的哪一组函数给出了总延迟，因此给所有线程相等的步数（如第一个example/constant）是更容易 "assume" 将有更大的机会在线程之间实现负载平衡。

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

对于 opencl 2.0 设备，您可以在内核线程中生成更多 threads/groups，这应该会使其更加动态。

OpenCL – 执行动态 for 循环

OpenCL – execution of dynamic for loops

parallel-processing

loops

gpgpu

cpu-architecture