CUDA学习笔记（一）——CUDA基础

最新推荐文章于 2024-04-14 12:32:00 发布

houjian2015

最新推荐文章于 2024-04-14 12:32:00 发布

阅读量2k

点赞数

研究生阶段开始接触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可读写	可在程序中保持