感谢大家的支持，这礼拜让我们来谈谈 CUDA 多执行绪的程式模型(SIMT)，好让大家 对这个平行化的 C++ 有更清楚的轮廓。 ※ 第二章 SIMT 概观 ※ 所谓 SIMT (single instruction multi threads) 指的是单一指令对应多执行绪的 计算机架构，利用硬体的 thread 来隐藏 I/O 的延迟 (效果有点类似之前 Intel 的 hyper-threading，不过那不是 single instruction)，nVidia 进一步地让这些执行绪 可由程式控制，用群组的方式让一堆执行绪执行相同的指令，并利用超多核心来强化它 (例如 8800 GTX 有 128 颗、GTX280 有 240 颗)。 简而言之，它是把超级电脑的平行架构，浓缩到单晶片中，所以产生这样的效能 (例如我实验室里的 kernel，在 GTX280上跑的效能是 Intel Q9300 的 30 多倍， 这测量的时间是实际跑完的时间，用 CPU 的高精度 timer 测量出来的，对照的是用 intel 自家的 compiler 进行 SSE3 最佳化过的)。 不过刚开始进入这多执行绪的模型，还真的有点不太习惯哩。 ◆ CUDA 的平行化程式设计模型 名词定义 网格(Grid) ：包含数个区块的执行单元 区块(Block) ：包含数个执行绪的执行单元 执行绪(Thread)：最小的处理单元 (实际写程式的环境) CUDA 的平行化模型是将核心交由一组网格执行，再将网格切成数个区块，然後每个区块 再分成数个执行绪，依次分发进行平行运算，如果用军队来比喻，将核心视为连任务， 那网格就是连队，区块就是排或班，执行绪就是小兵。 任务(kernel) | | +--> 区块(排or班) +--> 执行绪(小兵) | | +--> 执行绪(小兵) | | +--> 执行绪(小兵) | | +--> 执行绪(小兵) | | +--> 网格(连队) +--> 区块(排or班) +--> 执行绪(小兵) | +--> 执行绪(小兵) | +--> 执行绪(小兵) | +--> 执行绪(小兵) | +--> 区块(排or班) +--> 执行绪(小兵) +--> 执行绪(小兵) +--> 执行绪(小兵) +--> 执行绪(小兵) (图一) kernel、网格、区块、执行绪和军队的类比 ◆ 内建变数 我们可以透过内建变数来辨识每个执行绪，让每个小兵弄清楚要执行那一部份的任务， 基本的内建变数如下，它们只可以使用在 kernel 的程式码中： uint3 gridDim ：网格大小 (网格包含的区块数目) uint3 blockIdx ：区块索引 (区块的ID) uint3 blockDim ：区块大小 (每个区块包含的执行绪数目) uint3 threadIdx：执行绪索引 (执行绪的ID) 其中 uint3 为 3D 的正整数型态 struct uint3{ unsigned int x,y,z; }; 可以运用它来实做更高层次的平行运算结构，不过现阶段，先不要管这种复杂的结构， 把它当成单一正整数即可，也就是 y 和 z 都当成是 1，只用 uint3 的 x。 ps. 其实我平常在写的时候，也很少用到3D结构，因为我们的研究是4D或5D的 ~>_<~ 只好用1D载入kernel再自已去切。 ◆ 网格 & 区块大小 (gridDim, blockDim) CUDA 透过指定网格和区块的大小形成平行化的程式阵列，总执行绪数目为网格大小和 区块大小的乘积，而 gridDim, blockDim 这两个变数在 kernel 函式中为内建的唯读 变数，可直接读取 总执行绪数目 = 网格大小(gridDim) x 区块大小(blockDim) 例如下图为 (网格大小=3, 区块大小=4) 所形成的核心，它具有 12 个独立的执行绪 +-----------+-----------+--------------------+ | | | thread 0 (id 0) | | | +--------------------+ | | | thread 1 (id 1) | | | block 0 +--------------------+ | | | thread 2 (id 2) | | | +--------------------+ | | | thread 3 (id 3) | | +-----------+--------------------+ | | | thread 0 (id 4) | | | +--------------------+ | | | thread 1 (id 5) | | grid | block 1 +--------------------+ | (kernel) | | thread 2 (id 6) | | | +--------------------+ | | | thread 3 (id 7) | | +-----------+--------------------+ | | | thread 0 (id 8) | | | +--------------------+ | | | thread 1 (id 9) | | | block 2 +--------------------+ | | | thread 2 (id 10) | | | +--------------------+ | | | thread 3 (id 11) | +-----------+-----------+--------------------+ (图二) 网格、区块、执行绪 ID 的划分 ◆ 呼叫 kernel 的语法 在 CUDA 中呼叫 kernel 函式的语法和呼叫一般 C 函式并没什麽太大的差异， 只是多了延伸的语法来指定网格和区块大小而已： kernel_name <<<gridDim,blockDim>>> (arg1, arg2, ...); ^^^^^^^^^^^ ^^^^^ ^^^^^^ ^^^^^^^^^^^^^^^ 核心函式名 网格大小 区块大小 函式要传的引数(和C相同) 所以只是多了 <<<gridDim,blockDim>>> 指定大小而已 ^^y 其中 gridDim 和 blockDim 可以是固定数字或动态变数，例如 (1) 固定数字 ooxx_kernel<<<123,32>>>(result, in1, in2); (2) 动态变数 int grid = some_function_g(); //计算网格大小 int block = some_function_b(); //计算区块大小 ooxx_kernel<<<grid,block>>>(result, in1, in2); ◆ 区块 & 执行绪索引 (blockIdx, threadIdx) 我们可以用区块和执行绪索引来定出正在执行的程式位置，以决定该载入什麽样的资料， 而 blockIdx, threadIdx 这两个变数和 gridDim, blockDim 一样，在 kernel 中也是 内建的唯读变数，可直接读取 例如在(图二)中，我们要定出每一个小兵的唯一的 ID，可用下面这段程式码 int id = blockIdx.x*blockDim.x + threadIdx.x; 要产生(图二)配置的 kernel 呼叫为 kernel<<<3,4>>>(arguments); 其行为如下 (1) 传入的网格和区块大小为 1D 正整数，所以 uint3 中只有 x 有用到，其它 y=z=1 (2) 网格大小 gridDim.x = 3 (每个网格包含 3 个区块) (3) 区块大小 blockDim.x = 4 (每个区块包含 4 个执行绪) (4) 区块索引 blockIdx.x = 0,1,2 (每个 thread 看到的不一样) (5) 执行绪索引 threadIdx.x = 0,1,2,3 (每个 thread 看到的不一样) (6) 区块基底 blockIdx.x*blockDim.x = 0,4,8 (7) 区块基底加上执行绪索引 id = blockIdx.x*blockDim.x + threadIdx.x = 0,1,2,3, 4,5,6,7, 8,9,10,11 所以我们便可得到一个全域的索引，即每一个小兵的唯一的 ID (图二中的 id 栏)。 ◆ kernel 函式的语法 用 CUDA 写 kernel 函式写一般 C 函式也是没什麽太大的差异，只是多了延伸语法来 加入一些特殊功能，并且标明这个函式是 kernel 而已： __global__ void kernel_name(type1 arg1, type2 arg2, ...){ ...函式内容... }; 其中 (1) __global__ : 标明这是 kernel 的延伸语法 (2) void : kernel 传回值只能是 void (要传东西出来请透过引数) (3) kernel_name : 函式名称 (4) type1 arg1, type2 arg2, ... : 函式引数 (和 C 完全相同) (5) 函式内容 : 跟写 C 或 C++ 一样 (但不能够呼叫主机函式) (6) global 函式只能在 host 函式中呼叫，不能在其它 global 中呼叫。 ◆ 小结 以上是 CUDA 平行化程式设计的概观，和传统 C/C++ 的差异便是它这种的 SIMT 结构， 也许你会觉得奇怪，为什麽要分成两层的 grid/block 结构，直接一层就配多个 thread 不是更简单，这牵涉到它硬体上的细节以及成本问题(後面章节会解释)，再者单层结构 不见得有效率，会增加同步化时执行绪等待的问题，使用两层结构，可以使 block 单元 弹性的选择同时或者循序执行，增加往後硬体发展和软体重用的弹性。 ※ 後续章节 ※ CUDA 安装 简易 kernel 范例 CUDA 的记忆体分类 CUDA 的函式种类 CUDA API介绍 GPGPU 的硬体介绍 (顺序还在研究中... >_<) ※ 名词解释 ※ (1)SIMT(single instruction multi threads)：单一指令对应多执行绪的架构。 (2)网格 (Grid) ：包含数个区块的执行单元。 (3)区块 (Block) ：包含数个执行绪的执行单元。 (4)执行绪(Thread)：最小的处理单元 (实际写程式的环境)。 (5)核心 (Kernal)：并非执行单元，比较像是要执行某种任务的抽象归类。 (6)网格大小(gridDim, grid dimension)：网格包含的区块数目。 (7)区块大小(blockDim, block dimension)：区块包含的执行绪数目。 (8)区块索引(blockIdx, block index) ：区块在网格中的位置。 (9)执行绪索引(threadIdx, thread index)：执行绪在区块中的位置。 (10)唯读变数(read-only variable)：只可读取，不可写入的变数。 (11)延伸语法(extension)：在标准C/C++语法之外，外加的功能性语法。 (12)函式引数(arguments)：函式呼叫时传递的变数。 (13)基底(base) ：计算位址时的基准点，就像座标的原点一样。 (14)索引(index)：位址相对於基准点的偏移。 (15)同步化(synchronize)：使多执行单元的进度在某点上对齐(先到的要等待还没到的， 等全部到齐後才继续前进)，通常是为了交换共用资料，避免读写顺序错乱导致的 资料错误。 --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.208.87 ※ 编辑: a5000ml 来自: 114.45.208.87 (10/02 02:06) ※ AE2001:转录至看板 NTUT_MMRE86 10/05 01:05 ※ a5000ml:转录至看板 C_and_CPP 10/15 20:59 ※ uf2000uf:转录至某隐形看板 10/16 21:48