Su'S Blog

什么是异步并行

异步并行，即：控制GPU在没有完成任务请求之前返回CPU。

异步并行的意义

默认情况下，CPU在调用GPU函数时，主机将等待调用完成并返回结果，这就意味着在GPU计算的时候CPU并不会执行任何任务。

而在执行异步并行时，处于同一数据流内的计算与数据传输是依次进行的，但一个流内的计算可以和另一个流的数据传输可以同时进行。

1、通过异步执行就能够使GPU中的执行单元与存储器单元同时工作，更好地压榨GPU的性能。

2、当GPU在进行计算或者数据传输时就返回给主机线程，主机不必等待GPU运行完毕就可以进行进行一些计算，更好地压榨了CPU的性能。

例程

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "stdio.h"
#include "memory.h"
#include "time.h"
#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[])
{
    // CUDA device properties
    // CUDA 设备属性
    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;
    // Allocates page-locked memory on the host.
    // 在主机上申请 page-locked 内存
    cudaMallocHost((void**)&a, nbytes);
    // 为内存赋值为0
    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 << <blocks, threads, 0, 0 >> >(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 // 返回 0
        printf("Test FAILED\n");

    cudaEventDestroy(startGPU);
    cudaEventDestroy(stopGPU);

    // 释放内存
    cudaFreeHost(a);
    cudaFree(d_a);
    getchar();
    cudaThreadExit();
    return 0;
}

参考资料：

1、CUDA入门（六） 异步并行执行解析： https://blog.csdn.net/qq_25819827/article/details/52541813