CUDA 动态索引

Question

我注意到，使用动态索引会将 CUDA 代码的速度降低 12 倍 - 请参见以下示例：

__global__ void static3Ops(int start, int end, const float* p, const int* prog_dont_use, float* c)
{
    int i = threadIdx.x;
    float buf[5];
    buf[0] = 1.0e7;
    buf[1] = c[i];
    const int prog[] = { 0,1,2,3,4,5 };

    for (long j = start; j < end; j++) {
        buf[2] = p[j];
        buf[3] = buf[prog[0]] + buf[prog[1]];
        buf[4] = buf[prog[2]] - buf[prog[3]];
        buf[1] = buf[prog[4]] * buf[prog[5]];
    }
    c[i] = buf[1];
}

比

快 12 倍

__global__ void static3Ops(int start, int end, const float* p, const int* prog, float* c)
{
    int i = threadIdx.x;
    float buf[5];
    buf[0] = 1.0e7;
    buf[1] = c[i];

    for (long j = start; j < end; j++) {
        buf[2] = p[j];
        buf[3] = buf[prog[0]] + buf[prog[1]];
        buf[4] = buf[prog[2]] - buf[prog[3]];
        buf[1] = buf[prog[4]] * buf[prog[5]];
    }
    c[i] = buf[1];
}

任何提示如何最小化该开销？动态特性是我的代码的核心特性......所以没有它我几乎无法解决......

请注意，CPU 的开销仅为 20% 左右。

Answer 1

我能想到的两种可能：

如果 prog 是一个小数组：使用您自己的解决方案！即使用 prog 就像它在顶部示例中的定义一样，如果 prog 真的是一个包含少量元素的数组（如您的示例）。但是您对“动态性质是我的代码的核心功能”的评论使听起来这不是您的选择。当我将const int prog[] = { 0,1,2,3,4,5 } 更改为int prog_0 = 0, prog_1 = 1, ... 并使用prog_0、prog_1、...而不是prog[] 时，我得到了相同的性能。这表明prog[] 的值直接存储在寄存器中，而不涉及全局内存。如果prog 不是一个小数组或在编译时未知，这种方法可能会导致大量使用本地内存并显着降低性能。

如果 prog 是一个大数组：让线程将 prog 并行加载到共享内存中，然后在内核的其余部分相应地访问共享内存（块级 tiling） .

__shared__  int prog_sh[6]; // or dynamically allocate if size is not known
int i = threadIdx.x;
if (i < 6)
    prog_sh[i] = prog[i];
__syncthreads();

// and then use prog_sh instead of prog....

请注意，这对于像您的示例这样具有已知值的小数组确实没有意义，但是您会惊讶于在使用大数组时通过平铺获得了多少收益。尽管如此，在处理共享内存时，您应该确保能够为并发访问实现高内存带宽（请参阅this link）。

Answer 2

避免指针别名

首先：在所有指针上使用__restrict！超级重要！在此处阅读：

CUDA 专业提示：Optimize to avoid pointer aliasing

现在，除此之外......

使用您的访问模式并尝试改善内存局部性

如果：

• prog 的大小以一个小值为界，并且
• 对prog 的访问是对编译时已知的索引（即不是值，而是索引）

然后：

• 使用内核本地、纯 C 样式数组或 std::array-like 类，例如kat::array 来自the cuda-kat library's development branch （适当披露：这是我正在研究的一个库，所以我在这里有偏见。此外，数组实现非常稳定）。从作为参数获得的prog 指针加载它们的值。
• 仔细排列内存中的数据，以便可以合并加载到prog 数组中。因此，例如，第一个线程的第一个 prog 元素，然后是第二个线程的第一个元素等，直到第 31 个线程的第一个 prog 元素。
• 在使用任何值之前将所有负载加载到 prog 中。

如果：

• prog的大小不受小值的限制，而是
• 您可以安排 prog 的使用，以便对于它的每一个小的、有界长度的延伸，对它的访问都与某个基线有一个固定的偏移量（例如：循环 i，在迭代 i 时，我们访问 @987654333 @、prog[k*i + 3]、prog[k*i + 4] 仅限），

然后：

• 与前一种情况相同，但对于prog 的每个固定长度延伸。

如果：

• prog 不是那么小但也不是那么大（每个线程有几十个元素到几千个元素），并且
• 其中的访问模式是随机的、任意的或依赖于数据的

然后：

• 首先将prog 加载到共享内存中。
• 确保加载它，以免发生存储库冲突，即共享内存中每个块线程的 prog 等效项应完全包含在单个存储库中。

如果以上都不成立，那么：

• 保留允许合并访问的内存布局和访问模式。
• 尝试将大量数据集中在您期望相对较近的位置（或者它可能不那么重要，这更像是一周的提示）。

有条不紊的笔记

永远记住，无论您做什么 - 分析和分析它，不要只满足于底线数字。并尝试分解更改并分别对其进行分析。例如 - 首先添加 __restrict 看看它给了你什么。 CUDA“nSight 计算”还应该告诉你瓶颈在哪里（虽然不是如何处理它们......）

Answer 3

谢谢大家的提示！

目前我找到的最快的代码如下：

_global__ void static3OpsShared(int start, int end, const float* prices, const int* __restrict__ prog, float* c)
{
    int i = threadIdx.x;
    __shared__ float buf[5];
    buf[0] = 1.0e7;
    buf[1] = c[i];
    // I never use more than 6 values of prog in a single thread - but each thread has its own set
    // values of prog are ranging from 0...5 
    // Performance needs to focus on what happens within the following loop typically having over 10000 iterations
    for (long j = start; j < end; j++) { 
        buf[2] = prices[j];
        buf[3] = buf[prog[0]] + buf[prog[1]];
        buf[4] = buf[prog[2]] - buf[prog[3]];
        buf[1] = buf[prog[4]] * buf[prog[5]];
    }
    c[i] = buf[1];
}

（请暂时忽略共享内存索引 - 到目前为止我用一个线程运行它）

以

的形式使用 prog[0]...prog[5] 的寄存器

r0 = prog[0];

并使用buf[r0] instead of buf[prog[0]] 似乎是由优化器完成的。

对 buf[] 使用共享内存的最大改进。 restrict 在某种程度上没有帮助。尤其是限制不适用于 buf 作为应该被重复使用的值。

我的结论是： -- 如果可以使用寄存器而不是 buf[]，代码将快 5 倍左右。

Answer 4

提高速度的想法：

如果可能，使用 threadIdx 和 blockIdx 计算 CUDA 代码中 prog 的变化。计算比内存访问快。

小心共享内存（你提到它被忽略，但无论如何在这里）。您必须确保块的每个线程使用不同的索引。并确保 warp 中每个线程的索引转到不同的银行，否则会降低性能。

所以如果你有大小为 128 个线程的块并且 i 包含线程号：

__shared__ float buf[128 * 6];
buf[0] -> buf[0*128 + i];
buf[1] -> buf[1*128 + i];
buf[prog[0]] -> buf[prog[0]*128 + i];
...

由于块大小 (128) 可以被 32 整除，因此即使 prog 索引不同，warp 中的每个线程也会访问另一个共享内存库。线程 0 总是访问 bank0 等等。

另类

尝试将 buf 直接保存在寄存器中而不是共享内存中：buf0、buf1、buf2、...

如何使用索引访问它们？只需使用 switch case 编写内联函数或宏。

有 6*6*6*6*6*6 种可能性。您可以尝试通过为 36 或 216 种可能性生成代码来进行优化，然后调用适当的一种。例如

switch(prog01) {
case 0: buf3 = buf0 + buf0; break;
case 1: buf3 = buf0 + buf1; break;
...
case 6: buf3 = buf1 + buf0; break;
...
}

但可能会更快，如果您使用 6 个案例进行 6 次切换，那么您的案例/比较/跳转次数会更少。

最好的办法是：在循环外执行一半的开关（例如 216），在 216 个循环之一内执行一半的开关。

可能在循环外创建设备函数指针并调用相应的函数会更好，该函数会选择 bufs。但是这些函数必须从 bufs 中选择作为函数参数而不是局部变量。希望它们仍然有效地存储在寄存器中。

本地访问通常非常快。如果你有本地开关和跳转，你会失去计算时间。因此，请确保每个 warp 的线程都与相似的 prog 参数对齐。

请与完整的扭曲（不仅仅是一个线程）进行比较和测试，以获得更真实的结果，包括共享内存上的银行冲突。