在 CUDA 内核中使用 char 变量会受到惩罚吗？

Question

我似乎记得得到提示，我应该尽量避免在 CUDA 内核中使用 char，因为 SM 喜欢 32 位整数。使用它们会有一些速度损失吗？例如，是不是做的比较慢

int a[4];
int b = a[0] + a[1] + a[2] + a[3];
a[1] = a[3];
a2[0] = a[0]

比

char a[4];
char b = a[0] + a[1] + a[2] + a[3];
a[1] = a[3];
a2[0] = a[0]

在内核代码中？

注意事项：

• 我对使用 char 值进行算术运算、执行比较以及将它们读取和写入内存的惩罚感兴趣。

Answer 1

预先快速说明：在 C/C++ 中，char 的符号是实现定义的。因此，当使用char 执行 8 位整数运算时，强烈建议根据计算需要使用signed char 或unsigned char。

在 CUDA 中使用 char 类型可能会对性能产生负面影响。我不建议使用 char 类型，除非内存大小限制（包括共享内存大小限制）或计算的性质特别需要它。

CUDA 是遵循基本 C++ 语言规范的 C++ 派生语言。 C++（和 C）规定，在类型小于 int 的表达式中，在进入计算之前必须将其扩大到 int。除非底层硬件的整数指令带有内置转换，否则这意味着需要额外的转换指令，这将增加动态指令数并可能降低性能。

请注意，在“as-if”规则下，允许编译器偏离抽象 C++ 执行模型：只要生成的代码表现得好像它遵循抽象模型，即其语义相同，就可以消除这些转换操作。我最近的实验表明，CUDA 6.5 编译器正在积极地应用这种优化，因此能够彻底消除大多数转换操作，或者将它们合并到其他指令中。

但是，这并不总是可行的。一个简单的人为示例是以下内核，当使用T = char 与T = int 实例化时，它包含一个附加的转换指令I2I.S32.S8。我通过在可执行文件上运行cuobjdump --dump-sass 来转储机器代码来验证这一点。

template <class T>
__global__ void kernel (T *out, const T *in)
{
    int tid = threadIdx.x;
    if (threadIdx.x < 128) {
        T foo = 5 * in[tid] + 7 * in[tid+1];
        out [tid] = foo * foo;
    }
}

除了增加指令数之外，使用char 类型也会对​​性能产生负面影响，因为内存带宽较低。 GPU 内存子系统的设计使得总可实现的全局内存带宽通常随着访问的宽度而增加。对此的一种可能解释是跟踪内存访问的内部队列的深度有限，但可能还有其他因素在起作用。

char 类型由于用例的性质而自然出现，例如图像处理，人们可能希望研究 32 位复合类型的使用，例如 uchar4。在加载和存储操作期间使用更宽的类型可以提高内存带宽。 CUDA 有SIMD intrinsics 用于操作打包的char 数据，使用这些可以有益地减少动态指令数。请注意，SIMD 内在函数仅在 Kepler GPU 上由硬件完全支持，在 Fermi CPU 上完全模拟，在 Maxwell GPU 上部分模拟。我已经看到轶事证据表明，与单独处理每个字节相比，即使是模拟版本仍然可以提供性能优势。我建议在任何特定用例的上下文中验证这一点。

CUDA Best Practices Guide 的第 11.1.3 节中也有对此问题的非常简短的参考：

编译器有时必须插入转换指令，引入额外的 执行周期。情况就是这样……

• 在 char 或 short 上运行的函数，其操作数通常需要转换为 int。
• ...

Answer 2

算术

不可能从一般意义上说它是否会更快/更慢/不变，但通常我不认为会有太大差异。您说 chars 的算术将是 32 位是正确的，但这是否需要类型转换将取决于问题。在问题的示例中，我希望看到编译器将 a 和 b 存储在 32 位寄存器中，并且在我围绕这个问题的实验中（注意，如果没有完整的复制案例，很难保证这一点）我没有在SASS 中看不到区别。对于所有都在寄存器中完成的代码区域，我不希望性能受到影响。

但是，当char 变量被移动两次并从内存中移动时，会产生影响。由于char 在使用前必须转换为 32 位寄存器，因此会产生额外的指令。这可能会产生相当大的影响，也可能不会。

现在，还有一些边缘情况可能会有所作为。编译器可能能够将多个chars 打包到一个寄存器中并用算术提取它们（寄存器节省与算术成本）。您甚至可以使用联合来强制执行此行为。节省是否值得指示将根据具体情况而有所不同。目前我想不出任何其他会导致大量铸造开销的方法。

内存

很明显，如果您可以将变量存储在 1 个字节而不是 4 个字节中，那么您将节省 4 倍的内存和所需带宽。不过有一些事情需要考虑：

1. 共享内存。当前共享存储区大小为 4 字节或 8 字节。除非您正在读取每个线程至少 4/8 字节的事务，否则您无法实现峰值共享内存带宽。对于较小的交易，还需要考虑银行冲突。以 bank 大小的步长读取 1 字节将避免这些 bank 冲突，但会增加所需的内存并浪费带宽。
2. 全局内存。当您能够进行大型事务时，内存系统效率最高。 128 位事务往往比 64 位快，而 64 位往往比 32 位快。出于这个原因，打包（并对齐）数据是个好主意，这样您就可以通过一条指令将多个数据移动到一个线程中。

结论

我不知道有什么重要的理由不使用 char 而不是 int 进行算术运算，因为一切都在寄存器中，尽管在读取/写入内存时会支付转换成本。将数组存储为 char 而不是 int 应该，如果你小心的话，可以同时节省带宽和空间。