clBuildProgram 陷入嵌套循环答案

【问题标题】：clBuildProgram gets stuck with nested loopsclBuildProgram 陷入嵌套循环
【发布时间】：2014-02-25 19:15:14
【问题描述】：

在尝试这种 .cl 文件时，clBuildProgram 似乎卡住了，没有任何错误消息：

__local int bar(int a, int b, int c, int d, int e)
{
    return a*b*c*d; // 'e' not used
}

__kernel void foobar(__global int * notusedvariable)
{
    int foo=1;
    for (int a=1; a<=10; a++)
        for (int b=1; b<=10; b++)
            for (int c=1; c<=10; c++)
                for (int d=1; d<=10; d++)
                     for (int e=1; e<=10; e++)
                         foo *= bar(a,b,c,d,e);
}

当我删除最内层循环并将foo *= bar(a,b,c,d,e); 更改为foo *= bar(a,b,c,d,1); 时，它会编译。因此，存在某种过度优化或过度预计算。如果我有更多循环并且某些变量取自 get_global_id(...)，也会发生这种情况。

我能做什么？

我使用 Fedora Linux 20，并且已经安装

opencl-utils-0-12.svn16.fc20.x86_64
opencl-1.2-intel-cpu-3.2.1.16712-1.x86_64
opencl-utils-devel-0-12.svn16.fc20.x86_64
opencl-1.2-base-3.2.1.16712-1.x86_64

GPU 是 Geforce 210，即我能找到的最便宜的。

【问题讨论】：

标签： opencl

【解决方案1】：

它并没有真正“卡住”。它只是陷入了优化内核的地狱尝试中。主要是通过展开具有固定大小的循环（顺便说一句，通过发现根本没有使用 foo 变量！）

例如，当循环 a...d 被启用（并且 e 被关闭）时，为内核创建的二进制文件如下所示：

.entry foobar(
    .param .u32 .ptr .global .align 4 foobar_param_0
)
{
    .reg .pred  %p<4>;
    .reg .s32   %r<13>;


    mov.u32     %r10, 0;

BB0_1:
    add.s32     %r10, %r10, 1;
    mov.u32     %r11, 0;

BB0_2:
    mov.u32     %r12, 10;

BB0_3:
    add.s32     %r12, %r12, -2;
    setp.ne.s32     %p1, %r12, 0;
    @%p1 bra    BB0_3;

    add.s32     %r11, %r11, 1;
    setp.ne.s32     %p2, %r11, 10;
    @%p2 bra    BB0_2;

    setp.ne.s32     %p3, %r10, 10;
    @%p3 bra    BB0_1;

    ret;
}

你可以看到它实际上并没有计算任何东西 - 编译器已经很难发现实际上没有什么可做的。

将此与添加行时生成的输出进行比较

notusedvariable[0]=foo;

作为内核的最后一行：现在，计算可以不被跳过和优化掉。经过一段时间的编译，它产生了结果

.entry foobar(
    .param .u32 .ptr .global .align 4 foobar_param_0
)
{
    .reg .pred  %p<4>;
    .reg .s32   %r<80>;


    mov.u32     %r79, 1;
    mov.u32     %r73, 0;
    mov.u32     %r72, %r73;

BB0_1:
    add.s32     %r7, %r73, 1;
    add.s32     %r72, %r72, 2;
    mov.u32     %r76, 0;
    mov.u32     %r74, %r76;
    mov.u32     %r73, %r7;
    mov.u32     %r75, %r7;

BB0_2:
    mov.u32     %r9, %r75;
    add.s32     %r74, %r74, %r72;
    mov.u32     %r78, 10;
    mov.u32     %r77, 0;

BB0_3:
    add.s32     %r40, %r9, %r77;
    mul.lo.s32  %r41, %r40, %r79;
    mul.lo.s32  %r42, %r40, %r41;
    add.s32     %r43, %r74, %r77;
    mul.lo.s32  %r53, %r42, %r40;
    mul.lo.s32  %r54, %r53, %r40;
    mul.lo.s32  %r55, %r54, %r40;
    mul.lo.s32  %r56, %r55, %r40;
    mul.lo.s32  %r57, %r56, %r40;
    mul.lo.s32  %r58, %r57, %r40;
    mul.lo.s32  %r59, %r58, %r40;
    mul.lo.s32  %r60, %r59, %r40;
    mul.lo.s32  %r61, %r60, %r43;
    mul.lo.s32  %r62, %r61, %r43;
    mul.lo.s32  %r63, %r62, %r43;
    mul.lo.s32  %r64, %r63, %r43;
    mul.lo.s32  %r65, %r64, %r43;
    mul.lo.s32  %r66, %r65, %r43;
    mul.lo.s32  %r67, %r66, %r43;
    mul.lo.s32  %r68, %r67, %r43;
    mul.lo.s32  %r69, %r68, %r43;
    mul.lo.s32  %r70, %r69, %r43;
    mul.lo.s32  %r79, %r70, -180289536;
    add.s32     %r77, %r77, %r74;
    add.s32     %r78, %r78, -2;
    setp.ne.s32     %p1, %r78, 0;
    @%p1 bra    BB0_3;

    add.s32     %r76, %r76, 1;
    add.s32     %r30, %r9, %r7;
    setp.ne.s32     %p2, %r76, 10;
    mov.u32     %r75, %r30;
    @%p2 bra    BB0_2;

    setp.ne.s32     %p3, %r7, 10;
    @%p3 bra    BB0_1;

    ld.param.u32    %r71, [foobar_param_0];
    st.global.u32   [%r71], %r79;
    ret;
}

显然，它已经展开了一些循环，现在他无法再优化它们了。我假设当循环“e”也被激活时，这种展开（或优化未使用的循环）所需的时间至少呈二次方增加。所以如果你给他几个小时，他实际上可能也会完成编译......

正如 Tom Fenech 在 https://stackoverflow.com/a/22011454 中已经说过的那样，可以通过将 -cl-opt-disable 传递给 clBuildProgram 来缓解这个问题。

或者，您可以选择性地关闭每个循环的展开优化：当您插入时

#pragma unroll 1

直接在 for 循环之前，您实际上禁用了此特定循环的展开。

重要不要盲目地插入带有任意值的unroll pragma。使用1 是安全的，但对于其他值，您必须手动确保它不会影响程序的正确性。请参阅 CUDA 编程指南，“B.21.#pragma unroll”部分。

在这种情况下，在最里面的两个循环（d 和 e）之前插入 #pragma unroll 1 似乎就足够了，以便进行足够的优化以快速构建程序。

编辑：sigh 修剪快了 4 分钟... :-(

【讨论】：

谢谢，我会将此标记为已接受。我想我也可以将“10”作为参数传递给内核；那么编译器就不能（过度）优化代码。

【解决方案2】：

你正在做的计算确实会得到一个非常大的数字！

我在硬件（NVIDIA GTX480）上遇到了与您相同的问题，但我认为这与硬件无关。您只是在生成一个在编译时已知的数字，并且对于 int 变量类型来说太大而无法预先计算。我将int 更改为long，程序现在可以构建。

编辑

我刚刚尝试过，使用的是英特尔平台。它编译得很好。您还可以通过将开关 -cl-opt-disable 传递给 clBuildProgram 使其在 NVIDIA 上运行。这会禁用 all 优化 - 您可能会对其他一些编译器开关感到幸运。有关详细信息，请参阅clBuildProgram reference。

【讨论】：

但这也发生在 unsigned int 上。他们应该定义溢出。（实际上我想搜索函数的零点。通过首先尝试使用溢出算法，我可以找到候选者来检查更长的整数。）
@JoriMäntysalo 我已经用更多细节更新了我的答案。

【解决方案3】：

Nvidia OpenCL 编译器可能正在执行循环展开。如果它对你的每个嵌套循环都这样做，那将导致生成大量代码！

有 cl_nv_pragma_unroll Nvidia 特定的 OpenCL 扩展可用于更好地控制循环展开。至引用其文档：

用户可以指定展开源程序中的循环。这个是通过编译指示完成的。该pragma的语法如下

#pragma unroll [展开因子]

pragma unroll 可以选择指定一个展开因子。语用必须紧接在循环之前放置，并且仅适用于该循环循环。

如果未指定展开因子，则编译器将尝试执行完整或完全展开循环。如果循环展开因子是指定编译器将执行部分循环展开。循环因子，如果指定，必须是编译时非负整数常数。

循环展开因子为 1 表示编译器不应展开循环。

如果循环不是编译时可计算的。

默认情况下，听起来它会展开某些下限的最大限制（例如，在您的示例中为 10），但如果它们是嵌套的，它可能仍会展开（展开检查逻辑可能不够复杂，无法检查嵌套循环） .

您可以尝试#pragma unroll 1 禁用展开：

int foo=1;
#pragma unroll 1
for (int a=1; a<=10; a++)
    #pragma unroll 1
    for (int b=1; b<=10; b++)
        #pragma unroll 1
        for (int c=1; c<=10; c++)
           #pragma unroll 1
           for (int d=1; d<=10; d++)
              #pragma unroll 1
              for (int e=1; e<=10; e++)
                 foo *= bar(a,b,c,d,e);

您还需要通过以下方式启用扩展：

#pragma OPENCL EXTENSION cl_nv_pragma_unroll : enable

也在您的 OpenCL 源文件的顶部。

【讨论】：