第3章 CUDA 简介
《大规模并行处理器编程实战》学习,其他章节关注专栏 CUDA C
CUDA C 编程友情链接:
- 第三章 CUDA 简介-CUDA C编程向量加法
- 第四章 CUDA数据并行执行模型
- 第五章 CUDA 存储器
- 第六章 CUDA性能优化(内附原书链接)
- 核函数:CUDA编程入门(一)-以图片运算看线程的组织和核函数的使用
- 拓展:CUDA卷积计算及其优化——以一维卷积为例
这章主要以 向量加法 vector add 为切入点,讲述了如何把一个 c 语言向量加法代码 改写为 cuda扩展的c语言向量加法。
1.1 传统向量加法
传统向量加法是通过循环实现的
1.2 CUDA 加法加速
cuda向量加法是通过多线程控制的cuda加法并行实现的,即同时打开n个线程,每个线程计算1个加法,则长度为n的向量被同步计算。使用CUDA扩展的程序,需要三个步骤:
- 先申请设备(device, cuda)的内存(memory),将数据从主机(host)复制到设备(device)
- 使用 设备API 对申请到的内存进行操作。(设备上的操作是以kernel函数的形式进行的)
- 把计算结果复制回主机
加法函数的具体进行
在以多线程的形式对向量进行操作时,多个线程以线程块的方式存在。一个线程块中包含blockDim(例如256)个线程,每个线程执行相同的代码。
加法代码如下,其中threadIdx.x表示该线程在本线程块中的Index,即在第i个块中,该线程是第j个,0<=j<=blockDim-1。i的计算为该线程在所有块中的index,即i=本块index + 一个块blockDim个线程 * 第blockIdx线程块:
值得注意的是,在 CUDA 交叉编程环境中,默认代码是在主机 host 上执行,若要在设备上执行,则需要 global 修饰,具体如:
Kernel 函数参数
重看Kernel函数,配置参数在<<< >>>之间,分别表示
- ceil(n/256.0) 需要 n/256 个线程块
- 256 每个线程块有 256 个子线程