OpenCL不发散if()，后面跟着printf()，导致大量减速和内核代码干扰最大工作组大小

Question

我是OpenCL的新手，我在几个地方读到过应该避免if/else结构的文章，主要是因为当线程的计算不同时(发散分支)，就会出现明显的减速。

尽管如此，我还是使用了一个if(cond)，后面跟着一些打印，以保证当一个被禁止的条件被满足时，我可以调试导致它的原因。关键是，每当if()对于单个线程为真时，我就会终止进程，因此我并不担心不同的线程对条件进行不同的评估。

然而，我发现，即使在这个if()中所有线程都计算false时，与不使用if()相比，还是有一个大幅度的减速--我注释掉了条件语句和它的主体语句来验证。

一个观察：I有一个内核(128个工作组，每个工作组有128个工作项)，它调用函数foo()，if/位于foo()内。foo()的相关部分如下：

foo(){
     bool leftCorrect, rightCorrect, topCorrect, bottomCorrect, topLeftCorrect, topRightCorrect, bottomLeftCorrect, bottomRightCorrect;
     for(i=0;i<11;i++){
          for(j=0;j<11;j++){
               // Initial assignments. Using select() to avoid branch divergence
               leftCorrect = select(true,false,condition); 
               rightCorrect = select(true,false,condition); 
               topCorrect = select(true,false,condition); 
               bottomCorrect = select(true,false, condition);
            
               // Use the previous variables to compute the other bool variables and update them
               // Some boolean operations ...
           
               // This if basically tests if more than one is true
               if(leftCorrect + rightCorrect + topCorrect + bottomCorrect + topLeftCorrect + topRightCorrect + bottomLeftCorrect + bottomRightCorrect > 1){
                    printf("@@@@@\nFATAL ERROR: Multiple corrections in gid=%d\n",get_local_id(0));
                    printf("L  %d\n", leftCorrect);
                    printf("R  %d\n", rightCorrect);
                    printf("T  %d\n", topCorrect);
                    printf("B  %d\n", bottomCorrect);

                    printf("TL %d\n", topLeftCorrect);
                    printf("TR %d\n", topRightCorrect);
                    printf("BL %d\n", bottomLeftCorrect);
                    printf("BR %d\n", bottomRightCorrect);
              } 
              // Do something with the boolean variables and select() statements before return
         }
    }
}

相关信息：

1. 在我展示的代码中，如果注释掉整个if()语句(条件+打印)，整个程序(主主机+内核+ foo)需要大约90秒的来运行
2. ，如果在if()上删除所有打印并使用单个虚拟赋值(例如z=0)，则程序大约需要2 (2秒)来运行

< code >H 114，如果我删除所有打印并在if()上使用涉及到(0)(例如z=get_local_id(0))的单个赋值，则程序需要大约2秒的时间来运行

1. ，如果删除所有打印并使用涉及get_local_id(0)的单个赋值并在之后打印，则程序需要2秒运行
2. ，通过从原始代码中删除所有printf()语句并一次添加一个语句(只有第一次打印，然后是第一次和第二次，然后从第一次到第三次，.)，程序需要大约7秒的时间来运行seconds

• ，我发现到第四次打印，运行时间大约是7秒，但是在添加第5次打印之后，运行时间跳转到90

另一个相关信息，在前面的测试中发现：我正在打印主机代码中的最大工作组大小，以了解不同平台上的差异。在运行时间较短的情况下(情况2-5，最多7秒)，主机输出最大WG大小为8192。但是，在长时间运行的情况下，它输出的最大WG大小为4096。

clinfo命令向我展示了以下信息：

Platform Name                                   Intel(R) CPU Runtime for OpenCL(TM) 
Device Name                                     Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz
Max work item sizes                             8192x8192x8192
Max work group size                             8192

显示最大WG大小的主机代码是：

size_t size_ret;
cl_uint maximum_size;
error = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, 
size_ret, &maximum_size, NULL);
cout << "-- Maximum WG size " << maximum_size << endl;

所有这些信息归纳在以下两个问题中：

为什么代码中的这种微小修改会导致如此大规模的减速，如果if()代码从来不是executed?

• How，那么我的代码会干扰最大WG大小，这是设备的一个属性吗？

Answer 1

从printf内核实现OpenCL的方式必须总是包含某种缓冲区，以便将写好的文本发送到主机。但是必须在工作项之间进行协调，这样它们才不会写入缓冲区的同一部分。这意味着使用本地内存、atomics等。编译器不能静态地确定您的printf永远不会被调用，所以它仍然必须为所有可能的调试输出分配足够的内存以返回。如果协调涉及障碍，则还需要将这些障碍插入到else路径上，因为组中的所有工作项都必须同步到达障碍。随着最大工作组大小的变化取决于是否有打印语句，这很好地表明您的平台的printf实现在某种程度上使用了本地内存，这通常意味着它使用了障碍。因此，在else路径上的无形障碍可能是对你来说慢下来的原因。

因此，从本质上说，来自OpenCL内核的通用调试输出总是会减慢速度。您不应该在您正在分析或部署的内核中使用调试输出，因为它们会破坏您可能试图实现的任何优化。

如果您正在尝试调试或开发内核，并且发现必要的调试输出会使您的工作慢得太慢，那么您可能希望完全摆脱printf，使用输出结构化诊断:一个全局缓冲区，为您希望从工作项返回的所有不同的诊断值提供空间。因此，可能是一个包含每个返回值的数组的结构，每个工作项都有一个数组条目。然后在主机端将其转换为人类可读的形式(即printf)。这样做的好处是：

• printf本身实际上是一个相当复杂的函数，特别是对于GPU来说，所以它本身会减慢速度。即使文本格式本身是在主机上完成的，只有由设备写入共享缓冲区的数据参数，对该缓冲区的访问也需要以某种方式进行同步，除非所需缓冲区的大小可以并经编译器静态验证，否则也需要调整大小或刷新操作或类似操作。
• 自己管理诊断输出内存缓冲区意味着没有运行到输出缓冲区大小限制、重新分配等。如果没有两个工作项，则
• 写入相同的诊断值内存单元，您不需要任何同步或atomics.
• Simply将诊断输出存储为浮点数、整数等数组，这样内存效率更高，因此不会对GPU内存带宽造成太大压力。

Answer 2

看来您在调试方面遇到了问题。现在，使用ARM马里，你需要加入CCP

        /* Enable a printf callback function for this context. */
        CL_PRINTF_CALLBACK_ARM,   (cl_context_properties) printf_callback,

        /* Request a minimum printf buffer size of 4MiB for devices in the
         context that support this extension. */
        CL_PRINTF_BUFFERSIZE_ARM, (cl_context_properties) 0x1000,

在你的内核里

        pragma OPENCL EXTENSION cl_arm_printf : enable

因此，请检查GPU，并尝试查找您是否不需要一些特殊扩展；)