OpenCL 非发散 if() 后跟 printf() 导致大幅减速和内核代码干扰最大工作组大小
OpenCL non-diverging if() followed by printf() causing massive slowdown and kernel code interfering on maximum workgroup size
我是 OpenCL 的新手，我在几个地方读到应该避免 if/else 结构，主要是因为当线程的评估不同（发散分支）时，速度会显着降低。
尽管如此，我使用了一个 if(cond) 后跟一些打印来保证在满足禁止条件时我可以调试导致它的原因。关键是，只要单个线程的 if() 为真，我就会终止进程，因此我不担心不同的线程会以不同的方式评估条件。
然而，我发现即使所有线程都在这个 if() 中评估为 false，与不使用 if() 相比，速度也会大大降低——我注释掉了条件语句及其主体语句以验证.
观察： 我有一个内核（128 个工作组，每个工作组有 128 个工作项）调用函数 foo()，并且 if/prints 在 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
}
}
}
相关信息：
- 使用我展示的代码，整个程序（主机 + 内核 + foo）大约需要 90 秒 到 运行
- 如果我注释掉整个 if() 语句（条件 + 打印），程序大约需要 2（两）秒 到 运行
- 如果我删除所有打印件并在 if() 上使用单个虚拟分配（例如 z=0），程序大约需要 2 秒 到 运行
- 如果我删除所有打印件并在 if() 上使用涉及 get_local_id(0)（例如 z=get_local_id(0)）的单个赋值，程序会占用2 秒 到 运行
- 如果我删除所有打印件并使用涉及 get_local_id(0) 的单个分配并在之后打印它，程序大约需要 7 秒 到 运行
- 通过从原始代码中删除所有 printf() 语句并一次添加一个（仅打印第一个，然后是第 1 和第 2，然后从第 1 到第 3，...），我发现了到第四次打印 运行ning 时间 大约是 7 秒 ，但是在添加第 5 次打印之后 运行ning 时间 跳到 90 秒
另外一个相关资料，在之前的测试中发现：我在主机代码中打印最大工作组大小，以了解不同平台的差异。在运行ning时间小的情况下（case 2-5，最多7秒），host输出最大WG size为8192，但是在运行ning时间长的情况下，它输出的最大 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
显示工作组最大规模的主机代码是：
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() 代码从未执行时，为什么代码中的这种微小修改会导致如此巨大的速度下降？
- 我的代码如何影响最大 WG 大小，即设备的 属性？
看来你的调试有问题。
我知道使用 ARM mali 你需要添加你的 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 并尝试找出你是否不需要一些特殊的扩展 ;))
从 OpenCL 内核实现 printf 的方式必须总是包含某种缓冲区，用于将写入的文本发送到主机。但是工作项之间必须有协调，这样它们就不会写入缓冲区的同一部分。这意味着使用本地内存、原子等。编译器无法静态确定您的 printf 永远不会被调用，因此它仍然必须为所有可能的调试输出分配足够的内存 returned。如果协调涉及障碍，则这些障碍也需要插入 else 路径中，因为组中的所有工作项都必须步调一致地到达障碍。由于您的最大工作组大小根据您是否有打印语句而变化，这很好地表明您的平台的 printf 实现在某种程度上使用了本地内存，这通常意味着它使用了障碍。因此，else 路径上的无形障碍可能会严重拖慢您的速度。
因此，从本质上讲，OpenCL 内核的通用调试输出总是会减慢速度。你不应该在你正在分析或部署的内核中有调试输出，因为它们会破坏你可能试图实现的任何优化。
如果您正在尝试调试或开发内核，并且发现必要的调试输出正在减慢速度，以至于干扰了您的工作，您可能想要完全摆脱 printf 并且改用输出结构化诊断：一个带有 space 的全局缓冲区，用于您希望从工作项中 return 的所有不同诊断值。因此，也许每个 returned 值都有一个数组结构，每个工作项一个数组条目。然后在主机端将其转换为人类可读的形式（即 printf）。这样做的好处是：
printf 本身实际上是一个相当复杂的函数，特别是对于 GPU，因此它本身会减慢速度。即使文本格式化本身是在主机上完成的，只有数据参数被设备写入共享缓冲区，对该缓冲区的访问也需要以某种方式同步，除非缓冲区所需的大小可以而且是由编译器静态证明，也需要调整大小或刷新操作或类似操作。- 自己管理诊断输出内存缓冲区意味着没有 运行 输出缓冲区大小限制、重新分配等
- 如果没有两个工作项写入相同的诊断值内存单元，则不需要任何同步或原子操作。
- 将诊断输出简单地存储为浮点数组、整数数组等内存效率更高，因此这不会对 GPU 内存带宽造成太大压力。
我是 OpenCL 的新手，我在几个地方读到应该避免 if/else 结构，主要是因为当线程的评估不同（发散分支）时，速度会显着降低。
尽管如此，我使用了一个 if(cond) 后跟一些打印来保证在满足禁止条件时我可以调试导致它的原因。关键是，只要单个线程的 if() 为真，我就会终止进程，因此我不担心不同的线程会以不同的方式评估条件。
然而，我发现即使所有线程都在这个 if() 中评估为 false，与不使用 if() 相比，速度也会大大降低——我注释掉了条件语句及其主体语句以验证.
观察： 我有一个内核（128 个工作组，每个工作组有 128 个工作项）调用函数 foo()，并且 if/prints 在 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
}
}
}
相关信息：
- 使用我展示的代码，整个程序（主机 + 内核 + foo）大约需要 90 秒 到 运行
- 如果我注释掉整个 if() 语句（条件 + 打印），程序大约需要 2（两）秒 到 运行
- 如果我删除所有打印件并在 if() 上使用单个虚拟分配（例如 z=0），程序大约需要 2 秒 到 运行
- 如果我删除所有打印件并在 if() 上使用涉及 get_local_id(0)（例如 z=get_local_id(0)）的单个赋值，程序会占用2 秒 到 运行
- 如果我删除所有打印件并使用涉及 get_local_id(0) 的单个分配并在之后打印它，程序大约需要 7 秒 到 运行
- 通过从原始代码中删除所有 printf() 语句并一次添加一个（仅打印第一个，然后是第 1 和第 2，然后从第 1 到第 3，...），我发现了到第四次打印 运行ning 时间 大约是 7 秒 ，但是在添加第 5 次打印之后 运行ning 时间 跳到 90 秒
另外一个相关资料，在之前的测试中发现：我在主机代码中打印最大工作组大小，以了解不同平台的差异。在运行ning时间小的情况下（case 2-5，最多7秒），host输出最大WG size为8192，但是在运行ning时间长的情况下，它输出的最大 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
显示工作组最大规模的主机代码是：
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() 代码从未执行时，为什么代码中的这种微小修改会导致如此巨大的速度下降？
- 我的代码如何影响最大 WG 大小，即设备的 属性？
看来你的调试有问题。 我知道使用 ARM mali 你需要添加你的 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 并尝试找出你是否不需要一些特殊的扩展 ;))
从 OpenCL 内核实现 printf 的方式必须总是包含某种缓冲区，用于将写入的文本发送到主机。但是工作项之间必须有协调，这样它们就不会写入缓冲区的同一部分。这意味着使用本地内存、原子等。编译器无法静态确定您的 printf 永远不会被调用，因此它仍然必须为所有可能的调试输出分配足够的内存 returned。如果协调涉及障碍，则这些障碍也需要插入 else 路径中，因为组中的所有工作项都必须步调一致地到达障碍。由于您的最大工作组大小根据您是否有打印语句而变化，这很好地表明您的平台的 printf 实现在某种程度上使用了本地内存，这通常意味着它使用了障碍。因此，else 路径上的无形障碍可能会严重拖慢您的速度。
因此，从本质上讲，OpenCL 内核的通用调试输出总是会减慢速度。你不应该在你正在分析或部署的内核中有调试输出，因为它们会破坏你可能试图实现的任何优化。
如果您正在尝试调试或开发内核，并且发现必要的调试输出正在减慢速度，以至于干扰了您的工作，您可能想要完全摆脱 printf 并且改用输出结构化诊断：一个带有 space 的全局缓冲区，用于您希望从工作项中 return 的所有不同诊断值。因此，也许每个 returned 值都有一个数组结构，每个工作项一个数组条目。然后在主机端将其转换为人类可读的形式（即 printf）。这样做的好处是：
printf本身实际上是一个相当复杂的函数，特别是对于 GPU，因此它本身会减慢速度。即使文本格式化本身是在主机上完成的，只有数据参数被设备写入共享缓冲区，对该缓冲区的访问也需要以某种方式同步，除非缓冲区所需的大小可以而且是由编译器静态证明，也需要调整大小或刷新操作或类似操作。- 自己管理诊断输出内存缓冲区意味着没有 运行 输出缓冲区大小限制、重新分配等
- 如果没有两个工作项写入相同的诊断值内存单元，则不需要任何同步或原子操作。
- 将诊断输出简单地存储为浮点数组、整数数组等内存效率更高，因此这不会对 GPU 内存带宽造成太大压力。