CUDA 流——重叠GPUCPU基础

        在CUDA权威编程指南的第六讲学到了流在GPUCPU重叠当中所起的作用，但是对于这部分一直没有一个直观的体验，今天把Nsight Syste下载回来才有了系统的了解

        前面哪些传统的kernel函数就不写了，直接来看核心部分

for (int i = 0; i < N_SEGMENT; i++)
    {
        int ioffset = i * iElem;
        CHECK(cudaMemcpyAsync(&a_d[ioffset], &a_h[ioffset], nByte / N_SEGMENT, cudaMemcpyHostToDevice, stream[i]));
        CHECK(cudaMemcpyAsync(&b_d[ioffset], &b_h[ioffset], nByte / N_SEGMENT, cudaMemcpyHostToDevice, stream[i]));
        sumArraysGPU << <grid, block, 0, stream[i] >> > (&a_d[ioffset], &b_d[ioffset], &res_d[ioffset], iElem);
        CHECK(cudaMemcpyAsync(&res_from_gpu_h[ioffset], &res_d[ioffset], nByte / N_SEGMENT, cudaMemcpyDeviceToHost, stream[i]));
    }
    //timer
    CHECK(cudaEventRecord(stop, 0));
    int counter = 0;
    while (cudaEventQuery(stop) == cudaErrorNotReady)
    {
        counter++;
    }
    printf("cpu counter:%d\n", counter);

        这一块主要出自谭升大佬的博客，但是对于为什么for循环结束后还能触发cudaEventQuery(stop) == cudaErrorNotReady一直不太理解。

        这个现象的出现是因为作为异步传输，其实就是将命令传输给了device，之后紧接着就进行host的函数，我们来看System可能更加直观一点

 看上去可能是CudaMemcpyAsync是在后面无数个cudaeventQuery之前就结束的，但是这其实只是API的调用，也就是语句的调用

 这张图左边的竖线就代表的是我们上一张图所标注的位置，其实可以看到，在提交完DtoH的命令之后，HtoD的数据传输可能才刚刚开始。这一点是与之前的并行操作完完全全不一样的。并且可以最大化的节俭时间。在上面这个图也可以看到有大大小小的cudaeventquery，针对于为什么cudaeventquery的执行时间不一样，我通过搜索得出了三条可能性：

1. 操作是否完成：cudaEventQuery() 函数的主要作用是查询 CUDA 事件对象，判断其关联操作是否已经完成。当关联操作未完成时，该函数将一直处于阻塞状态，直到操作完成。因此，如果某个操作需要执行更长的时间，那么 cudaEventQuery() 函数的执行时间自然也会变得更长。

2. 程序运行环境：对于同样的 CUDA 代码，在不同的计算机、不同的 GPU 或者不同的设备驱动程序版本下，可能会有所不同的执行时间，这取决于硬件配置、设备驱动程序质量等因素。

3. 事件对象类型：CUDA 事件对象可以分为默认事件对象和手动记录事件对象两种类型。默认事件对象在 CUDA 操作完成后会自动触发，因此 cudaEventQuery() 的执行时间通常较短；而手动记录事件对象则需要在 CUDA 操作完成前手动调用 cudaEventRecord() 函数进行记录，该操作本身也可能需要较多的时间，从而导致 cudaEventQuery() 函数的执行时间延长。

         但很明显，在HtoD的过程中，cudaeventquert就已经出现了不同的时长。但这并不属于以上三条的任意一条，所以这点留个疑问。

        再者就是对于这个代码还有很大的改进空间了，比如可以在写入的同时读取数据，这些就是双缓冲的内容了，等过两天再看看搞一搞。用这个代码拿来入手感觉比cuda samples要好理解得多了。

就写到这里，Nsights System真的比Nsights Compute好用太多了，但我永远爱Nsights Compute。

学了一段时间的cuda了，今天才写下第一篇博客，希望自己能坚持下去吧。