opencl中的并行减和实现答案

【问题标题】：parallel sum reduction implementation in openclopencl中的并行减和实现
【发布时间】：2015-08-01 07:34:42
【问题描述】：

我正在浏览link提供的NVIDIA示例代码

在示例内核代码（文件oclReduction_kernel.c）中，reduce4 使用了

1) 展开和移除线程 id

2) 除此之外，代码使用块大小检查对本地内存中的数据求和。我认为在 OpenCL 中我们有 get_local_size(0/1) 来了解工作组的规模。块大小让我很困惑。

我无法理解上述两点。这些东西为什么以及如何帮助优化？任何关于 reduce5 和 reduce6 的解释都会有所帮助。

【问题讨论】：

docs.nvidia.com/cuda/samples/6_Advanced/reduction/doc/… 的演示文稿指的是 CUDA（诚然，我认为 OpenCL 示例基本上是通过尽可能将 CUDA 内核转换为 OpenCL 创建的），但它可能会带来一些见解，特别是关于进一步优化的 reduce5/reduce6 版本。
是的，他们只在 cuda 中。但它不回答任何问题
谢谢@Marco13 我在晚上错过了你的链接。可能是我困了:)

【解决方案1】：

@Marco13 在 cmets 中链接的 https://docs.nvidia.com/cuda/samples/6_Advanced/reduction/doc/reduction.pdf 的幻灯片 21 和 22 中已经对此进行了详细解释。

随着减少的进行，#“活动”线程减少

当 s

指令在 warp 中是 SIMD 同步的。

这意味着当 s

我们不需要 __syncthreads()

我们不需要“if (tid

在不展开的情况下，所有 warp 都会执行 for 循环的每次迭代和 if 语句

代码实际上是以 32 个线程为一组执行的，NVIDIA 是什么调用warp。

每个核心都可以执行一个顺序线程，但核心在 NVIDIA 称之为 SIMT（单指令多线程）方式；同一组中的所有内核同时执行相同的指令时间，很像经典的 SIMD 处理器。

Re 2) blockSize 看起来工作组的大小。

【讨论】：

当我们说代码是为 GPU 编译的 SIMD-4 时。这么说对吗:: 将有四个 SIMD 内核执行相同的指令，指令可以像 float/float4/float16 一样？？