我的CUDA学习笔记

研究生阶段开始接触CUDA。这里要感谢我的老板，感谢他的眼光和对我的信任，把这副担子交给我，我不想让他失望。在学习中慢慢陶醉于CUDA的神奇，被它强大的power所折服，这个过程实在太美妙，让我有很强烈的冲动把它记下来。这期间也受到一些前辈们博客的启发，我会在以下内容中特别提到。在项目实践中还有一些值得记下来的经验教训和典型bug，将另外专开一篇帖子记录。OK, let's begin.  

（一）、CUDA相关书籍：《GPU高性能编程CUDA实战》、《CUDA并行程序设计GPU编程指南》、《GPU高性能运算之CUDA》、《大规模并行处理器编程实战》、《OpenGL编程指南》、《GPGPU编程技术——从GLSL、CUDA到OpenCL》

（二）、程序体系：CUDA解构、线程模型、内存模型、

（三）、CUDA解构

一、代码结构：

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、CUDA runtime API & CUDA driver API 的用途、关系：必须用二者之一，才能实现管理GPU资源，在GPU上分配显存、启动核函数等功能。但是在一个程序中只能用二者中的一种，不能混合使用。

二、CUDA对C的扩展：（1）函数类型限定符。__global__、__device__等。（2）变量类型限定符。__shared__等。（3）内置矢量类型。dim3等。（4）引入了4个内建变量。blockIdx、threadIdx、gridDim、blockDim。（5）引入了<<<>>>运算符，用于指定grid和block维度。（6）引入了一些函数。__syncthreads等。

（四）、源文件&经典代码：

一、关于代码编写在源文件级别的区分：.cu文件里只编写与GPU计算有关的函数，.c或者.cpp文件里编写主函数和其他函数。

二、矩阵乘法：（1）若矩阵太大，可分割结果矩阵Pd成各个小块，由各个block计算。行索引为y，列索引为x。具体分法如下，

int Row = blockIdx.y * TILE_WIDTH + threadIdx.y;    \\Row：行。Col：列。

int Col  = blockIdx.x * TILE_WIDTH + threadIdx.x;

float Pvalue = 0;

for ( int k = 0; k < Width; k++ )

    Pvalue += Md[ Row * Width + k ] * Nd[ k * Width + Col ];

Pd [ Row * Width + Col ] = Pvalue;

dim3 dimGrid ( Width / TILE_WIDTH, Width / TILE_WIDTH );

dim3 dimBlock ( TILE_WIDTH, TILE_WIDTH );

MatrixMulKernel<<< dimGrid, dimBlock >>> (...);

（2）关键点：矩阵变量的索引。Md[ty * Width + k]；Nd[k * Width + tx]。

（五）、头文件：

一、#include <example_kernel.cu>//包含同一目录下的另一文件，<>中写的是文件名。

二、#ifndef...#dedine...#endif的作用：当存在多个文件时，避免头文件被重复引用。

-------------------------------------------------------------------------------------------------------

 #ifndef _EXAMPLE_KERNEL_H_

 #define _EXAMPLE_KERNEL_H_

 ...........................................................

#endif

-------------------------------------------------------------------------------------------------------

三、<stdlib.h>：标准库头文件。包含一些常用的系统函数。如malloc、free等。

四、<string.h>：包含字符数组函数的头文件。

（六）、宏：

（七）、函数：

一、GPU端引用数学函数：只能引用CUDA定义的一系列数学函数，而不能用hostt端的库函数。

二、CUDA事件函数：（1）本质：GPU时间戳。（2）用法：声明、创建、记录开始、记录结束、事件同步、显示时间、销毁

三、二维矩阵的表示：用一维表示，M[i*Width+k]。

四、__device__设备函数既不支持递归调用，也不支持指针进行间接的函数调用。

（八）、变量：

用来开辟空间的指针定义：float *h_odata = (float *)malloc(mem_size);

global函数内索引变量定义：

 const unsigned int bid = block;

 const unsigned int tid_in_block = blockIdx.x;

 const unsigned int tid_in_grid = blockIdx.x * blockDim.x + threadIdx.x;

共享内存定义：__shared__ float sdata[10];

共享内存赋值：

-------------------------------------------------------------------------------------------------------------------------------------------------

sdata[tid_in_block] = d_idata[tid_in_grid];  //把数据从global读入shared memory

g_odata[tid_in_grid] = sdata[tid_in_block];  //把数据从shared memory 读入global memory

--------------------------------------------------------------------------------------------------------------------------------------------------

kernel参数变量命名：threadsPerBlock；blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock

kernel函数中的中间变量即使设成一个，每个线程也会得出不同的结果值。

（九）、线程模型

一、线程：

 各个维度下threadID的索引：

二、线程束

0、线程：（1）tid的编译和多次读写时，第一次进行编译，之后就直接读取而不先编译。所以这时在改变tid值之后再读取的话是不可能的，解决这个问题的方法是对变量通过volatile关键字声明为敏感变量，编译器就会认为它可随时修改。这里注意在第二次读取前还要加上一条同步指令才能保证正确。（2）原子操作。

1、half-warp：指的是前16个线程，或者后16个线程。指令在硬件执行层涉及到 这个概念。

2、在算法上尽量做到warp内不分支。若不可避免的有分支时，尽量做到half-warp内的所有16个线程处在同一分支。

三、线程块：

四、线程网格：

五、CUDA通信机制

1、__syncthreads();

2、memory fence。__threadfence()：让线程在读写存储器的操作完成后，其结果对grid中的所有线程可见。__threadfence_block()：对block内的线程可见。代码示例：picture——博客图示——CUDA——线程模型——memory fence函数示例。

（十）、内存模型

一、register寄存器，由thread所有。物理位置在GPU片内，为高速存储器，访问延迟最低。生存周期与thread相同。

二、local mem，由thread所有，寄存器放不下时用它。物理位置在板载显存，很慢。生存周期与thread相同。对于一个小数组mt[3]，只定义不初始化==>local mem；定义的同时初始化==>register。

三、shared memory：（1）有shared必有同步__synthreads()，__shared__ float array[2];

四、global mem：（1）对二维数组分配线性空间：cudaMallocPitch()，三维用cudaMalloc3D()。对应的复制函数为cudaMemcpy2D()和cudaMemcpy3D()。注意，在访问由这两个函数分配的线性存储空间时，需要用到分配时返回的sride值或pitch值。

五、主机端内存：分为两种，可分页内存（pageable mem）和页锁定内存（page-locked 或 pinned）。可分页内存就是通过平常的malloc函数分配的存储空间；页锁定内存则是始终不会被分配到低速的虚拟内存中，而且是保证存在于物理内存中，并且能够通过DMA加速与设备端的通信。通过cudaHostAlloc()和cudaFreeHost()来分配和释放页锁定内存。可分页内存是双刃剑，有优有劣。好处是提高带宽。坏处是如果过多，会导致操作系统用于可分页的物理内存变小，导致系统性能下降。

六、常数存储器：（1）只读，物理位置位于显存，有缓存加速。（2）定义的话有两种方法。一是定义+赋值同时进行，如：__constant__ helloCUDA[2] = {1, 2}; ；二是先定义后赋值,如：__constant__ helloCUDA[2]; cudaMemcpyToSymbol(...);。（3）注意。定义时要定义在所有函数之外，作用范围为整个文件，并且对主机端和设备端都可见。

七、纹理存储器：（1）只读，物理位置位于显存，有缓存加速。（2）数据形式：一二三维数组。（3）操作。纹理绑定和纹理拾取。纹理绑定是把数据关联到参照系。可绑定的数据有普通的线性存储器和CUDA数组两种。绑定到纹理的线性存储器或CUDA数组中的元素称为像元texels（texture elements）。纹理拾取就是访问纹理存储器的操作。在这里，纹理拾取的坐标与数据在显存中的位置可以相同，也可以不同，这种映射方式用纹理绑定来描述。（4）纹理缓存有两个作用。一是重复利用数据，减少对显存的访问，节约带宽；二是纹理缓存一次预取拾取坐标对应位置附近的几个像元，可以实现滤波模式。（5）纹理存储器的特殊功能：有5个。①拾取坐标，可以是归一化或者非归一化两种；②寻址模式，有钳位模式（clamp）和循环模式（wrap）两种。③像元与纹理拾取后的返回值二者之间的数据类型转换。④滤波，针对纹理拾取后的返回值为浮点型的数据。滤波模式有最近点取样模式和线性滤波模式两种。最近点取样模式的返回值是与纹理拾取坐标对应位置最近像元的值，它不改变像元的值，适合用于查表。线性滤波模式是先取出附近几个像元的值，然后按照拾取坐标对应位置与这几个像元位置的距离进行线性插值。线性滤波的返回值是对最接近纹理坐标的两个像元、四个像元、或八个像元进行插值后得到的结果。使用线性滤波模式返回的值经过了插值处理，适合用于图像处理。（6）纹理存储器的使用。picture——博客图示——CUDA——内存模型——纹理内存——用法1,2,3,4,5

八、函数：

cudaMemcpy：不能用于多GPU系统中不同GPU之间的数据复制。

（十一）、算法

一、归约：（1）归约的循环控制变量 i /= 2; （2）归约代码里的线程之间需要通信，栅栏同步。共享存储，先开辟后存入。

二、分支：在算法上尽量做到warp内不分支。若不可避免的有分支时，尽量做到half-warp内的所有16个线程处在同一分支。

（十二）、经典代码示例

一、deviceQuery程序：（1）总体思路是获取设备数量，遍历每个设备，查询其属性。设备属性存在于结构体deviceProp中。