1. 页锁定主机内存
c库函数malloc()分配标准的,可分页(Pagable)的内存,cudaHostAlloc()分配页锁定的主机内存。页锁定内存也称为固定内存(Pinned Memory)或者不可分页内存,它有个重要属性:操作系统将不会对这块内存分页并交换到磁盘上,从而确保了该内存始终驻留在物理内存中。因此,操作系统能够安全的使某个应用程序访问该内存的物理地址,因为这块内存将不会被破坏或者重新定位。
由于GPU知道内存的物理地址,因此可以通过“直接内存访问(Direct Memory Access ,DMA)”技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU的介入,这也就同样意味着,CPU可能在DMA的执行过程中将目标内存交换到磁盘上,或者通过更新操作系统的可分页表来重新定位目标内存的物理地址。CPU可能会移动可分页的数据,这就可能对DMA操作造成延迟。因此,在DMA复制过程中使用固定内存是非常重要的。事实上,当使用可分页内存进行复制时,CUDA驱动程序仍然会通过DMA把数据传输给GPU。因此,复制操作将执行两遍,第一遍从可分页内存复制到一块临时的页锁定内存,然后再从这个页锁定内存复制到GPU上。因此,每当从可分页内存中执行复制操作时,复制速度将受限于PCIE传输速度和系统前端总线速度相对较低的一方。
当使用固定内存时,将失去虚拟内存的所有功能,导致系统更快的耗尽内存,降低系统整体性能。建议,仅对cudaMemcpy()调用中的源内存或者目标内存,才使用锁定内存,并且不再需要使用它们时立即释放,而不是等到应用程序关闭时才释放。
2. CUDA 流
CUDA流表示一个GPU操作队列,并且该队列中的操作将以指定的顺序执行。可以将每个流视为GPU的一个任务,并且这些任务可以并行执行。
选择一个支持设备重叠功能的设备:
cudaDeviceProp prop; int whichDevice; cudaGetDevice(&whichDevice); cudaGetDevice(&prop, whichDevice); if(!prop.deviceOverlap){ printf("Device will not handle overlaps , so no speed up from streams\n"); }
创建流:
cudaStream_t stream;
cudaStreamCreate(&stream);
使用流:
for(int i =0; i< FULL_DATA_SIZE; i+=N){ cudaMemcpyAsync(dev_a, host_a + i, N*sizeof(int),cudaMemcpyHostToDevice, stream); cudamemcpyAsync(dev_b, host_b + i, N*sizeof(int),cudaMemcpyHostToDevice, stream); kernel<<<N/256,256,0,stream>>>(dev_a, dev_b,dev_c); cudaMemcpyAsync(host_c + i, dev_c, N*sizeof(int),cudaMemcpyDeviceToHost, stream); }
任何传递给cudaMemcpyAsync的主机内存指针都必须已经通过cudaHostAlloc分配好内存,也就是说,你只能以异步方式对页锁定内存进行复制操作。
GPU与主机同步,调用cudaStreamSynchronize()并制定想要等待的流:
cudaStreamSynchronize(stream);
释放流:
cudaStreamDestroy(stream)
3. 使用多个流
for(int i =0; i< FULL_DATA_SIZE; i+= 2*N){ cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int),cudaMemcpyHostToDevice, stream0); cudamemcpyAsync(dev_b0, host_b + i, N*sizeof(int),cudaMemcpyHostToDevice, stream0); kernel<<<N/256,256,0,stream0>>>(dev_a0, dev_b0,dev_c0); cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int),cudaMemcpyDeviceToHost, stream0); //第二个流 cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int),cudaMemcpyHostToDevice, stream1); cudamemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int),cudaMemcpyHostToDevice, stream1); kernel<<<N/256,256,0,stream1>>>(dev_a1, dev_b1,dev_c1); cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int),cudaMemcpyDeviceToHost, stream1); } cudaStreamSynchronize(stream0); cudaStreamSynchronize(stream1);
4.GPU的工作调度机制
在硬件中并没有流的概念,而是包含一个或多个引擎(主机到设备,设备到主机可能是分开的两个引擎)来执行内存复制操作,以及一个引擎来执行核函数。这些引擎彼此独立的对操作进行排队,因此导致下图所示任务调度情形:
应用程序首先将第0个流的所有操作放入队列,然后是第一个流的所有操作。CUDA驱动程序负责按照这些操作的顺序把他们调度到硬件上执行,这就维持了流内部的依赖性。图10.3说明了这些依赖性,箭头表示复制操作要等核函数执行完成之后才能开始。
于是得到这些操作在硬件上执行的时间线:
由于第0个流中将c复制回主机的操作要等待核函数执行完成,因此第1个流中将a和b复制到GPU的操作虽然是完全独立的,但却被阻塞了,这是因为GPU引擎是按照指定的顺序来执行工作。记住,硬件在处理内存复制和核函数执行时分别采用了不同的引擎,因此我们需要知道,将操作放入流队列中的顺序将影响着CUDA驱动程序调度这些操作以及执行的方式。
5. 高效使用多个流
如果同时调度某个流的所有操作,那么很容易在无意中阻塞另一个流的复制操作或者核函数执行。要解决这个问题,在将操作放入流的队列时应采用宽度优先方式,而非深度优先方式。如下代码所示:
for(int i =0; i< FULL_DATA_SIZE; i+= 2*N){ cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int),cudaMemcpyHostToDevice, stream0); cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int),cudaMemcpyHostToDevice, stream1); cudamemcpyAsync(dev_b0, host_b + i, N*sizeof(int),cudaMemcpyHostToDevice, stream0); cudamemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int),cudaMemcpyHostToDevice, stream1); kernel<<<N/256,256,0,stream0>>>(dev_a0, dev_b0,dev_c0); kernel<<<N/256,256,0,stream1>>>(dev_a1, dev_b1,dev_c1); cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int),cudaMemcpyDeviceToHost, stream0); cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int),cudaMemcpyDeviceToHost, stream1); }
如果内存复制操作的时间与核函数执行的时间大致相当,那么新的执行时间线将如图10.5所示,在新的调度顺序中,依赖性仍然能得到满足:
由于采用了宽度优先方式将操作放入各个流的队列中,因此第0个流对c的复制操作将不会阻塞第1个流对a和b的内存复制操作。这使得GPU能够并行的执行复制操作和核函数,从而使应用程序的运行速度显著加快。