为什么会有 CL_DEVICE_MAX_WORK_GROUP_SIZE？

我试图了解 OpenCL 设备（例如 GPU）的体系结构，但我不明白为什么本地工作组中的工作项数量有明确的限制，即常量 CL_DEVICE_MAX_WORK_GROUP_SIZE。

在我看来，这应该由编译器处理，即，如果（为简单起见，一维）内核以本地工作组大小 500 执行，而其物理最大值为 100，并且内核看起来如下所示：

__kernel void test(float* input) {
    i = get_global_id(0);
    someCode(i);
    barrier();
    moreCode(i);
    barrier();
    finalCode(i);
}

那么它可以自动转换为该内核上工作组大小为 100 的执行：

__kernel void test(float* input) {
    i = get_global_id(0);
    someCode(5*i);
    someCode(5*i+1);
    someCode(5*i+2);
    someCode(5*i+3);
    someCode(5*i+4);
    barrier();
    moreCode(5*i);
    moreCode(5*i+1);
    moreCode(5*i+2);
    moreCode(5*i+3);
    moreCode(5*i+4);
    barrier();
    finalCode(5*i);
    finalCode(5*i+1);
    finalCode(5*i+2);
    finalCode(5*i+3);
    finalCode(5*i+4);
}

然而，这似乎不是默认情况下完成的。为什么不？有没有办法使这个过程自动化（除了我自己编写预编译器之外）？或者是否存在一个内在的问题，可能使我的方法在某些示例上失败（您能给我一个）吗？

我认为 CL_DEVICE_MAX_WORK_GROUP_SIZE 的起源在于底层硬件实现。

多个线程同时在计算单元上运行，每个线程都需要保持状态（用于调用、跳转等）。大多数实现为此使用堆栈，如果您查看 AMD Evergreen 系列，它们是可用堆栈条目数量的硬件限制（每个堆栈条目都有子条目）。这本质上限制了每个计算单元可以同时处理的线程数量。

至于编译器可以做到这一点，使其成为可能。它可以工作，但要明白这意味着再次重新编译内核。这并不总是可能的。我可以想象这样的情况：开发人员以二进制格式转储每个平台的编译内核，并将其与他们的软件一起发布，只是出于“不那么开源”的原因。