在 CUDA 内核中使用字符变量是否会受到惩罚

Is there a penalty to using char variables in CUDA kernels?

本文关键字：是否 惩罚 字符变量 CUDA 内核      更新时间：2023-10-16

我似乎记得得到提示，我应该尽量避免在 CUDA 内核中使用 char，因为 SMS 喜欢 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 值进行算术、执行比较以及将它们读取和写入内存的惩罚感兴趣。

预先快速说明：在 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 具有用于操作打包char数据的 SIMD 内联函数，使用这些内联函数可以有益地减少动态指令数。请注意，SIMD 内函数仅在开普勒 GPU 上完全由硬件支持，在费米 CPU 上完全仿真，在 Maxwell GPU 上部分仿真。我已经看到轶事证据表明，与单独处理每个字节相比，即使是模拟版本仍然可以提供性能优势。我建议在任何特定用例的上下文中验证这一点。

CUDA 最佳实践指南的第 11.1.3 节中也非常简要地提到了这个问题：

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

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

算术

不可能

在一般意义上说它是否会更快/更慢/不变，尽管通常我不会期望有太大差异。您说的是正确的，字符的算术将是 32 位的，但这是否需要类型转换将取决于问题。在问题的示例中，我希望看到编译器在 32 位寄存器中存储a和b，并且在我围绕此问题的实验中（注意，如果没有完整的重现情况，很难保证这一点）我没有看到 SASS 的区别。对于在寄存器中完成所有操作的代码区域，我不希望性能命中。

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

现在，也有一些边缘情况可能会有所作为。编译器可能能够将多个char打包到寄存器中，并使用算术提取它们（寄存器节省与算术成本）。您甚至可以使用联合来强制这种行为。节省是否值得说明将根据具体情况而有所不同。我想不出目前还有什么其他的会产生大量的铸造开销。

记忆

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

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

结论

我不知道有什么重要的原因如果不要使用char而不是int用于所有内容都在寄存器中的算术，尽管您在读取/写入内存时需要支付转换成本。如果您小心的话，将数组存储为char而不是int应该可以节省带宽和空间。

相关文章：