0_Simple__asyncAPI

▶ CPU - GPU 异步操作

▶ 源代码

#include <stdio.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <helper_cuda.h>
#include <helper_functions.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;
}

bool correct_output(int *data, const int n, const int x)
{
    for (int i = 0; i < n; i++)
    {
        if (data[i] != x)
        {
            printf("Error! data[%d] = %d, ref = %d
", i, data[i], x);
            return false;
        }
    }
    return true;
}

int main(int argc, char *argv[])
{
    printf("Start.
");
    int devID = findCudaDevice(argc, (const char **)argv);  // 通过命令行参数选择设备，可以为空
    cudaDeviceProp deviceProps;
    cudaGetDeviceProperties(&deviceProps, devID);
    printf("CUDA device [%s]
", deviceProps.name);

    const int n = 16 * 1024 * 1024;
    const int nbytes = n * sizeof(int);
    const int value = 26;

    int *a, *d_a;
    cudaMallocHost((void **)&a, nbytes);
    cudaMalloc((void **)&d_a, nbytes);
    memset(a, 0, nbytes);
    cudaMemset(d_a, 255, nbytes);

    cudaEvent_t start, stop;                // GPU 端计时器
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    StopWatchInterface *timer = NULL;       // CPU 端计时器
    sdkCreateTimer(&timer);
    sdkResetTimer(&timer);

    dim3 threads = dim3(512, 1, 1);
    dim3 blocks = dim3(n / threads.x, 1, 1);

    sdkStartTimer(&timer);                  // 注意 GPU 计时器是夹在 CPU 计时器内的，但是 GPU 函数都是异步的
    cudaEventRecord(start, 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(stop, 0);
    sdkStopTimer(&timer);

    unsigned long int counter = 0;          // 记录 GPU 运行完成以前 CPU 运行了多少次 while 的循环 
    while (cudaEventQuery(stop) == cudaErrorNotReady)
        counter++;

    float gpu_time = 0.0f;                  // 此时保证 GPU 运行完成，才能记录时间
    cudaEventElapsedTime(&gpu_time, start, stop);

    printf("time spent by GPU: %.2f
", gpu_time);
    printf("time spent by CPU: %.2f
", sdkGetTimerValue(&timer));
    printf("CPU executed %lu iterations while waiting for GPU to finish
", counter);
    printf("
	Finish: %s.", correct_output(a, n, value) ? "Pass" : "Fail");

    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    cudaFreeHost(a);
    cudaFree(d_a);
    getchar();
    return 0;
}

● 输出结果：

GPU Device 0: "GeForce GTX 1070" with compute capability 6.1

CUDA device [GeForce GTX 1070]
time spent by GPU: 11.50
time spent by CPU: 0.05
CPU executed 3026 iterations while waiting for GPU to finish

        Finish!

▶ 新姿势：

● 调用主函数时的第0个参数作为程序名字符串，可以用于输出。

int main(int argc, char *argv[])
...
printf("%s", argv[0]);

● 在没有附加 flag 的情况下申请主机内存，注意使用cudaFreeHost释放

int *a, nbytes = n * sizeof(int);
cudaMallocHost((void **)&a, nbytes);
...
cudaFreeHost(a);

● 记录 CPU 调用 CUDA 所用的时间

StopWatchInterface *timer = NULL;
sdkCreateTimer(&timer);
sdkResetTimer(&timer);
sdkStartTimer(&timer);
    
...// 核函数调用
    
sdkStopTimer(&timer);
printf("%.2f ms", sdkGetTimerValue(&timer));

● 查看GPU队列状态的函数

extern __host__ cudaError_t CUDARTAPI cudaEventQuery(cudaEvent_t event);

■ stop为放置到流中的一个事件，cudaEventQuery(stop)返回该事件的状态，等于cudaSuccess（值等于0）表示已经发生；等于cudaErrorNotReady（值等于35）表示尚未发生。源代码中利用这段时间让CPU空转，记录了迭代次数。

while (cudaEventQuery(stop) == cudaErrorNotReady) counter++;

● stdlib.h 中关于返回成功和失败的宏

#define EXIT_SUCCESS 0
#define EXIT_FAILURE  1

● 示例文件中的错误检查函数（定义在helper_cuda.h中），报告出错文件、行号、函数名，并且重启cudaDevice。

#define checkCudaErrors(val)  check((val), #val, __FILE__, __LINE__)

template< typename T >
void check(T result, char const *const func, const char *const file, int const line)
{
    if (result)
    {
        fprintf(stderr, "CUDA error at %s:%d code=%d(%s) "%s" 
",
            file, line, static_cast<unsigned int>(result), _cudaGetErrorEnum(result), func);
        DEVICE_RESET// Make sure we call CUDA Device Reset before exiting
        exit(EXIT_FAILURE);
    }
}

#define DEVICE_RESET  cudaDeviceReset();