【GPU高性能编程 CUDA实战】学习笔记

第1章 为什么需要CUDA

在2000年早期，GPU的主要目标都是通过可编程计算单元为屏幕上的每个像素计算出一个颜色值，这些计算单元也称为像素着色器（Pixel Shader）。

第2章 入门

开发环境：

• 支持CUDA的图形处理器
• NVIDIA设备驱动程序
• CUDA开发工具箱
• 标准C编译器

第3章 CUDA C

• 了解为**主机（Host）编写的代码与为设备（Device）**编写的代码之间的区别。

• 如何从主机上运行设备代码

• 了解如何在支持CUDA的设备上使用设备内存

• 一个空的函数kernel()，并且带有修饰符_global_。

• 对这个空函数的调用，并且带有修饰符<<<1,1>>>。

• 可以像调用C函数那样将参数传递给核函数

• 当设备执行任何有用的操作时，都需要分配内存，例如将计算值返回给主机。

设备信息，被课程全在NVIDIA 750Ti下测试。

   --- General Information for device 0 ---
Name:  NVIDIA GeForce GTX 750 Ti
Compute capability:  5.0
Clock rate:  1084500
Device copy overlap:  Enabled
Kernel execution timeout :  Enabled
   --- Memory Information for device 0 ---
Total global mem:  2097414144
Total constant Mem:  65536
Max mem pitch:  2147483647
Texture Alignment:  512
   --- MP Information for device 0 ---
Multiprocessor count:  5
Shared mem per mp:  49152
Registers per mp:  65536
Threads in warp:  32
Max threads per block:  1024
Max thread dimensions:  (1024, 1024, 64)
Max grid dimensions:  (2147483647, 65535, 65535)

第4章 CUDA C并行编程

<<<[线程块数量], [每个线程块中线程数量]>>>

blockIdx，一个内置变量，描述线程块的编号，即上面三尖括号中的第一个参数

gridDim，也是内置变量

__global__，启动，

__device__，在设备（GPU）上运行的代码，只能从其他__device__或者__global__函数调用它们

第5章 线程协作

• 了解CUDA C中的线程
• 了解不同线程之间的通信机制
• 了解并行执行线程的同步机制

<<<[线程块数量], [每个线程块中线程数量]>>>

CUDA程序中如何计算线程号：

[线程块数量] dim3 dimGrid([], [])

[每个线程块中线程数量] dim3 dimBlock([], [])

线程块用blockIdx标识，并且是列优先的；线程块中的线程用threadIdx标识，也是列优先的。

线程块的维度用gridDim标识，单个线程块内线程的维度用blockDim标识。

1. 使用N个线程块，每个线程块只有一个线程，即

   dim3 dimGrid(N);
   dim3 dimBlock(1);
   threadID = blockIdx.x;	// 找到线程块，就是找到了线程的编号
   

2. 使用$M\times N$个线程块，每个线程块一个线程，显然是列优先

   dim3 dimGrid(M, N);
   dim3 dimBlock(1);
   threadID = blockIdx.y * gridDim.x + blockIdx.x;	// gridDim.x = M
   

3. 使用一个线程块，该线程块具有N个线程

   dim3 dimGrid(1);
   dim3 dimBlock(N);
   threadID = threadIdx.x;
   

4. 使用M个线程块，每个线程块内有N个线程，

   dim3 dimGrid(M);
   dim3 dimBlock(N);
   threadID = blockIdx.x * blockDim.x + threadIdx.x;	// blockDim.x = N
   

5. 使用$M\times N$的二维线程块，每个线程块中有$P\times Q$个线程，索引有两个维度

   dim3 dimGrid(M, N);
   dim3 dimBlock(P, Q);
   blockID = blockIdx.y * gridDim.x + blockIdx.x;	// gridDim.x = M, 找到属于哪个线程块
   threadID = threadIdx.y * blockDim.x + threadIdx.x;		// blockDim.x = P,找到在当前线程块中的编号
   idx = blockID * blockDim.x * blockDim.y + threadID;
   

总结，对于一个三维网格，和一个三维的线程块

<<<[dim3 grid], [dim3 block]>>>

总的线程数量$\mathrm{N}$为：
$$\mathrm{N}=\mathrm{g}\mathrm{r}\mathrm{i}\mathrm{d}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{x}\times \mathrm{g}\mathrm{r}\mathrm{i}\mathrm{d}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{y}\times \mathrm{g}\mathrm{r}\mathrm{i}\mathrm{d}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{z}\times \mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{x}\times \mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{y}\times \mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{z}$$
线程标号按照下面的顺序求：

1. 先找到当前线程位于那一个线程块中
   $$\mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{I}\mathrm{D}=\mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{I}\mathrm{d}\mathrm{x}.\mathrm{x}+\mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{I}\mathrm{d}\mathrm{x}.\mathrm{y}\times \mathrm{g}\mathrm{r}\mathrm{i}\mathrm{d}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{x}+\mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{I}\mathrm{d}\mathrm{x}.\mathrm{z}\times \mathrm{g}\mathrm{r}\mathrm{i}\mathrm{d}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{x}\times \mathrm{g}\mathrm{r}\mathrm{i}\mathrm{d}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{y}$$

2. 找到当前线程位于当前线程块中的位置
   $$\mathrm{t}\mathrm{h}\mathrm{r}\mathrm{e}\mathrm{a}\mathrm{d}\mathrm{I}\mathrm{D}=\mathrm{t}\mathrm{h}\mathrm{r}\mathrm{e}\mathrm{a}\mathrm{d}\mathrm{I}\mathrm{d}\mathrm{x}.\mathrm{x}+\mathrm{t}\mathrm{h}\mathrm{r}\mathrm{e}\mathrm{a}\mathrm{d}\mathrm{I}\mathrm{d}\mathrm{x}.\mathrm{y}\times \mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{x}+\mathrm{t}\mathrm{h}\mathrm{r}\mathrm{e}\mathrm{a}\mathrm{d}\mathrm{I}\mathrm{d}\mathrm{x}.\mathrm{z}\times \mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{x}\times \mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{y}$$

3. 计算一个线程块中一共有多少线程
   $$\mathrm{M}=\mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{x}\times \mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{y}\times \mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{D}\mathrm{i}\mathrm{m}.\mathrm{z}$$

4. 求得当前的线程序列号idx
   $$\mathrm{i}\mathrm{d}\mathrm{x}=\mathrm{t}\mathrm{h}\mathrm{r}\mathrm{e}\mathrm{a}\mathrm{d}\mathrm{I}\mathrm{D}+\mathrm{M}\times \mathrm{b}\mathrm{l}\mathrm{o}\mathrm{c}\mathrm{k}\mathrm{I}\mathrm{D}$$

gridDim, blockDim, blockIdx, threadIdx都是内置常量，一旦kernel启动，他们就是确定的了

__shared__ 声明一个驻留在共享内存中的变量，注意，是线程块中的线程共享，不同线程块不能共享。

// 声明一个驻留在共享内存中的变量，注意，是线程块中的线程共享，不同线程块不能共享。
__shared__ float cache[threadPerBlock];
// 对线程块中的线程进行同步
__syncthreads();

__syncthreads(); 也和MPI的**MPI_Barrier()**一样，需要组内所有的线程（进程）都执行到这条指令，才会继续运行，所以要谨慎的将其加入条件语句中，最好是不加。

对一个输入数组执行某种计算，然后产生一个更小的结果数组，这种过程也称为归约（Reduction）。

第6章 常量内存与事件

常量内存（Constant Memory）

通过事件来测量CUDA应用程序的性能。通过这些测量放，你可以定量地分析对应用程序某个修改是否会带来性能提升（或者性能下降）

• 了解如何在CUDA C中使用常量内存
• 了解常量内存的性能特性
• 学习如何使用CUDA事件来测量应用程序的性能。

常量内存（Constant Memory） __constant__

• 申请的是device内存，cudaMemcpyToSymbol拷贝就是从host拷贝到global memory。

• 申请的是constant内存，cudaMemcpyToSymbol拷贝就是从host拷贝到constant memory。

与从全局内存中读取数据相比，从常量内存中读取相同的数据可以节约内存带宽，主要有两个方面：

• 对常量内存的单词读操作可以广播到其他了“邻近（Nearby）”线程，将节约15次读取操作。
• 常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。

线程束warp，Warp可以看成是一组线程通过交织而形成的一个整体。在CUDA架构中，Warp是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致（Lockstep）“的形式执行。在程序的每一行，Warp中的每个线程都将在不同的数据上执行相同的指令。

当处理常量内存时，NVIDIA硬件把单次内存读取操作广播到每个半线程束（Half-Warp）。在半线程束中包含了16个线程，即线程束中线程数量的一半。如果在半线程束中的每个线程都从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量的数据，那么这种方法产出的内存流量只是使用全局内存时的1/16（大约6%）。

由于这块内存的内容不会改变，因此硬件将主动将这个常量数据缓存在GPU上。因此只有第一次读取会产生内存流量，后面都会命中缓存。这将进一步减少额外的内存流量。

负面影响：当半线程束同时读相同地址时，这个功能可以极大提升性能，但是当所有16个线程分别读取不同的地址时，他实际上会降低性能

使用事件来测量性能

为了测量GPU在某个任务上花费的时间，我们将使用CUDA的事件API。

CUDA中的事件本质上是一个GPU时间戳。获得时间戳的两个步骤：首先创建一个事件，然后记录一个事件。

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);		// 第二个参数在讨论Stream（流）时再解释
// 在GPU上执行一些工作
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);	// 运行时，阻塞之后的语句，知道GPU执行stop事件。当返回时，我们知道stop事件之前的所有GPU工作都已完成，可以读stop中保存的值了。
float elapsedTime;	// 单位 ms
cudaEventElapsedTime(&elapsedTime, start, stop);	
cudaEventDestroy(start);
cudaEventDestroy(stop);

值得注意的是，由于CUDA事件是直接在GPU上实现的，因此他们不适用于对同时包含设备代码和主机代码的混合代码的计时。也就是说，如果试图通过CUDA事件对核函数和设备内存复制之外的代码进行计时，将得到不可靠的结果。

第7章 纹理内存

纹理内存（Texture Memory），与常量内存一样，是另一种类型的制度内存，在特定的访问模式中，纹理内存同样能够提升性能并减少内存流量。

• 了解纹理内存的性能特性
• 了解如何在CUDA C中使用一维、二维纹理内存

纹理缓存是专门为那些在内存访问模式中存在大量空间局部性（Spatial Locality）的图形应用程序而设计的。

首先，将输入的数据声明为texture类型的引用。

texture<float> texConstSrc;
texture<float> texIn;
texture<float> texOut;

在为这三个缓冲区分配了GPU内存后，需要通过cudaBindTexture()将这些变量绑定到内存缓冲区，这等于告诉CUDA运行时两件事：

• 我们希望将指定的缓冲区作为纹理来使用
• 我们希望将纹理引用作为纹理的名字

HANDLE_ERROR( cudaBindTexture( NULL, texConstSrc,
                               data.dev_constSrc,
                               imageSize ) );

tex1Dfetch()，两个参数，第一参数表明缓存位置，第二参数表明取缓存中的哪个元素。这个虽然看上去像是一个函数，但它其实是编译器内置函数(Intrinsic)。由于纹理引用必须声明为文件作用域内的全局变量，因此不再将输入缓冲区和输出缓存区作为参数传递给blend_kernel()，因为编译器需要在编译时知道tex1Dfetch()应该对哪些纹理采样。

// this kernel takes in a 2-d array of floats
// it updates the value-of-interest by a scaled value based
// on itself and its nearest neighbors
__global__ void blend_kernel( float *dst,
                              bool dstOut ) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    int left = offset - 1;
    int right = offset + 1;
    if (x == 0)   left++;
    if (x == DIM-1) right--; 

    int top = offset - DIM;
    int bottom = offset + DIM;
    if (y == 0)   top += DIM;
    if (y == DIM-1) bottom -= DIM;

    float   t, l, c, r, b;
    if (dstOut) {
        t = tex1Dfetch(texIn,top);
        l = tex1Dfetch(texIn,left);
        c = tex1Dfetch(texIn,offset);
        r = tex1Dfetch(texIn,right);
        b = tex1Dfetch(texIn,bottom);

    } else {
        t = tex1Dfetch(texOut,top);
        l = tex1Dfetch(texOut,left);
        c = tex1Dfetch(texOut,offset);
        r = tex1Dfetch(texOut,right);
        b = tex1Dfetch(texOut,bottom);
    }
    dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
}

清理工作，不仅要释放全局缓冲区，还需清除与纹理的绑定

// clean up memory allocated on the GPU
void anim_exit( DataBlock *d ) {
    cudaUnbindTexture( texIn );	 // 取消纹理的绑定
    cudaUnbindTexture( texOut );
    cudaUnbindTexture( texConstSrc );
    HANDLE_ERROR( cudaFree( d->dev_inSrc ) );
    HANDLE_ERROR( cudaFree( d->dev_outSrc ) );
    HANDLE_ERROR( cudaFree( d->dev_constSrc ) );

    HANDLE_ERROR( cudaEventDestroy( d->start ) );
    HANDLE_ERROR( cudaEventDestroy( d->stop ) );
}

二维纹理

tex2D()，可以直接通过x，y坐标来访问纹理，而且不用担心溢出的问题，如果x小于0，那么返回0处的值，大于宽度，返回宽度处的值，y同理。

绑定二维纹理时，CUDA运行时要求提供一个cudaCreateChannelDesc。通道格式描述符（Channel Format Description）的声明。

cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
HANDLE_ERROR( cudaBindTexture2D( NULL, texConstSrc,
                                data.dev_constSrc,
                                desc, DIM, DIM,	// 纹理维数
                                sizeof(float) * DIM ) );

一维二维同样的函数取消绑定。

第8章 图形互操作性

能否在同一应用程序中GPU既执行渲染运算，又执行通用计算？如果要渲染的图像依赖通用计算的结果，那么该如何处理？或者，如果想要在已经渲染的帧上执行某种图像处理或者同级，又该如何实现？

假设已经具备一些其他的技术背景知识，因为在示例代码中包含了大量的OpenGL和GLUT（OpenGL Utility Toolkit）代码，但是没有对其解释。

这章暂时搁置，里面用的图形库没用过，不懂，目前重点不在这。

第9章 原子性

• 了解不同NVIDIA GPU的计算功能集
• 了解原子操作以及为什么需要他们
• 了解如何在CUDA C核函数中执行带有原子操作的运算

1. 使用全局内存原子操作的直方图核函数

cudaMemset()和memset()的行为基本相同。

原子操作atomicAdd( addr, y);，包括读取addr处的值，将y加到这个值，然后结果保存回addr

atomicAdd( &(histo[buffer[i]]), 1);

由于在核函数中只包含了非常少的计算工作，因此很可能是全局内存上的原子操作导致了性能的降低。当数千个线程尝试访问少量的内存位置时，将发生大量的竞争。为了确保递增操作的原子性，对相同内存位置的操作都将被硬件串行化。这可能导致保存未完成操作的队列非常长，因此会抵消通过并行运行线程而获得的性能提升。

2. 使用共享内存原子操作和全局内存原子操作的直方图核函数

也就是先共享内存算好一部分的，最后给加到全局的上面。

第10章 流

**任务并行性（Task Parallelism）**是指并行执行两个或多个不同的任务，而并不是在大量数据上执行同一个任务。

• 了解如何分配**页锁定（Page-Locked）**类型的主机内存
• 了解CUDA流的概念
• 了解如何使用CUDA流来加速应用程序

cudaHostAlloc()分配页锁定的主机内存，页锁定内存也称为**固定内存（Pinned Memory）**或者不可分页内存，操作系统将不会对这块内存分页交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此，操作系统能够安全地使某个应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位。

由于GPU知道物理地址，那么可用DMA来复制数据。

cudaHostAlloc( (void**)&a, size * sizeof( *a ), cudaHostAllocDefault );
cudaFreeHost( a );
// up 从主机到设备，	down 从设备到主机
// 使用页锁定内存有四五倍的加速
Time using cudaMalloc:  10072.8 ms
        MB/s during copy up:  2541.5
Time using cudaMalloc:  14715.3 ms
        MB/s during copy down:  1739.7
Time using cudaHostAlloc:  2323.9 ms
        MB/s during copy up:  11016.1
Time using cudaHostAlloc:  2041.3 ms
        MB/s during copy down:  12541.0

cudaEventRecord(cudaEvent_t, stream)第二个参数用于指定插入事件的流（Stream）。

CUDA流在加速应用程序方面起着重要作用。CUDA流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。我们可以在流中添加一些操作，例如核函数启动、内存复制，以及事件的启动和结束等。将这些操作添加到流的顺讯也就是他们执行的顺序。你可以将每个流视为GPU上的一个任务，并且这些任务可以并行执行。

支持**设备重叠（Device Overlap）**功能的GPU能够在执行一个CUDA C核函数的同时，还能在设备与主机之间执行复制操作。

int main(){
	cudaStream_t    stream;
	// initialize the stream
	cudaStreamCreate( &stream );
    
	/****** 初始化工作 ************/
    
    // now loop over full data, in bite-sized chunks
    for (int i=0; i<FULL_DATA_SIZE; i+= N) {
        // copy the locked memory to the device, async
        // 用于在GPU与主机之间复制数据, 没用之前的cudaMemcpy(),cudaMemcpy()同步的，返回时，复制操作便已经完成，并且在输出缓存区中包含了刚复制进去的内容
        HANDLE_ERROR( cudaMemcpyAsync( dev_a, host_a+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_b, host_b+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream ) );

        kernel<<<N/256,256,0,stream>>>( dev_a, dev_b, dev_c );

        // copy the data from device to locked memory
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c,
                                       N * sizeof(int),
                                       cudaMemcpyDeviceToHost,
                                       stream ) );

    }
}

cudaMemcoyAsync()用于在GPU与主机之间复制数据。没用之前的cudaMemcpy()，cudaMemcpy()同步的，返回时，复制操作便已经完成，并且在输出缓存区中包含了刚复制进去的内容。cudaMemcoyAsync()是异步的，调用时，只是放置一个请求，表示在流中执行一次内存复制操作，这个流是通过参数stream来指定的。当函数返回时，我们无法确保复制操作是否已经启动，更无法保证它是否已经结束。**我们能够得到的保证是：复制操作肯定会在下一个被放入流中的操作之前执行。**任何传递给cudaMemcoyAsync()的主机内存指针都必须已经通过cudaHostAlloc()分配，也就是说，你只能以异步方式对页锁定内存进行复制操作（注解：这里翻译应该不大准确，从异步操作来看，这块内存不应该被置换的，那么这句话应该这么说：异步方式的复制只能对页锁定内存操作）。

核函数的尖括号，现在我们知道他能有四个参数了。

kernel<<<[gridDim3], [blockDim3], [stream number], [cudaStream_t]>>>();

从上面的流来看，很像MPI中的非阻塞通信。所以我猜测，stream这个参数中，肯定包括标识流中的操作是否完成的标志。

果然有用于同步的函数。cudaStreamSynchronize([stream])

// copy result chunk from locked to full buffer
cudaStreamSynchronize( stream );

当要确保相互独立的流能够真正地并行执行时，我们自己要起到一定的作用。记住，硬件在处理内存复制和核函数执行时分别采用了不同的引擎，因此我们需要知道，将操作放入流中队列中的顺序将影响着CUDA驱动程序调度这些操作以及执行的方式。

高效地使用多个CUDA流

如果同时调度某个流的所有操作，那么很容易在无意中阻塞另一个流的复制操作或者核函数的执行。要解决这个问题，在将操作放入流的队列时应该采用宽度优先方式，而非深度优先方式。

    // now loop over full data, in bite-sized chunks
    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 ) );
    }

第11章 多GPU系统上的CUDA C

• 了解如何分配和使用零拷贝内存（Zero-Copy Memory）
• 了解如何在同一个应用程序中使用多个GPU
• 了解如何分配和使用可移动的固定内存（Portable Pinned Memory）

cudaHostAlloc()的cudaHostAllocDefault参数来获得默认的固定内存。本章会介绍除此之外的其他参数值。cudaHostAllocMapped分配的主机内存也是固定的。可以在CUDA C核函数中直接访问这种类型的主机内存。由于这种内存不需要复制到GPU，因此也称为零拷贝内存。

矢量点积运算，cudaMalloc()和cudaHostAlloc(, cudaHostAllocMapped)使用，运算时间的对比。

// allocate the memory on the CPU
cudaHostAlloc( (void**)&a,size*sizeof(float),cudaHostAllocWriteCombined | cudaHostAllocMapped );
cudaHostAlloc( (void**)&b,size*sizeof(float),cudaHostAllocWriteCombined | cudaHostAllocMapped );
cudaHostAlloc( (void**)&partial_c,blocksPerGrid*sizeof(float),cudaHostAllocMapped );
                              
Value calculated:  27621697910970467221504.000000
Time using cudaMalloc:  147.3 ms
Value calculated:  27621697910970467221504.000000
Time using cudaHostAlloc:  24.8 ms

cudaHostAllocWriteCombined 表示运行时应该将内存分配为**“合并式写入（Write-Combined）”**内存。这个标志不会改变应用程序的功能，但却可以显著地提升GPU读取内存时的性能。然而，当CPU也要读取这块内存时，会显得低效，因此使用之前，必须首先考虑应用程序可能的访问模式。

cudaHostAlloc()返回CPU上的指针，需要cudaHostGetDevicePointer()来获得这块内存在GPU上的有效指针。

// find out the GPU pointers
cudaHostGetDevicePointer( &dev_a, a, 0 );
cudaHostGetDevicePointer( &dev_b, b, 0 );
cudaHostGetDevicePointer( &dev_partial_c, partial_c, 0 );

cudaThreadSynchronize()将CPU与GPU同步。

当输入内存和输出内存都只使用一次时，那么在独立GPU上使用零拷贝内存将带来性能提升。由于GPU在设计时考虑了隐藏内存访问带来的延迟，因此这种机制在某种程度上将减轻PCIe总线上读取和写入等操作的延迟，从而会带来可观的性能提升。但由于GPU不会缓存零拷贝内存的内容，如果多次读取内存，那么最终会得不偿失，还不如一开始就将数据复制到GPU。

集成GPU在物理上是与CPU共享内存的，那么将缓冲区声明为零拷贝内存的作用就是避免不必要的数据复制。

零拷贝内存也是一种固定内存（页锁定），每个固定内存都是占用系统的可用物理内存，并且不会被交换到磁盘中，这会降低系统的性能。

固定内存只对于分配它的线程是页锁定的，对于其他线程而言，就是一块普通内存对待，所以不能用cudaMemcpyAsync()，且不能放入流中，而使用cudaMemcpy()这种速率大约为最高传输速率的50%。

cudaHostAlloc(, cudaHostAllocPortable)固定内存分配为可移动的，cudaFreeHost()

第12章 后记

• CUDA 工具箱

  • CUFFT
  • CUBLAS，线性代数库，其中包含了著名的基本线性代数子程序(Basic LInear Algebra Subprograms, BLAS)，列优先

• NVIDIA GPU Computing SDK

  包含许多GPU计算示例程序

  • CUDA基本主题
  • CUDA高级主题
  • CUDA系统集成
  • 数据并行算法
  • 图形互操作
  • 纹理
  • 性能策略
  • 线性代数
  • 图形/视频处理
  • 计算金融
  • 数据压缩
  • 物理模拟

• NVIDIA性能原语（NVIDIA Performance Primitives，NPP）。

• 调试CUDA C

  • CUDA-GDB，除了调试器功能，还提供了CUDA内存检查器（CUDA Memory Checker）
  • NVIDIA Parallel Nsight，集成在VS中的GPU/CPU调试器

• CUDA Visual Profiler

  可视化分析工具来运行核函数，根据profiling调优。写出真正的高性能计算程序

• 参考资料

  • 《Programming Massively Parallel Processors: a Hands-On Approach》，2012年出的第二版，国内有翻译的了，目前最新是2022年的第四版，目前国内还没有翻译版本，第三版也没有翻译版本。伊利诺伊大学的课程可一起看，我看过这本第二版的翻译版，中间更多的讲思想与原理。

• CUDA U

  大学课程，最好的之一是伊利诺伊大学的课程

• NVIDIA 论坛

• 代码资源

  • CUDA 数据并行原语库（CUDA Data Parallel Primitives Library，CUDPP），这些原语为许多数据并行算法提供了重要基础，包括排序、流压缩、构建数据结构以及其他并行算法等等，如果你正在编写某个算法，CUDPP可能已经提供了这个算法或这个算法的大部分功能
  • 语言封装器

附录 高级原子操作

原子锁

#ifndef __LOCK_H__
#define __LOCK_H__

struct Lock {
    int *mutex;
    Lock( void ) {
        HANDLE_ERROR( cudaMalloc( (void**)&mutex, sizeof(int) ) );
        HANDLE_ERROR( cudaMemset( mutex, 0, sizeof(int) ) );
    }
    ~Lock( void ) {
        cudaFree( mutex );
    }
    __device__ void lock( void ) {
        while( atomicCAS( mutex, 0, 1 ) != 0 );
	__threadfence();
    }
    __device__ void unlock( void ) {
        __threadfence();
        atomicExch( mutex, 0 );
    }
};

#endif

原子锁和CPP中学到的差不多，可以迁移过来使用。

实现散列表

多线程散列表，要考虑多个线程同时将节点加入散列表中的同一个位置。

GPU中散列表的构造还是使用**Lock**