1、CUDA性能测量
虽然可以通过CPU上的计时器来测量CUDA应用的性能。但这种方式并不会给出很精确的结果。CPU上的时间测量还需要CPU端具有高精度的定时器。很多时候，在GPU内核异步运行的同时，主机上正在执行着计算，所以CPU端的计时器可能无法给出正确的内核执行时间。
CUDA事件等价于CUDA应用运行的特定时刻被记录的时间戳。通过使用CUDA事件API，由GPU来记录这个时间戳，因此消除了CPU端的计时器测量性能时所会受到的影响。
使用CUDA测量时间需要两个步骤：
创建事件
记录事件
我们将会记录两个事件：
一个记录我们的代码运行开始的时刻；
另一个记录我们的代码运行结束的时刻。
接着我们会通过两个事件记录的时刻相减来计算出运行时间，这将会给出我们的代码整体性能的参考信息。
在CUDA代码中可以使用CUDA事件API导入以下代码来测量性能：
#include "stdio.h" #include<iostream> #include <cuda.h> #include <cuda_runtime.h> //循环次数 #define N50000 //定义kernel函数 __global__ void gpuAdd(int *d_a, int *d_b, int *d_c) { //Getting Thread index of current kernel int tid = threadIdx.x + blockIdx.x * blockDim.x; while (tid < N) { d_c[tid] = d_a[tid] + d_b[tid]; tid += blockDim.x * gridDim.x; } } int main(void) { //定义主机数组 int h_a[N], h_b[N], h_c[N]; //定义设备指针 int *d_a, *d_b, *d_c; cudaEvent_t e_start, e_stop; cudaEventCreate(&e_start); cudaEventCreate(&e_stop); cudaEventRecord(e_start, 0); // 申请显存 cudaMalloc((void**)&d_a, N * sizeof(int)); cudaMalloc((void**)&d_b, N * sizeof(int)); cudaMalloc((void**)&d_c, N * sizeof(int)); //初始化数组 for (int i = 0; i < N; i++) { h_a[i] = 2 * i*i; h_b[i] = i; } // 把数据传给device cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice); //调用kernel函数 gpuAdd << <512, 512 >> > (d_a, d_b, d_c); //把数据传递给host主机 cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost); cudaDeviceSynchronize(); cudaEventRecord(e_stop, 0); cudaEventSynchronize(e_stop); float elapsedTime; cudaEventElapsedTime(&elapsedTime, e_start, e_stop); printf("Time to add %d numbers: %3.1f ms\n", N, elapsedTime); int Correct = 1; printf("Vector addition on GPU \n"); for (int i = 0; i < N; i++) { if ((h_a[i] + h_b[i] != h_c[i])) { Correct = 0; } } if (Correct == 1) { printf("GPU has computed Sum Correctly\n"); } else { printf("There is an Error in GPU Computation\n"); } //释放内存 cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); return 0; }
建立两个事件对象，e_start和e_stop，用来测量一段代码的开始和结束区间（的执行性能）。通过cudaEvent_t类型来定义相关的变量，然后通过cudaEventCreate建立事件，保存在刚才定义好的变量上。我们通过将cudaEvent_t类型的变量地址传递给cudaEventCreate来得到对应的事件对象，而非cudaEventCreate通过返回值直接返回它。
在代码的开始，我们通过调用cudaEventRecord，用刚才建立的第一个事件对象e_start记录一次GPU时间戳。第二个参数我们指定为0，这里代表了CUDA流。
当这些代码结束后，我们会再记录一次时刻，这次是记录在事件对象e_stop里。这次通过cudaEvent-Record(e_stop,0)这行代码进行记录。一旦我们记录了开始的时刻和结束的时刻，两者的差值将会给出用来衡量你的代码性能的时间。但直接计算两个事件对象的差值时间的话，还存在问题。
CUDA C（代码）的执行可能是异步的。当GPU正在执行内核的时候，没等它执行完毕，CPU可能就已经继续执行后续的代码行了。同样的，类似内核，事件的记录命令从它在CPU上发出到它实际在GPU上执行也是异步的。所以，CPU和GPU不进行同步，直接认为发出了事件的记录命令后就可以立刻使用测量值的想法，可能会给出错误的结果。
因为cudaEventRecord()作用的事件对象将会在它之前所有发出GPU命令都执行完毕后才会记录一个时刻值。我们必须等待事件对象e_stop实际的已经完成了时刻值记录，也就是它之前的所有GPU工作都完成后才应该试图读取使用它。这样，对于异步的cuda-EventRecord()操作，等它完成后再访问才是安全的。
通过CUDA提供的cudaEventElapsedTime这个API函数来计算两个时间戳之间的差值。该API函数具有3个参数，第一个参数是用来返回时间差结果的，第二个参数是起始时刻的事件，第三个参数则是结束时刻的事件。用这个函数计算出时间后，我们在下一行代码从控制台上显示出来它。
GPU用来累加50000个元素的程序耗时0.5ms。这个输出和系统配置有关。
2、CUDA中错误处理
定义一个cudaError_t类型的变量用来保存CUDA函数的返回值。例如这里的代码就用cudaStatus=cudaMalloc()来保存它的返回值。如果该返回值不等于cudaSuccess，则说明设备上的显存分配出错了。然后我们用if语句处理错误情况。具体这里将错误信息显示到控制台上，并且跳转到程序结尾的Error标号处。cudaMemcpy也同样具有类似cudaMalloc的错误处理代码。
无论在（上述检测位置的）何处侦测到错误发生，我们都会跳转到（Error标号）这里。在这里，我们释放设备上的显存分配，并退出main函数。这是一种非常高效的编写CUDA代码的方式，我们建议你用这种方式来编写你的CUDA代码。虽然在我们之前的代码例子中说过要避免不必要的复杂性，而在CUDA程序中加上错误处理代码会让程序变得（很）长，但这样做能够指出代码里是什么操作造成了问题。
#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h> __global__ void gpuAdd(int *d_a, int *d_b, int *d_c) { *d_c = *d_a + *d_b; } int main() { //Defining host variables int h_a, h_b, h_c; //Defining Device Pointers int *d_a, *d_b, *d_c; //Initializing host variables h_a = 1; h_b = 4; cudaError_t cudaStatus; // Allocate GPU buffers for three vectors (two input, one output) . cudaStatus = cudaMalloc((void**)&d_c, sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&d_a, sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&d_b, sizeof(int)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } // Copy input vectors from host memory to GPU buffers. cudaStatus = cudaMemcpy(d_a, &h_a, sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } cudaStatus = cudaMemcpy(d_b, &h_b, sizeof(int), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } // Launch a kernel on the GPU with one thread for each element. gpuAdd << <1, 1 >> > (d_a, d_b, d_c); // Check for any errors launching the kernel cudaStatus = cudaGetLastError(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus)); goto Error; } // Copy output vector from GPU buffer to host memory. cudaStatus = cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } printf("Passing Parameter by Reference Output: %d + %d = %d\n", h_a, h_b, h_c); Error: cudaFree(d_c); cudaFree(d_a); cudaFree(d_b); return 0; }
3、CUDA程序的性能优化
1、使用适当的块数量和线程数量
启动内核的时候，需要指定两个参数：（Grid里）的块数量，和块里的线程数量。在内核执行期间，GPU资源不应当存在空闲，只有这样才能给出较优性能。如果存在空闲的资源，则可能会降低应用程序的性能。合适的块和线程数量有助于让GPU资源保持充分忙碌。
研究表明：如果块数量是GPU的流多处理器数量的两倍，则会给出最佳性能，不过，块和线程的数量和具体的算法实现有关。GPU的流多处理器数量则可以通过第2章中获取设备属性的方法取得。有人认为块中的线程数量应当被设定等于设备属性中每个块所能支持的最大线程数量，但实际上这些数值只是作为一种基本的准则来说的。你可以适当微调这些数值，来取得你的程序的优化性能。需要反复实现，试探可能的形状组合。
2、最大化数学运算效率
数学运算效率的定义是，数学运算操作和访存操作的比率。但我们不能认为直接通过最大化每个线程的运算量和最小化访存时间就可以取得最好的数学运算效率，就可以提升性能。常见的内核执行有3个瓶颈：卡在计算瓶颈上，卡在访存上和卡在延迟掩盖上。对于特定的内核，如果卡在计算上，则应当考虑将一些计算等效地转换成访存，例如一些运算可以尝试转换成存储器查表；而卡在访存上，则可以将一些访存转换成对应的计算，例如一些数据不是重新载入，而是直接计算出来。这需要检查具体代码，在具体显卡上通过Profiler分析。哪种资源先达到瓶颈，就减少这种资源的使用（计算或者访存），而增加另外一种，并非一味地增加计算，或者减少访存。
缓存的使用也有助于减少存储器访问时间，最终一定程度地辅助达成减少（内核的）全局内存的带宽需求就能减少花费在访存上的时间的目的。高效地使用存储器对提升CUDA程序性能非常重要，当显存带宽是瓶颈的时候，减少带宽需求有助于提升性能。
3、使用合并的或跨步式的访存
合并访存大致上意味着线程束（warp）整体读取或者写入连续的存储器区域。这种对存储器的访问对GPU来说是最高效的。如果warp中的线程固定步长地离散式访问某段存储器区域，则这叫跨步式访存。跨步式访存的效果不如合并访存好，但依然比随机访存要好。所以，如果你尝试在程序中使用合并访存的话，它有时会对提升性能有帮助。
4、避免warp内分支
当warp内的线程发生了分别转向执行不同的代码路径的时候，我们叫它warp内分支。它可能发生在下面的内核代码场景中：
在第一个代码片段中，因为if语句的判断，奇数和偶数ID的线程将会分别执行不同的代码。在GPU上，特别是计算能力小于7.0的卡上，一个warp中的32个线程总是同步伐的执行，所有的同一个warp内的线程，都必须执行相同的指令。所以对于这个例子来说，偶数线程的路径和奇数线程的路径都会被warp分别执行一遍，这就造成了性能损失。
在第二个代码片段中，使用for循环，每个线程都以不同的迭代次数运行for循环，因此所有线程将花费不同的时间完成。因为warp的同步执行机制，对于同一个warp中的线程来说，整体执行时间以最长的那个线程为准，所以warp内其他线程造成了时间浪费，影响性能。
5、使用锁定页面的内存
在之前的所有例子中，我们都是用malloc函数在CPU上分配内存，该函数分配的是可换页的标准内存。CUDA提供了另外一个叫作cudaHostAlloc的API函数，该函数分配的是锁定页面的内存。这种内存也叫Pinned内存。操作系统会保证永远不会将这种内存换页到磁盘上，总是在物理内存中。所以，系统内的所有设备都可以直接用该段内存缓冲区的物理地址来访问。此属性帮助GPU通过直接内存访问（DMA）将数据复制到主机或从主机复制数据，而无需CPU干预。但是锁定页面的内存应当正确地使用，不能使用过多，因为这种内存不能被换页到磁盘上，分配的过多，你的系统可能会物理内存不足，从而其他在这个系统上运行的应用程序可能会受到影响。你可以通过该API来分配适合（高效）传输的内存。使用该API函数的语法如下：
cudaHostAlloc函数的语法类似普通的malloc函数。注意cudaHostAlloc的第三个参数如果不是指定为cudaHostAllocDefault，则可以用来调节分配到的锁定页面的内存的属性。
cudaFreeHost函数用来释放通过cudaHostAlloc分配到的内存。
