CUDA 学习（十八）、优化策略3：传输

一、锁页内存

        主机内存有较大数量的锁页内存。锁页内存允许GPU上的DMA（直接内存访问）控制器请求主机内存传输而不需要CPU主机处理器的参与。因此，在管理传输或从磁盘将换出的页面调回时，没有加载操作需要劳烦主机处理器处理。

         在GPU上分配的内存默认情况下为锁页内存，这只是因为GPU不支持将内存交换到磁盘上。我们关心的是在主机处理器上如何分配锁页内存。为了分配锁页内存，我们需要使用特殊的函数,cudaHostMalloc 函数或者使用常规的malloc函数，然后将其注册为锁页内存。

        注册内存只是设置一些内部标志以确保内存从不被换出，并且告诉CUDA 驱动程序，该内存为锁页内存，所以能够直接使用它，而不需要使用一个临时的缓冲区。

        选择当前可用的硬件时，锁页内存的传输速度大约为非锁页内存（换页内存）传输的两倍。

二、零复制内存

       零复制内存是一种特殊形式的内存映射，它允许将主机内存直接映射到GPU 内存空间上。零复制内存有一个非常有用的使用场合。就是将CPU 应用程序移植到GPU的初始阶段。在这个开发阶段，经常会有主机上的若代码没有被移植到GPU。将这样的数据声明为零复制内存区域就能够允许代码整段地移植并且仍然能工作。在所有代码都真正移植到GPU之前，程序的性能通常是很差的。

       使用零复制或主机映射内存要做三件重要的事情。第一是启用它，第二是使用它分配内存，最后将常规的主机指针转换成指向设备内存空间的指针。

       我们需要在任何CUDA 上下文创建之前进行下面的调用：

       CUDA_CALL（cudaSetDeviceFlags(cudaDeviceMapHost)）;

       当CUDA 上下文被创建是，驱动程序会知道它需要支持主机映射内存。没有驱动程序的支持，主机映射（零复制）内存将无法工作。如果该支持在CUDA 上下文创建之后完成，内存也无法工作。请注意对cudaHostAlloc 这样的函数的调用，尽管在主机内存上执行，也仍然创建一个GPU上下文。

        需要显式的检查：

        struct cudaDeviceProp device_prop

        CUDA_CALL(cudaGetDeviceProperties(&device_prop, device_num));

        zero_copy_supported = device_prop.canMapHostMemory;

        下一个阶段是分配主机上的内存，这样它就可以被映射到设备内存。我们对cudaHostAlloc 函数使用额外的标准cudaHostAllocMapped 就可以实现。

         CUDA_CALL(cudaHostAlloc((void **) &host_data_to_device, size_in_bytes, cudaHostAllocWriteCombined | cudaHostAllocMapped) );

         最后，我们需要通过cudaHostGetDevicePointer 函数将主机指针转换成指向设备指针：

         CUDA_CALL(cudaHostGetDevicePointer(&dev_host_data_to_device, host_data_to_device, 0));

         这样调用中，我们将之前在主机内存空间分配的host_data_to_device 转换成GPU内存空间指针。不要混淆这两个指针。在GPU 内核中，只使用转换后的指针；原始的指针只出现在主机上执行的代码中。

         实际上，使用零复制内存，我们将传输和内核操作分解成更小的块，然后以流水线的方式执行它们。整体时间减少得非常显著。

三、带宽限制

        对于绝大多数的程序而言，最终的带宽限制来源于设备获取输入数据和写回输出数据的 I/O 速度。这也是应用程序无法继续加速的限制。 

        在带宽方面我们遇到第一个问题是：简单的从机器上存入和取出数据。对于这个问题最好的解决方案是使用一个包含多个高速固态硬盘的高速SATA3 RAID 控制器。

        另外一个带宽限制是主机内存速度的限制。但如果在单个计算机节点引入多个GPU，考虑到我们能够以6GB/s 的速度通过PCI-E 总线从高速固态硬盘系统获取数据，这将不是什么大的问题。

四、GPU 计时

1、单核GPU 计时

        流是一种高效得任务队列。当没有向CUDA API 定义流时，0号流用于默认任务队列。然而，0号流包含很多隐式的与主机同步操作。若要使用异步操作，我们需要先建立如下流：

        cudaStream_t   stream;

        CUDA_CALL(cudaStreamaCreate(& stream));

        我们创建了一个事件数组：

        #define MAX_NUM_TEST 16

        cudaEvent_t kernel_start[MAX_NUM_TEST];

        cudaEvent_t kernel_stop[MAX_NUM_TEST];

        GPU提供了一些由GPU硬件赋予时间截的事件。因此，想要在GPU上计算特定操作的执行时间，需要向任务队列中先添加一个启动事件，然后添加想要的计时的操作，最后添加停止事件。GPU上执行的流是简单的FIFO（先进先出）的操作队列。

                                              

         创建流后，我们需要创建一个或多个事件。

         for (u32 test = 0; test < MAX_NUM_TEST; test++)

         {

                 CUDA_CALL(cudaEventCreate(&kernel_start[test]));

                 CUDA_CALL(cudaEventCreate(&kernel_stop[test]));

         }

        我们使用一个简单的循环来建立MAX_NUM_TESTS 个事件，一个启动事件和一个停止事件。当然我们需要在带计时的任务的任一端将这些事件插入流中。

        // Start envet

       CUDA_CALL(cudaEventRecord(kernel_start[test].stream));

       // Run the kernel

       kernel_copy_single<data_T><<<num_blocks, num_threads, dynamic_shared_memory_usage, stream>>>(s_data_in, s_data_out, num_elements);

      // Stop event

     CUDA_CALL(cudaEventRecord(kernel_stop[test].stream));

     为了计算时间，要么每个CUDA 调用要么全部一同调用 cudaEventElapsedTime 函数来获取两个带时间戳事件的时间差。

      // Extract the total time

      for (u32 test = 0; test < MAX_NUM_TESTS; test++)

     {

        float delta;

        // Wait for the event to complete

        CUDA_CALL(cudaEventSynchronize(kernel_stop[test]));

        //Get the time difference

        CUDA_CALL(cudaEventElapsedTime(&delta, kernel_start[test], kernel_stop[test]));

        kernel_time += delta;

      }