zoukankan      html  css  js  c++  java
  • CUDA: 流

    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
    ");  
    }

    创建流:

    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能够并行的执行复制操作和核函数,从而使应用程序的运行速度显著加快。

  • 相关阅读:
    织梦开发——相关阅读likeart应用
    织梦标签教程
    织梦专题调用代码
    HIT 2543 Stone IV
    POJ 3680 Intervals
    HIT 2739 The Chinese Postman Problem
    POJ 1273 Drainage Ditches
    POJ 2455 Secret Milking Machine
    SPOJ 371 Boxes
    HIT 2715 Matrix3
  • 原文地址:https://www.cnblogs.com/programmer-wfq/p/6744878.html
Copyright © 2011-2022 走看看