DAY39:阅读扩展数据类型

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第39天，我们正在讲解CUDA C语法，希望在接下来的61天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。

本文共计317字，阅读时间15分钟

前情回顾：

DAY36：阅读”执行空间"扩展修饰符

DAY37：阅读不同存储器的修饰符

DAY38：阅读存储器修饰符

B.3. Built-in Vector Types

B.3.1. char, short, int, long, longlong, float, double

These are vector types derived from the basic integer and floating-point types. They are structures and the 1st, 2nd, 3rd, and 4th components are accessible through the fields x, y, z, and w, respectively. They all come with a constructor function of the form make_<type name>; for example,

int2 make_int2(int x, int y);

which creates a vector of type int2 with value(x, y).

The alignment requirements of the vector types are detailed in Table 3.

B.3.2. dim3

This type is an integer vector type based on uint3 that is used to specify dimensions. When defining a variable of type dim3, any component left unspecified is initialized to 1.

B.4. Built-in Variables

Built-in variables specify the grid and block dimensions and the block and thread indices. They are only valid within functions that are executed on the device.

B.4.1. gridDim

This variable is of type dim3 (see dim3) and contains the dimensions of the grid.

B.4.2. blockIdx

This variable is of type uint3 (see char, short, int, long, longlong, float, double) and contains the block index within the grid.

B.4.3. blockDim

This variable is of type dim3 (see dim3) and contains the dimensions of the block.

B.4.4. threadIdx

This variable is of type uint3 (see char, short, int, long, longlong, float, double ) and contains the thread index within the block.

B.4.5. warpSize

This variable is of type int and contains the warp size in threads (see SIMT Architecture for the definition of a warp).

本文备注/经验分享：

這一章節主要說了CUDA所引入的一些擴展數據類型, 他們的大小和對齊要求,以及,在kernel內部所能訪問到的一些重要內置變量(主要用於任務切分和線程ID判定) 。

本章节提供了一个表格(如图), 提供了一些内置的向量/矢量(vector)类型.这些类型基本上是普通的标量类型(C里面原本有的基本类型), 外加一个表示又几个分量的数字构成. 例如:float4, 这是连续的4个float.具体有哪些可以见此表. 但是需要说明的是：这些类型, 和连续的n个对应的标量类型的访问, 还是有区别的, 主要需要注意这些地方:

(1) 对齐要求发生了变化. 例如float4要求对齐到16B(4*4B)的边界. 而普通的4次float读取, 可以每次都放宽到4B的边界. 目前的卡(截至到Pascal), 对于这种类型的不对齐读取(例如你自己生成或者推导出来的指针), 在global memory上会直接挂掉kernel. 而shared memory上的则可以没事(但手册没说)。如果来自OpenCL的用户, 可以理解成普通的读取(float4 *)和vload4()的方式读取的区别。

(2) 具有更好的性能. 还是用这4个float和float4做对比. 虽然(1)中要求严格了, 但是往往float4具有更好的性能, 因为一次性的LD.128(16B)风格的读取, 比.32(4B)的读取有更好的性能: 编译器生成更少的指令(1条vs 4条访存指令), 对cache之类的更加友好.所以你看, 有弊就有利. 事情都是在代价和收益上权衡的.

（3）一些卡, 必须使用float4之类的向量类型才有性能, 主要是Kepler这一代:还记得之前的章节中, 我们说过, TLP和ILP是主要的掩盖延迟的方式吗? (前者是平铺一大堆线程, 靠线程间的自动切换来掩盖延迟, 发挥GPU的并行性能),而后者ILP则是线程内部的前后(或者说上下)的指令间,进行无关操作，发挥性能。而Kepler上必须使用这种向量类型, 因为Kepler太难掩盖延迟和发挥性能了. 因此使用float4这种, 等于在这代卡上, 前后能提供4条无关的依赖链(依赖链是一条前后相关的指令序列), 4个分量之间很多计算是无关的,一下子等于减轻了延迟掩盖要求4X，所以这些向量类型不仅仅能提高读写性能, 有的时候在一定的卡上,是必须的.

大致这三点针对向量类型的.

但有一些东西是必须需要额外说明的:

(1)有些数据本身放在一起比较容易理解, 例如一个复数的实部和虚部, 此时用float2之类的表示比较符合人类的认知.

(2)这里的float3或者int3这种比较奇特, 和OpenCL中的对应名字的占据空间大小, 以及, 对齐要求均不相同.(包括NV自家的OpenCL和自家的CUDA中的同名类型)。这个实际上之前提到过, 之前有个float3的例子, 当时我们说过,这个float3和内置的float3向量类型不同.

CUDA内置的float3是从节省空间的角度设计的, 大小只有12B, 对齐只需要到4B，而之前章节里面的那个float3, 则是从性能的角度考虑的, 大小是16B, 对齐也需要16B.前者节省了1/4的空间, 但需要被拆分成3次32-bit(4B)访问.后者浪费了1/4的空间, 但可以一次性128-bit访问(16B)，其实这是因为CUDA没有global memory上的不对齐访问支持, 也没有96-bit访问能力的缘故.

一些更好的硬件(但更耗电, 能效不高), 例如AMD的GCN 1.1(GCN2)+的卡,具有奇特的96-bit访问, 和不对齐支持,但CUDA不行, 所以你必须在2种布局(内置的这种, 和前面章节手册提供的那种)中, 2选1.这并不代表CUDA所依赖的硬件(N卡)不好, 从Kepler以后(不包含), N卡的能效得到了长足长进. 同样的一台N卡和装满A卡的GPU服务器对比, 如果24小时满载,放在非具有国家电费补贴的单位运行, 则在这台服务器出保前,光N卡节省下来的电费, 相比A卡, 就足够节省出来一台服务器了.因此虽然有轻微不便(例如你不对齐会挂kernel), 但总体来说还是划算很多的.所以无需计较.此外, 关于这些类型的对齐, 有个记忆上的小诀窍: (1)含有3结尾的类型是特殊的, 对齐到对于没有3结尾(或者等于1结尾的)；

(2)其他的所有对齐要求都是元素大小*数量.

(3)超过16B的只需要16B即可.

第三点是因为N卡的最大线程访存大小只有16B(warp聚合可以更大, 例如一次的L1的128B传输)

这三点就可以轻松记忆(例如从5.0开始, N卡都有上次使用的寄存器的缓冲(可以理解成L0 D-Cache), 节省了寄存器的bank conflict, 也省电了---而此时A卡没有这些特性)

然后说下这个表格对移植性的影响,很多人说, 我感觉Windows下特别方便, 我将一个kernel在Windows下调好了,然后直接在Linux编译这kernel，然后发现挂了居然, 这是为何?请注意这个表格里面有一些类型是没有固定大小的.主要是含有long的这些类型.而归根到底, 则是因为:有各种不同系统上的C的内置类型(例如ILP64和LP64模型中的基本数据类型的大小差异),也就是说, 不是因为CUDA C的原因, 而是因为你的Host C编译器, 不同的系统上这些类型的sizeof()大小都不一样.导致Windows和Linux下默认的一些数据类型虽然是同名的， 但发生了变化。而CUDA C总是要跟随Host编译器来提供最大的兼容性和无缝host/device代码链接和调用，从而引入的问题.

这可以说是一个易用性和轻微的小坑之间的矛盾和选择。OpenCL没有这个问题, 但它也没有这种无缝调用的方便(例如OpenCL上不能支持直接struct在host和device上传递, 等等)。

对于这个问题考虑到非常常见, 建议使用如下方式解决:

(1) 改用uint32_t这种明确的大小的能跨平台的类型.

(2) 总是明确不同的系统上, 不同的C类型区别(主要有Windows/Linux, Mac我不懂). 并建议总是使用sizeof(), 而不要硬编码你认为的代码.通过搜索刚才说的LP64, ILP64, LP32之类的字样, 可以获取更多信息.(实际上这个问题, 在当年的16-bit和32-bit过渡期间, 严重很多. 就如同现在的在32-bit到64-bit的过渡一样).

关于kernel内部能使用的重要内置变量, 例如threadIdx这种，还有本章节后续提到的blockIdx, blockDim, gridDim 。这些都是提供给线程用来获取自己的坐标信息, 以及, 全局的伙伴线程之类的信息的重要内置变量.你知道我们的CUDA是SIMT模型的, 从程序员的角度看, 我们写的代码实际上只是看上去,在处理一个大问题的一个局部，例如处理一张图片的一个或者邻近即可点. 而CUDA则为我们通过启动海量线程的方式, 将局部的这种处理, 自动的应用到了全部的整体问题——这就是为了说CUDA简单, 它的模型总是只需要你会处理一部分即可。而此时, 我们的所有线程, 执行的代码都是一样的, 他们唯一的区别就是这些内置变量的值上,每个线程都一样的执行代码, 但得到了不同的这些内置的值, 通过简单或者复杂的对这些值的变换(例如int tid = threadIdx.x + blockDim.x * blockIdx.x;），然后每个线程处理p[tid]这种形式, 就自动的将一个巨大的问题像蚂蚁啃象一样的吃掉了.所以这里应当需要注意这些. 我建议可以看一下NV的例子(安装CUDA开发包, 例如CUDA 9的时候会自带),你会很快的熟悉这些内置变量的.其他没有说的都是字面意思.

本章节还是基本无坑的(轻微小坑正常)

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖