在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具体按计算要求.

使用charCUDA类型可能会对性能产生负面影响.我不建议使用char类型,除非内存大小限制(包括共享内存大小限制)或计算的性质特别需要它.

CUDA是一种遵循基本C++语言规范的C++派生语言.C++(和C)指定在进入计算之前,int必须加宽的类型的表达式数据int.除非底层硬件的整数指令带有内置转换,否则这意味着需要额外的转换指令,这将增加动态指令数量并可能降低性能.

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

但是,这并不总是可行的.一个简单的人为例子是以下内核,它I2I.S32.S8在使用T = charvs 实例化时包含一个额外的转换指令T = int.我通过运行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内在函数char,使用它们可以有利地减少动态指令数量.请注意,SIMD内在函数仅在Kepler GPU上由硬件完全支持,在Fermi CPU上完全仿真,并在Maxwell GPU上进行部分仿真.我已经看到轶事证据表明,与分别处理每个字节相比,即使是模拟版本仍然可以提供性能优势.我建议在任何特定用例的上下文中验证.

CUDA最佳实践指南第11.1.3节还对此问题进行了非常简短的参考:

编译器必须偶尔插入转换指令,引入额外的执行周期.这是......的情况

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