问题标签 [reduction]

我是 OpenCl 的新手。

我需要对一维双精度数组进行归约（求和运算符）。

我一直在网上徘徊，但我发现的例子很混乱。任何人都可以发布一个易于阅读（并且可能有效）的教程实现吗？

附加信息： - 我可以使用一台 GPU 设备；- 我使用 C 作为内核代码

以下代码来自amd官网

我对其进行了调整以使其作为总和减少工作：

当我只使用一个工作组时，它就像一种魅力（意思是我给予NULL）local_work_size，clEnqueueNDRangeKernel()但是当我尝试改变工作组维度时，事情就超出了我的控制。（我应该说我是 OpenCl 的新手）

我所做的如下

我做错了吗？

此外，我注意到如果我设置#define GLOBAL_DIM 60000（这是我需要的）我会用完本地内存。如果我使用多个工作组，我是否会获得“更多”本地内存，或者本地内存在工作组之间均匀分布？

我想做并行减少，但在我的内核内部，数据在共享内存中。推力库可以做到这一点吗？就像是

但这在内核中不起作用。可能吗？谢谢你。

我试图用这段代码对一个数组求和，但我被卡住了。我可能需要一些“CUDA for dummies tutorial”，因为我花了很多时间在这样的基本操作上，但我无法让它工作。

以下是我不明白或不确定的事情的清单：

1. 我应该使用多少块（dimGrid）？我认为应该是N/dimBlock.x/2（N=输入数组的长度），因为在内核开始时，数据被加载并从全局内存的两个“块”添加到共享内存

2. 在原始代码中有blockSize. 我将其替换为blockDim.x因为我不知道这些变量有何不同。但是当blockSize=时blockDim.x，gridSize = blockDim.x*2*gridDim.x对我来说没有意义 -gridSize将大于 N。在一维数组的上下文中 *Dim.x 和 *Size 有什么区别？

3. 主要逻辑 - 在内核中，每个块总和 2*dimBlock（块中的线程）个数。当 N = 262144 和 dimBlock = 128 时，内核返回 1024 个部分和数组。然后我再次运行内核，结果 = 4 部分和。最后，在最后一次运行中，返回单个和，因为数组是由单个块处理的。

4. 我对二进制数组求和。在第一次运行中，我可以uchar4用于输入数据。在第二次和第三次运行中，我将使用int.

请告诉我我错过了什么

谢谢

对于下面的代码，我得到一个错误：unrecognized #pragma: #pragma omp reduction (+: sum)。请注意，函数内部的 for 循环不是并行循环，因为函数本身已经并行化。你能说问题在哪里吗？

主cpp文件：

另一个 cpp 文件，其中定义了函数

我正在使用 QR 码条形码将 UUID 存储在我的系统中，我需要检查生成的条形码是否是我的，而不是其他人的。我还需要保持编码数据简短，以便 QR 码保持在较低版本范围内并易于扫描。

我的方法是采用 UUID 原始值编号（128 位值）和 16 位校验和，然后在转换为 QR 码之前对该数据进行 Base64 编码。到目前为止一切顺利，这非常有效。

为了生成校验和，我采用 UUID 的字符串版本并将其与一个长秘密字符串组合，并对奇数字节进行异或运算以生成 SHA-1 哈希。但是这个散列太长了，所以我将所有旧字节异或在一起以产生一半的校验和，同样与偶数字节产生另一半。

让我担心的是，我通过 XORing 不必要地损害了 SHA-1 系统。从结果中的某处获取两个未经处理的字节会更好吗？我承认 16 位校验和不如 160 位校验和安全，但这是我必须为条形码的可用性付出的代价。我真正不想找到的是，我现在提供了一个易于破解的校验和，因为 UUID 是明文传输的。

如果有更好的方法来生成校验和，这也将是该问题的合适答案。一如既往地非常感谢您抽出宝贵的时间或只是阅读本文，如果您发布答案，请加倍感谢。

我有一个包含以下信息的数据集：

• 工作负载名称
• 使用的配置
• 测量性能

在这里，您有一个玩具数据集来说明我的问题（性能数据根本没有意义，我只是选择了不同的整数以使示例易于理解。实际上，数据将是来自性能测量的浮点值）：

您可以使用以下方法生成它：

我正在尝试计算不同配置的谐波加速。为此，需要一个基本配置（在本例中为 cfg = 1）。然后谐波加速计算如下：

例如，对于配置 2，它将是：

我想计算每个工作负载对和配置的谐波加速。通过使用示例数据集，结果将是：

我正在努力寻找一个不使用循环aggregate的ddply解决方案，但我一直无法想出一个可行的解决方案。所以，我面临的基本问题是：

• 如何处理工作负载和配置之间的关系。给定工作负载对 (AB) 和给定配置的结果必须一起处理（谐波加速公式分母中的前两个性能测量来自工作负载 A，而其他两个来自工作负载 B）
• 对于每个工作负载对和配置，我需要使用配置库中的值“规范化”性能值（示例中的 cfg 1）

我真的不知道如何用一些 R 函数来表达这一点，比如aggregateor ddply（如果可能的话）。

有谁知道如何解决这个问题？

编辑：我有点害怕使用 1..8 asperf会导致一些混乱。为了简单起见，我这样做了，但值不需要是那些值（例如，想象像这样初始化它们：）dframe$perf <- runif(8)。詹姆斯和扎克的回答都理解我的部分问题是错误的，所以我认为最好在问题中澄清这一点。无论如何，我概括了这两个答案来处理配置 1 的性能不是 (1, 2) 的情况

我有关于基于扭曲的并行减少的想法，因为根据定义，扭曲的所有线程都是同步的。

所以想法是输入数据可以减少 64 倍（每个线程减少两个元素），而无需任何同步。

与 Mark Harris 的原始实现相同，减少应用于块级，数据位于共享内存上。 http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf

我创建了一个内核来测试他的版本和我的基于 warp 的版本。
内核本身完全相同地将 BLOCK_SIZE 元素存储在共享内存中，并在输出数组中的唯一块索引处输出其结果。

该算法本身运行良好。用完整的数组进行测试以测试“计数”。

实现的函数体：

1.他的版本的实现：

资源：

使用了 4 个同步线程 使用了
12 个 if 语句
11 个读取 + 添加 + 写入操作
1 个最终写入操作
5 个寄存器使用

表现：

五次测试运行平均：~ 19.54 ms

2.基于Warp的方法：（与上面相同的函数体）

资源：

使用了 1 个同步线程
7 if 语句
10 读取添加写入操作
1 最终写入操作
5 寄存器使用

5 位移位
1 加
1 子

表现：

五次测试运行平均：~ 20.82 ms

在具有256 mb 浮点值的Geforce 8800 GT 512 mb上多次测试两个内核。并以每块 256 个线程（100% 占用率）运行内核。

基于 warp 的版本慢了 ~ 1.28毫秒。

如果未来的卡允许更大的块大小，则基于 warp 的方法仍然不需要进一步的同步语句，因为最大值为 4096，它会减少到 64，而最终的 warp 会减少到 1

为什么它不快？或者这个想法的缺陷在哪里，内核？

从资源使用情况来看，翘曲方法应该领先吗？

Edit1：更正了内核只有一半线程处于活动状态，不会导致超出范围的读取，添加了新的性能数据

模幂运算的典型方程是 (a + b) MOD n = ((a MOD n) + (b MOD n)) MOD n。如果 a 和 b 非常大，那就太棒了。但是我被要求用一个非常大的 n (2^31 -1) 来做这个幂运算，a 和 b 没问题。

我只需要一种减少 n 的方法。

下面是我发现难以减少的 lambda 表达式，即我无法理解如何解决这个问题。

(λmn(λsz.ms(nsz)))(λsz.sz)(λsz.sz)

我迷失了。

如果有人能带领我朝着正确的方向前进，将不胜感激