有效评估 CUDA 中的索引映射函数答案

【问题标题】：Efficiently evaluating an index mapping function in CUDA有效评估 CUDA 中的索引映射函数
【发布时间】：2015-01-18 07:07:33
【问题描述】：

在 CUDA 内核中，我需要找到一个映射到 threadIdx 的键。

映射可能如下所示：

key -> threadIdx

0 -> {0,1,2,3,4}

1 -> {5,6,7}

2 -> {8,9,10}

...

每个键 k_i 都映射到 n_i（可变的、任意的 n_i 和 n_i>0）线程。该键将用于检索全局数组中的相应值。然后在这个内核的后续计算中使用这个值。

映射可以绘制为分段常数函数：

键的数量不限于 3 个（这只是一个示例！）并且仅在运行时知道，以及每个键的各自“宽度”。

如何有效地在CUDA内核中找出对应的key？我想到了以下两种选择：

在内核中使用二分查找（内存高效）
预先计算每个 threadIDx 的映射，然后启动内核（运行时高效）

0 0 0 0 0 1 1 1 2 2 2 ...

有没有更好的方法来实现这一点？

【问题讨论】：

将每个元素读取一次（到共享内存中）可能更有效。虽然缓存会在重复访问相同地址时为您提供帮助，但共享内存对于大量读取可能会更好。如果每个键的值大致为偶数，那么最好让每个线程处理其键的所有值。
您的映射似乎只是int key = ((int)threadIdx.x-2) / 3;，每个线程都可以轻松有效地计算它。如果除数是运行时变量而不是模板参数，这将需要一个完整的 32 位整数除法，但它仍然应该非常有效。你试过并计时了吗？
@njuffa 正如我上面指出的：每个“宽度”都可以不同！
@m.s.我不清楚这到底是什么意思。我认为“宽度”是不同的，而是键数的线性函数，正如您的示例所建议的那样[具有恒定“斜率”的阶跃函数]。在这种情况下，应该使用基于整数除法的公式。您是说“宽度”是完全任意的，即它不能由每个线程可以在运行时评估的封闭式表达式来描述？
@njuffa 是的，“宽度”是任意的。

标签： c++ cuda gpgpu gpu

【解决方案1】：

还有另一种算法可以在内存和运行时效率方面为您提供中值：假设线程总数为N。让我们取接近sqrt(N) 的数字M 并将所有线程按M 线程分组（最后一个将不完整）。现在，只为每个组中的第一个线程预先计算密钥（它们的 idxes 将是0、M、2M 等等）。它给了我们O(sqrt(N)) 记忆渐近。现在，在内核中，我们可以很容易地找到当前组（groupIdx = threadIdx / M）和下一组（groupIdx + 1）的索引。对于它们中的每一个，我们都知道预先计算的密钥key[groupIdx] 和key[groupIdx + 1]。现在你可以做 BS，但使用 [key[groupIdx]; key[groupIdx + 1]] 片段而不是 [1; MAX_KEY_VALUES] 进行搜索。

【讨论】：

该算法的复杂度与传统的二分查找相比如何？
这很难估计，因为它取决于关键线程的分布，但是，对于每个线程，类似O(log_2(MAX_KEY_VALUE/sqrt(N))) 而不是O(log_2(MAX_KEY_VALUE))。

【解决方案2】：

您可以组合 - 每个线程在启动时，使用 BS 找到它自己的密钥并将其存储在数组中。

【讨论】：

请看我的编辑，目标不仅是找到密钥，而且密钥用于检索正确的内存位置以进行进一步计算。所以我不需要存储密钥，我只是在内核中临时需要它们。