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

  • 相关阅读:
    配置JAVA的环境变量
    Navicat Premium v12.0.23.0 安装,使用激活码激活
    mysql 5.7.33安装教程
    查看mysql服务
    postman测试post请求参数是List entity后端用@RequestBody接受lIst entity方式
    最简单的MySQL安装教程(数据库免安装版,免安装版,亲测好用) mysql-5.7.33-winx64.zip
    mysql数据库安装5.7.33
    java删除附件图片
    java生成条形码
    最高的牛
  • 原文地址:https://www.cnblogs.com/programmer-wfq/p/6744878.html
Copyright © 2011-2022 走看看