CUDA入门（六）异步并行执行解析-白红宇

为了更好地压榨GPU和CPU，很多时候都使用异步并行的方法让主机端与设备端并行执行，即控制在设备没有完成任务请求前就被返回给主机端。

异步执行的意义在于：首先，处于同一数据流内的计算与数据拷贝都是依次进行的，但一个流内的计算可以和另一个流的数据传输同时进行，因此通过异步执行就能够使GPU中的执行单元与存储器单元同时工作，更好地压榨GPU的性能；其次，当GPU在进行计算或者数据传输时就返回给主机线程，主机不必等待GPU运行完毕就可以进行进行一些计算，更好地压榨了CPU的性能。

之前，单单使用GPU的同步函数，会使设备在完成请求任务前，不会返回主机线程，主机线程将进入让步（yield）、阻滞（block）或者自旋（spin）。

下面一个SDK简单的异步并行执行的程序：

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include 
     
      #include 
      
       #include
       
        #include"helper_timer.h"__global__ void increment_kernel(int *g_data, int inc_value){    int idx = blockIdx.x*blockDim.x + threadIdx.x;    g_data[idx] = g_data[idx] + inc_value;}int correct_output(int *data, const int n, const int x){    for (int i = 0; i < n; i++)        if (data[i] != x)            return 0;}int main(int argc, char *argv[]){    cudaDeviceProp deviceProps;    cudaGetDeviceProperties(&deviceProps, 0);    printf("CUDA device [%s]\n", deviceProps.name);    int n = 16 * 1024 * 1024;    int nbytes = n*sizeof(int);    int value = 26;    int *a = 0;    cudaMallocHost((void**)&a, nbytes);    memset(a, 0, nbytes);    int *d_a = 0;    cudaMalloc((void**)&d_a, nbytes);    cudaMemset(d_a, 255, nbytes);    dim3 threads = dim3(512, 1);    dim3 blocks = dim3(n / threads.x, 1);    cudaEvent_t startGPU, stopGPU;    cudaEventCreate(&startGPU);    cudaEventCreate(&stopGPU);    StopWatchInterface *timer = NULL;    sdkCreateTimer(&timer);    sdkResetTimer(&timer);    cudaThreadSynchronize();    float gpu_time = 0.0f;    sdkStartTimer(&timer);    cudaEventRecord(startGPU, 0);    cudaMemcpyAsync(d_a, a, nbytes, cudaMemcpyHostToDevice, 0);    increment_kernel << 
        
         > >(d_a, value);    cudaMemcpyAsync(a, d_a, nbytes, cudaMemcpyDeviceToHost, 0);    cudaEventRecord(stopGPU, 0);    sdkStopTimer(&timer);    unsigned long int counter = 0;    while (cudaEventQuery(stopGPU) == cudaErrorNotReady)    {        counter++;    }    cudaEventElapsedTime(&gpu_time, startGPU, stopGPU);    printf("time spent executing by the GPU:%.2f\n", gpu_time);    printf("time spent by CPU in CUDA calls:%.8f\n", sdkGetTimerValue(&timer));    printf("GPU execute %d iteration while waiting forGPU to finish\n", counter);    printf("-------------------------------------------------------------------------\n");    if (correct_output(a, n, value))        printf("TEST PASSED\n");    else        printf("Test FAILED\n");    cudaEventDestroy(startGPU);    cudaEventDestroy(stopGPU);    cudaFreeHost(a);    cudaFree(d_a);    getchar();    cudaThreadExit();    return 0;}

执行结果：

代码中

GPU上执行的是对每一个数据进行加26的处理，在CPU上执行加1迭代，

执行结果依次是GPU执行的时间，异步时CPU调用GPU所用的时间，以及在GPU上进行计算时CPU进行迭代的次数。

其中：

void *memset(void *s, int ch, size_t n);

函数解释：将s中前n个字节替换为ch并返回s；

memset:作用是在一段内存块中填充某个给定的值，它是对较大的结构体或数组进行清零操作的一种最快方法。

helper_timer.h在C:\ProgramData\NVIDIA Corporation\CUDA Samples\v7.5\common\inc下是在CUDA取消了cutil.h后包含计时的一个文件

cudaThreadSynchronize()函数是是CPU与GPU同步