研究生阶段开始接触CUDA。这里要感谢我的老板,感谢他的眼光和对我的信任,把这副担子交给我,我不想让他失望。在学习中慢慢陶醉于CUDA的神奇,被它强大的power所折服,这个过程实在太美妙,让我有很强烈的冲动把它记下来。这期间也受到一些前辈们博客的启发,我会在以下内容中特别提到。在项目实践中还有一些值得记下来的经验教训和典型bug,将另外专开一篇帖子记录。OK, let's begin。
CUDA相关书籍不多,这里列举出一些经典书目。《GPU高性能编程CUDA实战》、《CUDA并行程序设计GPU编程指南》、《GPU高性能运算之CUDA》、《大规模并行处理器编程实战》(1、2版)、《OpenGL编程指南》、《GPGPU编程技术——从GLSL、CUDA到OpenCL》。
这一部分主要记录CUDA的软硬框架、线程模型和内存模型。
第一章 CUDA的总体框架
1 CUDA对C的扩展
CUDA对C的扩展主要体现在如下6个方面:
(1)函数类型限定符。__global__、__device__等。(2)变量类型限定符。__shared__等。(3)内置矢量类型。dim3等。(4)引入了4个内建变量。blockIdx、threadIdx、gridDim、blockDim。(5)引入了<<<>>>运算符,用于指定grid和block维度。(6)引入了一些函数。__syncthreads等。
2 API的一些事项
CUDA runtime API & CUDA driver API 的用途、关系:必须用二者之一,才能实现管理GPU资源,在GPU上分配显存、启动核函数等功能。但是在一个程序中只能用二者中的一种,不能混合使用。
3 代码框架
1 主机端代码框架
(1)启动CUDA,使用多卡时加上设备号,或使用cudaSetDevice()设置GPU设备;(2)为输入数据分配内存空间;(3)初始化输入数据;(4)为GPU分配显存,用于存放输入数据;(5)将内存中的数据拷贝到显存;(6)在GPU上分配显存,用于存放输出数据;(7)调用device端的kernel进行计算,将结果写到显存中对应区域;(8)为CPU分配内存,用于存放GPU传回来的输出数据;(9)将显存中的结果读回到内存;(10)使用CPU对数据进行其他处理;(11)释放内存和显存空间;(12)退出CUDA。
2 设备端代码框架
(1)存显存读数据刀GPU内;(2)对数据进行处理;(3)将处理后的数据写回显存。
3 代码编写上的风格规范
从源文件的层面说,.cu文件里只编写与GPU计算有关的函数,.c或者.cpp文件里编写主函数和其他函数。
从源代码的层面说,以下列出一些从书本上学到的良好的风格规范,提醒自己注意以后的培养。
1 头文件
(1)文件之间的包含关系在代码上的体现
#include <example_kernel.cu>//包含同一目录下的另一文件,<>中写入的是文件名。
(2)当存在多个文件时,避免头文件被重复引用的方法,那就是宏定义。
#ifndef _EXAMPLE_KERNEL_H_
#define _EXAMPLE_KERNEL_H_
...//正常工作代码
#endif
(3)一些C头文件
<stdlib.h>:标准库头文件。包含一些常用的系统函数。如malloc、free等。
<string.h>:包含字符数组函数的头文件。
2 C函数
3 C变量
(1)用来开辟空间的指针:float *h_odata = (float *)malloc(mem_size);
第二章 CUDA的线程模型
1 线程
1 线程的索引
这里主要搞懂各个维度下各种线程的索引就足够了。
tid_in_block = threadIdx.x;
tx_in_grid = threadIdx.x + blockIdx.x * blockDim.x;
ty_in_grid = threadIdx.y + blockIdx.y * blockDim.y;
tid_in_grid = x + y * blockDim.x * gridDim.x;
2 线程的多次读写
tid的编译和多次读写时,第一次进行编译,之后就直接读取而不先编译。所以这时在改变tid值之后再读取的话是不可能的,解决这个问题的方法是对变量通过volatile关键字声明为敏感变量,编译器就会认为它可随时修改。这里注意在第二次读取前还要加上一条同步指令才能保证正确。
2 线程束
在算法上尽量做到warp内不分支。若不可避免的有分支时,尽量做到half-warp内的所有16个线程处在同一分支。half-warp指的是前16个线程,或者后16个线程。指令在硬件执行层涉及到 这个概念。
3 线程块
1 线程块的分配
这里的重点在线程块的分配,这要依据要处理的问题和GPU本身的硬件参数,最终会影响到到GPU的利用率、代码的效率。
2 线程块在代码上的设置
一种专业设法是用 #define 三目运算符
#define imin( a, b ) ( a<b?a:b )
const int N = 33 * 1024;
const int threadsPerBlock = 256;
const int blocksPerGrid = imin(32, (N + threadsPerBlock - 1) / threadsPerBlock );
4 线程格
这里的重点在于线程格的布局和代码表示。
1 布局
横条形,以保证内存访问连续。且尽量保证32的整数倍。
2 代码表示
dim3 blocksPerGrid( 1, 1)
第三章 CUDA的内存模型
CUDA中涉及到8中存储器类型,其中device端有6种,host端有2种。
存储器 | 位置 | 拥有缓存 | 访问权限 | 变量生存周期 |
---|---|---|---|---|
register | GPU片内 | 它本身就是,速度快 | device可读写 | 与thread相同 |
local memory | 显存 | 无,慢 | device可读写 | 与thread相同 |
shared memory | GPU片内 | 它本身就是,快 | device可读写 | 与block相同 |
constant memory | 显存 | 有,快 | device可读,host可读写 | 可在程序中保持 |
texture memory | 显存 | 有,快 | device可读,host可读写 | 可在程序中保持 |
global memory | 显存 | 无,慢 | device可读写,host可读写 | 可在程序中保持 |
host memory | 内存 | 无 | host可读写 | 可在程序中保持 |
pinned memory | 内存 | 无 | host可读写 | 可在程序中保持 |
第四章 CUDA实践
1 各种并行架构概览
各种并行架构之间并不等同,许多并行程序和并行语言是针对MIMD模型的,即线程是独立的,不需要像GPU上那样成组的执行。而且,并行程序都是需要做一些修改之后,才能在特定的GPU硬件上工作。目前,对具体架构的优化通常需要将应用程序与特定的硬件相绑定。
MPI和OpenMP的标准就不适合GPU模型,二者之中OpenMP可能更相近些,因为它有共享内存的概念。OpenACC是一种基于指令的GPU计算方法。
2 GPU在递归运算上的盲点
许多CPU算法都用到了递归,递归可以很容易的将问题分解成越来越小的问题直到成为一个非常简单的问题,二分查找就是一个典型的例子。但是在GPU上,只有计算能力为2.x的GPU才支持递归,并且只有带__device__前缀的函数可以使用,带__global__前缀的函数无法使用。
再回来说说递归,任何递归算法都可以描述成迭代算法,比如二分查找,快速排序。快速排序就是典型的采用递归方式实现的,它的实现过程是先选择一个基准点,把小于或等于这个基准点的数放在左侧,大于的放在右侧,然后得到两个独立的数据集,接下来就可以通过两个单独的线程再分别对它们排序。