参考文档：
学习目的：学习CUDA 就是为了发挥其并行结构的特点，加速源程序的速度，提升性能(CUDA有关的性能提升一般均指速度提升了)。
基础前言
几条基础知识(掌握了一下几条再去看下面的内容会容易些)：
- CUDA C与C/C++高度兼容，学习CUDA C只需要有C语言基础就够了。他们两在绝大部分上没有区别，这也是CUDA C流行的主要原因。
- Host指CPU，Device指GPU。
- GPU是图形处理器，是集成在显卡上的。GPU类似于CPU，显卡类似于计算机。显存类似于内存。
- 在Linux系统上，C/C++的编译采用GNU gcc/g++，而CUDA C采用
nvcc编译。 - CUDA C编写的文件以
.cu为后缀。 - CUDA C中最重要的函数叫核函数(Kernels)，用
__global__关键字定义，意味着编译器将编译为在设备上运行，在主机中调用。 - 第二个重要的是修饰字符
<<<M,N>>>，表示告诉运行时如何启动设备代码，而非传递给设备参数(传参和普通C一样，用圆括号里的参数)。其中M表示块block的数量；N表示一个block里线程threads的数量。可以认为运行时创建核函数的M*N个副本，并以并行的方式运行他们。此外CPU在调用这个核函数之后，将会义无反顾的立刻返回执行接下来的程序，这个核函数将会在设备上异步执行。 - 1一个格子grid有
gridDim.x个block，一个block中有blockDim.x个threads。因为CUDA支持一维、二维、三维的线程块数组，可以使用一维、二维、三维的索引，因此就有了.x、.y、.z。 - 块索引
blockIdx.x(当然还有yz)线程索引threadIdx.x。通过这些索引值，我们就知道当前是哪个线程块，哪个块在运行。具体来说：
//省略main中核函数调用dot<<<1,4>>>(d_a, d_b, d_c);
__global__ void dot(float *a, float *b, float *c)
{
int tid = threadIdx.x;
c[tid] = a[tid] * b[tid];
}
那么实际上，4个线程相当于在同步执行以下代码：
-
CPU和GPU共同协作必定存在同步问题。CUDA一些API函数有些是同步函数(即阻塞CPU运行，强行让CPU等待GPU完成再继续往下走)，比如你查询官方文档的时候，会出现这样一条：
这就是典型的同步函数了。
这就是典型的异步函数了。
此外，对于那些异步函数。CPU在调用之后，就会马上离开，这些异步函数会在自己的设备上运行。 -
因为显存和内存是两套不同的地址空间，因此指针虽然可以作为参数传递、执行计算、转换类型，但是绝不可以进行解引用，也就是说Host端的指针只能访问主机内存，Device端的指针只能访问设备显存。
-
不管什么类型的语言，都逃不过写入内存、计算、读出内存三个步骤，CUDA C也是如此，在遵循这个大框架进行学习会容易的多。
-
由于硬件的限制使得我们每个block的线程数有限；每个grid的块数量有限，那么当计算对象长度很长时侯，我们需要结合block和thread来索引：
代码采用固定的写法：
int tid = threadIdx.x + blockDim.x * blockIdx.x;
- 那如果结合之后还是不够怎么办？代码中仍是有固定的写法，以后都可以这么用：
int tid = threadIdx.x + blockDim.x * blockIdx.x;
while(tid<N)
{
temp += a[tid] * b[tid];
tid += blockDim.x * gridDim.x;
}
blockDim.x和gridDim.x其实就是尖括号中M、N的值，即实际执行的线程数和块数。但这样做其实是损失了一定的并行度的，但有了这个技术之后，对于尖括号中M、N的赋值就随意多了。比如你可以写：
dot<<<256,128>>>();
- 对于尖括号中block数目，一般有2种做法，一种是：
#define blockspergrid imin(32, (N+threadsperblock-1)/threadsperblock)
这样做就是为了避免若N小于threadperblock的时候，由于整除而使得线程数为0。并且当N大于threadperblock的时候，会分配足够的线程，不至于超过数组的边界。
另一种是取SM数量的2倍。
16. 一个中可以定义多个不同的Kernels。
三个基础函数
为了方便(CSDN编辑器字数超过一定，写起来很慢)写作，我把这部分内容放在另一篇中，请点击。
除此之外，还需要补充几点：
- cudaMemcpy中最后一个kind参数有四种，其是用来指定设备内存指针究竟是源指针还是目标指针，
比如
kind=cudaMemcpyHostToDevice，那就是主机复制到设备。因此源指针a就是主机指针，d_a就是设备指针，指向的是设备中的地址。
设备查询
这点其实还蛮重要的，因为设备查询可以找到你显卡的极限在哪。如果你不知道你硬件设备的极限，那么写代码的时候就容易出错。
常用的操作：
查询显卡中插了几块GPU：
int count;
cudaGetDeviceCount(&count);
printf("%s\n", count);
查询设备极限：
cudaDeviceProp prop;
for (int i=0; i< count; i++) {
HANDLE_ERROR( cudaGetDeviceProperties( &prop, i ) );
printf( " --- General Information for device %d ---\n", i );
printf( "Name: %s\n", prop.name );
printf( "Compute capability: %d.%d\n", prop.major, prop.minor );
printf( "Clock rate: %d\n", prop.clockRate );
printf( "Device copy overlap: " );
if (prop.deviceOverlap)//是否支持设备覆盖，用于CUDA流
printf( "Enabled\n" );
else
printf( "Disabled\n");
printf( "Kernel execution timeout : " );
if (prop.kernelExecTimeoutEnabled)
printf( "Enabled\n" );
else
printf( "Disabled\n" );
printf( "whether can map host memory : " );
if (prop.canMapHostMemory)//是否支持buffer映射，用于零拷贝内存
printf( "Enabled\n" );
else
printf( "Disabled\n" );
printf( " --- Memory Information for device %d ---\n", i );
printf( "Total global mem: %ld\n", prop.totalGlobalMem );
printf( "Total constant Mem: %ld\n", prop.totalConstMem );
printf( "Max mem pitch: %ld\n", prop.memPitch );
printf( "Texture Alignment: %ld\n", prop.textureAlignment );
printf( " --- MP Information for device %d ---\n", i );
printf( "Multiprocessor count: %d\n",
prop.multiProcessorCount );
printf( "Shared mem per mp: %ld\n", prop.sharedMemPerBlock );
printf( "Registers per mp: %d\n", prop.regsPerBlock );
printf( "Threads in warp: %d\n", prop.warpSize );
printf( "Max threads per block: %d\n",
prop.maxThreadsPerBlock );
printf( "Max thread dimensions: (%d, %d, %d)\n",
prop.maxThreadsDim[0], prop.maxThreadsDim[1],
prop.maxThreadsDim[2] );
printf( "Max grid dimensions: (%d, %d, %d)\n",
prop.maxGridSize[0], prop.maxGridSize[1],
prop.maxGridSize[2] );
printf( "\n" );
}
事件
事件用于测量CUDA C代码的性能，即执行速度，可以用来评价优化源代码的程度。
事件是CUDA的一个API，本质是一个GPU的时间表。用户可以指定何时开始记录，何时结束记录。
需要注意的是：CUDA事件只能用于对核函数和设备内存复制的代码进行计时。如果包含其他代码，那么计时结果就不如你所愿。
使用方式(先后顺序)：
- 创建事件：cudaEventCreate()
- 记录事件：cudaEventRecord()
注意第二个参数是和CUDA流有关。 - 同步事件：cudaEventSynchronize()
- 记录时间差：cudaEventElapseTime()，如果没有同步事件的话，CPU会直接调用cudaEventElapseTime()，但此时GPU上可能还在跑程序，还没有记录结束事件，这时候你记录的时间差就会有问题。单位为毫秒。
- 销毁事件：cudaEventDestroy()
下面是固定搭配：
int main( void ) {
// capture the start time
cudaEvent_t start, stop;
HANDLE_ERROR( cudaEventCreate( &start ) );
HANDLE_ERROR( cudaEventCreate( &stop ) );
HANDLE_ERROR( cudaEventRecord( start, 0 ) );
...
...
...
// get stop time, and display the timing results
HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
HANDLE_ERROR( cudaEventSynchronize( stop ) );
float elapsedTime;
HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime,
start, stop ) );
printf( "Time to generate: %3.1f ms\n", elapsedTime );
HANDLE_ERROR( cudaEventDestroy( start ) );
HANDLE_ERROR( cudaEventDestroy( stop ) );
...
...
...
return 0;
}
共享内存
每一块SM都会一块共享内存区域，共享内存的好处：
- 共享内存驻留在
物理GPU上，而不是Global Memory(显存)，因此访问共享内存的延迟要远远低于访问显存的延迟，故共享内存就像每个块的高速缓存。 - CUDA编译器将创建一个共享内存变量的一个
副本，每个block块一个副本，一个block中的所有线程共享这块内存。不同block是无法访问其他block共享内存中的数据。使得同个块内的线程之间可以协作通信而不受外界干扰。 - 共享内存可以减弱原子操作中的竞争程度，从而进一步优化性能。
共享内存的创建：
__shared__ cache[threadsperblock];
Note：
- 编译器将为每个线程块block
自动默认生成共享变量的一个副本，因此只需根据线程块中线程的数量来分配即可。 - 共享内存的访问只需要线程索引，不需要块索引，因为每个块之间是独立的，是block私有的，不同块之间无关联。
- 如果输入矢量的长度不是线程块中线程数量的整数倍，那么cache中将会有一些threads啥事都不做。这样资源就浪费了，因为这部分内存本来这个时间可以用来做其他事。
- 共享内存中一定会出现线程之间不同步的现象，那么当我线程A要读取线程B计算完的数，而线程B此时并未完成计算，若不加阻塞，就会导致出错。我们可以通过加入__syncthreads()来同步共享内存中不同线程进度不一致的情况。
- 共享内存的共享指的是我一个实验室(block)的人(threads)用这块地，其余实验室用别的地方。一个实验室的人待在一块，不同实验室的人不户干涉。而不是说一个block只有一块内存，是有
blockDim.x块的。各个线程并不是读取同一个地址，并不是类似于pytorch中的共享内存那个共享。
关于__syncthreads()需要注意的地方：
首先要解释下线程发散的概念：当某些线程执行一条指令，而其他线程不需要执行。比如if分支语句中，就可能产生一部分线程执行if中的语句，另一部分线程只能空着等待。
CUDA架构确保除非每个线程都执行了__syncthreads()，否则没有任何线程能执行__syncthreads()之后的指令。比如说__syncthreads()在if语句中，那么有一部分线程无法执行__syncthreads()，因为进不去if里面。这会使得处理器将挂起，因为GPU在等待某个永远不会发生的事。
常量内存
和共享内存一样，常量内存也可以加速应用程序的执行。
需要介绍两个概念：内存带宽、内存通信量、线程束。
内存带宽：
简单地理解就是读取内存数据的速率。
内存通信量=内存带宽
⋅
\,\cdot\,
⋅时间。
线程束指的是一个包含32个线程的集合，半线程束中每个线程只有在读取同一个地址上的内容时，才会极大提升性能，反之若每个线程分别读取不同地址时，那么这16次读取将会被串行化，需要16倍的时间来读取。而全局内存读取起码是同时读取的，因此其还不如在全局上读呢。
GPU可以加速应用程序，但是其瓶颈在于内存读取速率跟不上CUDA core的计算速率。CUDA core(SP)是一种ALU，GPU进行并行计算，也就是很多个ALU同时处理，当ALU这边快速计算完之后，需要去内存中取数，但内存的延迟使得无法立即取到数据，因此会影响到整个程序的速度。
常量内存给出了一个解决办法，就是将一部分在核函数期间不会发生变化的数据存在特定内存区域，这样的话就使得内存通信量降低，得到了缓解，从而降低了内存带宽。
为何内存带宽的降低了呢？
- 对常量内存的单次读操作可以广播到半个线程束，半个线程束是16个线程，故其实你只需要读一次就好，这将节约15次读取操作，从而将内存通信量节约94%。
- 在第一次对某个常量内存地址处的数据读取之后，硬件将会把该数据缓存到GPU上。之后的半线程束读取同一个地址时，就会往这个缓存中去读取，这相当于又一次减少了内存通信量。不仅如此高速缓存使得半程束读取得更快。
接下来，我们就可以正式引入常量内存：
常量内存不能用于修改，故在我们自己的应用程序设计中，应该用于保存不变的数据。
__constant__ int array[N];
Note:
- 常量内存不再需要用cudaMalloc和cudaFree，而只需要分配一个固定大小的数组即可，一般定义为具有文件作用域的全局变量。
- __constant__将变量的访问设置为已读。
- 常量内存的引入，使得内存通信大大降低，同时又满足了超高的计算速率。因此这是一种可以降低内存带宽，减少内存请求，提供更高效内存带宽的技术。
纹理内存
纹理内存是一种只读内存，其缓存在SM中，能够减少对内存的请求并提供高效的内存带宽。纹理内存适用于相邻线程访问空间位置临近的数据。
图中四个地址并非连续，但是空间上是连续的，将这样的数据存在纹理内存中，就可以加速访问，获得性能提升。因此对于具有空间局部性的待访问对象，我们可以将它放到纹理内存上。
纹理内存的应用场所讲完了，接下来就是怎么用？主要分4步，用①②③④表示。
①首先是声明一个纹理引用texture<class> xx
texture<int> age
- 文件作用域。
- 32位整型变量age声明为纹理内存的引用。
②接下来需要为缓冲区分配GPU内存，比如cudaMalloc，然后通过cudaBindTexture()将这些变量绑到内存缓冲区。
比如:
texture <float> num;
...
int main(){
cudaMalloc((void **)&data, imagesize);
cudaBindTexture(NULL, num, data, imagesize);
...
...
}
这样一来就相当于将指定缓冲区data作为纹理内存来使用，这块纹理内存的“名字”是纹理引用num。两者共用一个内存地址。
③可以开始启动核函数，使用tex1Dfetch():
- 告诉GPU将读取请求转发到纹理内存而不是标准的global memory(显存)。
- 编译器内置函数。
- 编译器需要在编译时知道tex1Dfetch应该对哪些纹理采样。
- 在kernels里是这样使用tex1Dfetch(num, offset)，其中offset是这样定义的：
int x = threadIdx.x + blockIdx.x * blockDim.x;//横向有多少个线程
int y = threadIdx.y + blockIdx.y * blockDim.y;//纵向有多少个线程
int offset = x + y * blockDim.x * gridDim.x;//
其实offset就是访问位置的二维索引，可见在底层是按一维数组展开。
- 格式：
value = tex1Dfetch(num,top)
一般情况下访问你数组元素是这样的：
num[top];
④清楚纹理引用于缓冲区的绑定：
cudaUnbindTexture(num)；
以上都是一维纹理内存，接下来讲下二维纹理内存。二维仅仅只是让代码变得简洁，性能和一维的效果差不多的。有三处需要修改，用①②③表示：
①二维纹理声明：
texture<int, 2> num;
②tex1Dfetch改为tex2D:
- 不需要再定义offset来表示空间相邻的数据，直接通过x，y来访问纹理：
value = tex2D(num, x, y-1)//表示访问上面那个数据。
- 不用担心边界溢出问题，如果x或y超出0或宽度，那么就会返回0处的值或宽度处的值。
③修改绑定：
cudaMalloc( (void**)&a, imageSize );
//描述通道格式
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
cudaBindTexture2D( NULL, texConstSrc,
a,
desc, DIM, DIM,//纹理的维数(DIM*DIM)
sizeof(float) * DIM );
cudaBindTexture2D这个API的详细如下：
但是对于这个pitch形参不是很清楚啥意思，有知道的麻烦告知下。
原子操作
GPU高性能这本书中举的递增例子很好解释了什么是原子操作：
C语言中对于这个操作x++的解释是：
- 先读取x中的值。
- 再2将访问到的值增加1.
- 最后将递增后的结果写回x中。
即读取-修改-写入。
如果我们现在在多个线程上执行递增操作，那么就会因为线程之间调度方式的不正确，使得最后的结果并不如意，因为读取-修改-写入这个操作必须是一次性的，不能被中断。因此为了产生“除非已经完成这三个操作，否则其他线程都不能读取或写入x的值”。这个操作(读取-修改-写入)被称之为原子操作。
原子操作带来了什么？
GPU的并行使得加速程序变得很容易，但是如果出现多个线程需要对同一地址进行读取写入，那么势必要有先后顺序，即串行性。虽然这样会引入竞，特别是线程比较多的时候，这个时候可以借助共享内存来缓解这个强竞争环境。
如何使用原子操作？
在官方文档把所有的原子操作API都写出来了，在GPU高性能这本书中举了atomicAdd()这个例子，以这个例子来说明原子操作如何使用。
__global__ void histo_kernel( unsigned char *buffer,
long size,
unsigned int *histo ) {
int i = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
while (i < size) {
atomicAdd( &histo[buffer[i]], 1 );
i += stride;
}
}
最重要的当然是：
atomicAdd(*dev, value)
- 函数调用这个API将会生成一个原子的操作序列，包括读取地址处的值，然后加上value的值，最后将这个值返回给地址处。
- 最重要的是，底层硬件将确保执行这些操作时，其他任何线程都不会干扰到这个原子操作，从而得到预计的结果。
通常这样都会产生一个问题：我们的优化项目一般都是有成百上千的线程，如果这么多线程访问少量的内存，原子操作会使得线程之间产生竞争，那么对相同地址的原子操作们会被硬件并行化，使得一帮线程都会在等待，失去了加速的效果。
主要原因：太多线程在少量内存山产生竞争。
解决办法：分散竞争，让少量的线程在少量的内存上进行原子操作，这样竞争就大大减少。一个办法就是使用共享内存，天然的绝缘线程联系的特性很适合原子操作中的竞争问题。因此，通过使用共享内存而不是全局内存，可以降低竞争。
CUDA流、固定内存
主机内存之页锁定内存(固定内存、不可分页内存)：
相对应的是CPU上的虚拟内存技术，可被分页到磁盘上。而固定内存是不可被分页到磁盘上的，因此其占用的是物理内存。
详细内容参考：
计算机内存和磁盘的关系
计算机一些内存分类
CUDA学习笔记之内存
固定内存的创建，CUDA提供自己独有的机制来分配主机内存。相比malloc()分配的是可分页内存，固定内存的分配则是由cudaHostAlloc()来做。
调用格式如：
int *a;
cudaHostAlloc((void **)&a, 8*sizeof(*a), cudaHostAllocDefault);
...
...
cudaFreeHost(a)
这里的flag=cudaHostAllocDefault表示默认值，用于分配不同形式的固定主机内存。这里可以理解为就是分配普通的固定内存。
固定内存的好处：
- 操作系统不会j将这部分内存交换到磁盘上，因此这块内存上的数据不会被重新定位。GPU知道内存的物理地址，因此可以通过直接内存访问(DMA，简单说就是GPU抢下系统总线的控制权。关于系统总线和系统前端总线：
)技术在GPU和主机之间复制数据。由于DMA不需要CPU的介入，故CPU可能在DMA的执行过程中将目标内存交换到磁盘上。CPU对可分页内存的移动会造成DMA的延迟。除此之外，DMA对可分页内存进行操作时，会先把可分页内容拷贝到临时页锁定内存，然后再从临时拷贝到GPU上。 - 固定内存的
访问速度比可分页内存更快，除了第1点中的以外，很重要的一点是因为每次从可分页内存执行复制操作时，复制速度受限于PCIE传输速度(CPU与GPU传输的协议)和Host系统前端总线速度较低的一方。 - 处理
CUDA流的时候，需要使用cudaMemcpyAsync()，这个复制的Host端内存必须是由cudaHostAlloc()分配的固定内存。
固定内存的坏处：
- 由于其一定是一直占用物理内存的(除非及时手动释放)，故一些需要内存的程序可能会因为Host内存不足而运行失败。
- 因此GPU高性能的作者建议：仅对cudaMemcpy的源内存或目标内存使用页锁定，并且在不需要时立即释放而不是等到main函数的结尾。
介绍完固定内存之后就可以引出CUDA流
作用：用于加速应用程序。
关于CUDA流：
- CUDA流表示一个GPU操作队列，队列中可以的操作可以是核函数启动、内存复制、事件的启动和结束…并以指定的顺序执行，所谓执行的顺序就是流操作中添加的顺序。
- 只有支持设备重叠(通过
prop.deviceOverlap可查询)才可以使用CUDA流。支持设备重叠的GPU可以在执行核函数的同时，还能在Host和Device端执行复制操作。故我们可以通过使用多个流来实现这种计算和数据传输的重叠。 - 将流视作GPU上一个任务，那么这些任务(流与流之间)就可以并行执行。
- CUDA只能保证在流中的各个操作必须是在上一个操作完成之后才会开始执行下个操作，比如第一个复制操作会在第二次复制操作开始之前执行；第二个复制操作会在核函数启动之前完成；第三次复制操作会在核函数执行完才开始执行。GPU依次从流中取操作执行。
- 从第4点看出在指定了流之后，核函数不仅是异步执行还是异步调用。那么核函数如何加入流中呢？举例如下：
add<<<128, 128, 0, stream0>>>();
add<<<128, 128, 0, stream1>>>();
- 比如你把流放在了for循环里，可能CPU在for循环结束之后，CPU接下来不等你了，去干别的指令去了，而在这个流里面复制、核函数可能没有一个执行完的。因此在队列中会包含许多等待GPU执行的工作。
CUDA流的使用
①流的创建：
cudaStearm_t stream0, stream1;
cudaStreamCreate(&stream0);
cudaStreamCreate(&stream1);
②cudaHostAlloc()与cudaMemcpyAsync()：
- 在需要使用流的时候，主机内存的分配必须得用
cudaHostAlloc()，因为复制操作必须得用cudaMemcpyAsync()，这是一个异步复制(cudaMemcpy是同步复制)并且可以处理流的API；并且这种异步复制使得CPU可以做更多的事。 - cudaMemcpyAsync()官方文档:
从形参可以看出，这个复制只是比cudaMemcpy多了一个流的控制，这个参数的意思就是表示在CUDA流中执行一次内存复制操作。并且需要注意其中的主机内存指针必须是通过cudaHostAlloc()分配的固定内存。
③同步CPU与GPU
cudaStreamSynchronize(stream0);
cudaStreamSynchronize(stream1);
并指定要等待的流。
④流的销毁
在结束流的部分之后，还要销毁CUDA流。
cudaStreamDestroy(stream0);
cudaStreamDestroy(stream1);
我们可以看出单个CUDA流就可以使得程序得到加速，但CUDA流真正厉害的地方在于多个流的协同调用。但是流与流之间操作的顺序是不能乱放的，如果放的不好的话，会导致一些操作得到阻塞，到头来并没有得到有效地加速。因此想要在重叠中获得加速，就必须遵从GPU的工作调度机制。
GPU调度机制:
程序员可以将流视为有序的操作序列，但硬件不知道，硬件只知道用一个或多个引擎来执行内存复制操作；另一个与之独立的引擎来执行核函数。引擎会对这些操作进行排队。程序员与硬件对于排队方式有着不同的理解，这时候就需要CUDA驱动程序负责在两者之间协调。
我们用2个例子来说明如何结合硬件将操作放入流中。
深度优先：
for (int i=0; i<FULL_DATA_SIZE; i+= N*2) {
// copy the locked memory to the device, async
HANDLE_ERROR( cudaMemcpyAsync( dev_a0, host_a+i,
N * sizeof(int),
cudaMemcpyHostToDevice,
stream0 ) );
HANDLE_ERROR( cudaMemcpyAsync( dev_b0, host_b+i,
N * sizeof(int),
cudaMemcpyHostToDevice,
stream0 ) );
kernel<<<N/256,256,0,stream0>>>( dev_a0, dev_b0, dev_c0 );
// copy the data from device to locked memory
HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c0,
N * sizeof(int),
cudaMemcpyDeviceToHost,
stream0 ) );
// copy the locked memory to the device, async
HANDLE_ERROR( cudaMemcpyAsync( dev_a1, host_a+i+N,
N * sizeof(int),
cudaMemcpyHostToDevice,
stream1 ) );
HANDLE_ERROR( cudaMemcpyAsync( dev_b1, host_b+i+N,
N * sizeof(int),
cudaMemcpyHostToDevice,
stream1 ) );
kernel<<<N/256,256,0,stream1>>>( dev_a1, dev_b1, dev_c1 );
// copy the data from device to locked memory
HANDLE_ERROR( cudaMemcpyAsync( host_c+i+N, dev_c1,
N * sizeof(int),
cudaMemcpyDeviceToHost,
stream1 ) );
}
硬件上的执行时间线：
从图中可以看出，驱动程序中操作的放入顺序就是先把第一个流的4个操作放完，然后放第二个流的，但是不仅浪费了2个时间段，而且没有利用流与流之间的独立性——调用一个流的时候无意间阻塞了另一个流。
正确的做法：宽度优先：
for (int i=0; i<FULL_DATA_SIZE; i+= N*2) {
// enqueue copies of a in stream0 and stream1
HANDLE_ERROR( cudaMemcpyAsync( dev_a0, host_a+i,
N * sizeof(int),
cudaMemcpyHostToDevice,
stream0 ) );
HANDLE_ERROR( cudaMemcpyAsync( dev_a1, host_a+i+N,
N * sizeof(int),
cudaMemcpyHostToDevice,
stream1 ) );
// enqueue copies of b in stream0 and stream1
HANDLE_ERROR( cudaMemcpyAsync( dev_b0, host_b+i,
N * sizeof(int),
cudaMemcpyHostToDevice,
stream0 ) );
HANDLE_ERROR( cudaMemcpyAsync( dev_b1, host_b+i+N,
N * sizeof(int),
cudaMemcpyHostToDevice,
stream1 ) );
// enqueue kernels in stream0 and stream1
kernel<<<N/256,256,0,stream0>>>( dev_a0, dev_b0, dev_c0 );
kernel<<<N/256,256,0,stream1>>>( dev_a1, dev_b1, dev_c1 );
// enqueue copies of c from device to locked memory
HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c0,
N * sizeof(int),
cudaMemcpyDeviceToHost,
stream0 ) );
HANDLE_ERROR( cudaMemcpyAsync( host_c+i+N, dev_c1,
N * sizeof(int),
cudaMemcpyDeviceToHost,
stream1 ) );
}
硬件上的时间线：
从图中可以看出，宽度优先使得总体的时间从8格减小到了6格。宽度优先就是利用2个独立性：①流与流之间的独立性。②引擎之间的独立性。
这里有2个疑问：
- 最大的疑问就是两个红色为何是同步执行的，其实就是利用了2个流之间的独立性，第0个流的复制B已经结束了，那么接下去就是第0个流的核函数了。
- 第二个疑问就是既然两个流独立，那么为何第0个流的复制B在第一个流的复制A后面：那是因为CUDA驱动程序会按照这些操作的顺序把他们调度到硬件上按指定顺序执行。而内存复制引擎只有一个，因此只能按先后放了。
CUDA流的加速总结：
- 这种将流与流之间操作的进行重叠，第0个流从GPU复制回CPU的时候，第2个流在执行核函数，高度并行化。
- 每个流内部有Kernels在执行，本身就是并行化实现。
- CUDA流采用的API都是异步执行的，CPU可以去干更多的事情。
- CUDA流采用的cudaHostAlloc来分配主机内存，其拥有更高的带宽。
零拷贝内存、多GPU
零拷贝内存：如果GPU要访问CPU上的数据，需要用到cudaMemcpy()将数据从CPU拷贝到GPU内存上，然后执行核函数解引内存上的数据。零拷贝内存就是一种不需要你拷贝这个步骤，在核函数中直接访问这种类型的主机内存(这里的“零”就是不需要拷贝的意思)。因此，有了零拷贝内存的存在，我们就可以通过GPU直接访问CPU的固定内存了。
零拷贝内存的使用：
cudaHostAlloc( (void**)&a,
size*sizeof(float),
cudaHostAllocWriteCombined |
cudaHostAllocMapped ) ;
Note：
- 首先需要判断设备是否支持映射：
prop.canMapHostMemory()。 - 拷贝内存的调用：用cudaHostAlloc()分配固定内存，指定
flag=cudaHostAllocMapped，意味着将CPU上指定的固定内存映射到GPU上，并从GPU中访问这块内存。 flag=cudaHostAllocWriteCombined表示将内存分配为WC内存，WC内存在PCIE总线上会有更快的传输速度。注意点①CPU内存映射到GPU后，由GPU访问会提升性能。②主机内存复制到GPU上。③不要让CPU去读取这块内存。④这个标志不会改变应用程序的功能，在需要提升GPU性能的时候可以使用。- GPU想要调用这块内存还需要被告知地址所在。故需要使用API：
cudaHostGetDevicePointer()来获取这块内存在GPU上的有效指针。官方文档
如下：
调用格式：
cudaHostGetDevicePointer( &dev_a, a, 0 );
和cudaMalloc一样，dev_a在核函数调用的时候是设备指针。这里a是主机指针。
- 映射前后的主设内存共享。
- 核函数在执行期间可能会修改零拷贝内存的数据，故在核函数结束之后，需要同步(调用
cudaThreadSynchronize())一下，以此来告知CPU，类似于cudaDeviceSynchronize()。
零拷贝内存的作用：
- 对于集显来说，CPU和GPU使用同一块内存，零拷贝内存的唯一作用就是避免了数据复制。
- 对于独显来说，对于只会使用一次的输入和输出内存而言，可以带来性能提升。这样可以减轻复制带来的PCIE读入写入的延迟。但是零拷贝内存这种隐藏内存访问也会有延迟，故不要重复访问太多次，否则还不如复制。
多GPU运算:
一个CPU线程控制着一个GPU，CPU不同线程也是并行的，因此实现多GPU需要主机进行线程管理。除此之外，每个GPU都需要调用：
HANDLE_ERROR( cudaSetDevice( data->deviceID ) );
HANDLE_ERROR( cudaSetDeviceFlags( cudaDeviceMapHost ) );
来告知GPU的编号，如0,1…。并且这种操作对于一个Host线程只能使用1次。
多GPU主要复杂在主机线程管理上，不是GPU的事，就不写了。
另外需要注意的是，在多GPU操作中，固定内存有些许的改变。我们之前单GPU时候说的固定内存在多GPU的时候是这样的：分配它的GPU将它视作固定内存，其余GPU访问这块内存时候将视作为可分页内存，麻烦的是，用户却使用固定内存的方式去使用它，就会出错。
解决方式：可移动的固定内存
将固定内存配置为可移动的，故可以在主机线程之间移动这块内存，并且每个CPU线程都将视之为固定内存。
如何配置呢？
使用cudaHostAlloc()，并配置flag=cudaHostAllocportable。并且通过“|”和cudaHostAllocMapped、cudaHostAllocWriteCombined共同使用。
规约运算(reduction)
是在GPU高并行化常用的操作，通过某种运算将大数组输出为小数组。假设1个256个元素的数组，C语言下相加需要256次操作。而在规约运算下，只需 log 256 \log256 log256次，节省了很多时间，具体的：
int i = blockDim.x/2;
while(i!=0)
{
if(cacheIndex<i)
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
i/=2;
}
常用的CUDA API
这里记录一下上面没有提及的常用的API：
①cudaMemset()：
相当于做了cudaMemcpy()的工作，这是一个给GPU指定内存填充value的API。比如:
cudaMemset(设备指针=dev, 填充的数=0, 填充区域大小=128*sizeof(int))
370
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
