stream 流式负载优化 （CUDA 版本）

Aalan157
304 阅读6分钟

Performance benchmark 中包含一类 stream workload 常被用来测量 GPU 与 HOST 传输数据的带宽性能。本文主要通过学习 《CUDA Handbook》总结整理对于 stream kernel 常见的优化思路，该优化思路不仅适用于 CUDA 编程体系以及 NVidia 的 GPU, 同时对其他硬件加速器也提供了优化 stream kernel 的方向。 本文主要包含：

  • stream kernel 介绍以及 bottleneck 分析
  • kernel 优化以及数据传输优化
  • 利用 nsys debug stream 在硬件上真实的执行状态

什么是 stream kernel？

流式负载（streaming workload）是每个数据元素可以被独立地计算的任务，它是可以移植到CUDA中的最简单的负载。这类低计算密度（computational density）的负载通常属于带宽受限型（bandwidth-bound）。流式负载无须使用太多的GPU硬件资源，例如用于优化数据重用的高速缓存和共享存储器等。

鉴于GPU对于高计算密度的负载具有最大效力，我们主要从 Kernel 本身 以及 数据传输两个角度进行 stream workload 的优化。

stream kernel 具体名称用 saxpyGPU 代替。

"Talk is cheap, show me the code"

优化流程分析

image.png

Kernel 优化

一般的 stream Kernel 写法

__global__ void saxpyGPU( float *out, const float *x, const float *y, size_t N, float alpha ) 
{ 
    for ( size_t i = blockIdx.x*blockDim.x + threadIdx.x;
    i < N; 
    i += blockDim.x*gridDim.x ) { 
        out[i] = alpha*x[i]+y[i]; 
    } 
    }

循环展开 Unroll stream Kernel

__device__ void
saxpy_unrolled(  float *out,  const float *px, const float *py,  size_t N, float alpha )
{
    float x[n], y[n];
    size_t i;
    for ( i = n*blockIdx.x*blockDim.x+threadIdx.x; 
          i < N-n*blockDim.x*gridDim.x; 
          i += n*blockDim.x*gridDim.x ) {
        for ( int j = 0; j < n; j++ ) {
            size_t index = i+j*blockDim.x;
            x[j] = px[index];
            y[j] = py[index];
        }
        for ( int j = 0; j < n; j++ ) {
            size_t index = i+j*blockDim.x;
            out[index] = alpha*x[j]+y[j];
        }
    }
    // to avoid the (index<N) conditional in the inner loop, 
    // we left off some work at the end
    for ( int j = 0; j < n; j++ ) {
        for ( int j = 0; j < n; j++ ) {
            size_t index = i+j*blockDim.x;
            if ( index<N ) {
                x[j] = px[index];
                y[j] = py[index];
            }
        }
        for ( int j = 0; j < n; j++ ) {
            size_t index = i+j*blockDim.x;
            if ( index<N ) out[index] = alpha*x[j]+y[j];
        }
    }

数据传输优化

stream1Device

Memcpy 都是同步拷贝，拷贝性能较低

Code：

    chTimerGetTime( &chStart );
    cuda(EventRecord( evStart, 0 ) );
    cuda(Memcpy( dptrX, hptrX, N*sizeof(float), cudaMemcpyHostToDevice ) );
    cuda(Memcpy( dptrY, hptrY, N*sizeof(float), cudaMemcpyHostToDevice ) );
    cuda(EventRecord( evHtoD, 0 ) );
        saxpyGPU<<<nBlocks, nThreads>>>( dptrOut, dptrX, dptrY, N, alpha );
    cuda(EventRecord( evKernel, 0 ) );
    cuda(Memcpy( hptrOut, dptrOut, N*sizeof(float), cudaMemcpyDeviceToHost ) );
    cuda(EventRecord( evDtoH, 0 ) );
    cuda(DeviceSynchronize() );

Performance：

    ./stream1Device -N 1024
Measuring times with 1024M floats
Memcpy( host->device ): 2053.36 ms (4183.36 MB/s)
Kernel processing     : 29.25 ms (440522.25 MB/s)
Memcpy (device->host ): 1113.78 ms (3856.21 MB/s)

nsys profile:

image.png

unpinned mem copy

stream2Async

Kernel 不变，使用cudaMemcpyAsync ==> 从而利用Pinned memory 以及 异步拷贝，从而提升数据传输效率。但是本身无法掩盖 数据传输 以及 Kernel 执行的时间。

和 stream1Device 相比，Kernel 执行时间大致相同 28~29ms， 但是Memcpy 的时间显著减少，带宽从 4183.36 MB/s 增加到 12981.42 MB/s 提升了 3 倍的数据传输性能。

为什么锁页内存比非锁页内存更快：

CUDA 驱动程序通过检查内存范围判断某个地址是锁页内存还是分页内存。锁页内存存储在物理内存中，因此 device 可以在没有 CPU 帮助的情况下获取它（通过DMA）。分页内存在通过 DMA 访问时会产生缺页中断，并且它有可能在磁盘上。在这种情况下，device 需要访问分页内存的每一页，将其拷贝到锁页内存缓冲区，然后再将其通过 DMA 一页页拷贝到 device 上。所以，使用锁页内存更快是因为省掉了从分页内存拷贝到锁页内存的时间。

锁页内存的缺点：由于内存的物理页无法被置换，非常容易产生 out of memory。

Code：

   chTimerGetTime( &chStart );
    cuda(EventRecord( evStart, 0 ) );
    cuda(MemcpyAsync( dptrX, hptrX, N*sizeof(float), cudaMemcpyHostToDevice, NULL ) );
    cuda(MemcpyAsync( dptrY, hptrY, N*sizeof(float), cudaMemcpyHostToDevice, NULL ) );
    cuda(EventRecord( evHtoD, 0 ) );
        saxpyGPU<<<nBlocks, nThreads>>>( dptrOut, dptrX, dptrY, N, alpha );
    cuda(EventRecord( evKernel, 0 ) );
    cuda(MemcpyAsync( hptrOut, dptrOut, N*sizeof(float), cudaMemcpyDeviceToHost, NULL ) );
    cuda(EventRecord( evDtoH, 0 ) );
    cuda(DeviceSynchronize() );

Performance：

./stream2Async -N 1024
Measuring times with 1024M floats
Memcpy( host->device ): 661.71 ms (12981.42 MB/s)
Kernel processing     : 28.78 ms (447758.10 MB/s)
Memcpy (device->host ): 325.72 ms (13186.26 MB/s)

nsys profile:

image.png

可以发现 CUDA API 调用完立马结束了，cudaMemcpyAsync 的返回不代表数据拷贝完成了，具体的拷贝会在 GPU 上执行一段时间，有点是 GPU 可以自主的完成数据拷贝，不需要 HOST 参与与等待，同时由于事先 pin 住了 memory 数据传输也会更快一点。

image.png

pinned mem copy.png

对于 锁页内存 传输相同的数据两大小 1GB 也从 1s 下降到 0.33085 s 了（相对于 unpinned mem copy）

stream3Streams

启用 nStreams 实现 task paralle，加上 异步的内存拷贝，从而可以用 数据传输 掩盖部分Kernel 的执行时间， 对于 nstream 的 code 具体有两种写法。

  1. 不同的流同时调用 memcpy 然后调用 Kernel 到不同 stream 上， 最后将数据拷贝回 CPU 效率更高

  2. 单个流完成 host->gpu copy + kernel + gpu->host mem copy

Code：

   for ( int iStream = 0; iStream < nStreams; iStream++ ) {
        cuda(MemcpyAsync( 
                          dptrX+iStream*streamStep, 
                          hptrX+iStream*streamStep, 
                          streamStep*sizeof(float), 
                          cudaMemcpyHostToDevice, 
                          streams[iStream] ) );
        cuda(MemcpyAsync( 
                          dptrY+iStream*streamStep, 
                          hptrY+iStream*streamStep, 
                          streamStep*sizeof(float), 
                          cudaMemcpyHostToDevice, 
                          streams[iStream] ) );
    }

    for ( int iStream = 0; iStream < nStreams; iStream++ ) {
        saxpyGPU<<<nBlocks, nThreads, 0, streams[iStream]>>>( 
            dptrOut+iStream*streamStep, 
            dptrX+iStream*streamStep, 
            dptrY+iStream*streamStep, 
            streamStep, 
            alpha );
    }

Performance：

./stream3Streams -N 1024
Measuring times with 1024M floats
Testing with default max of 8 streams (set with --maxStreams <count>)
Streams Time (ms)       MB/s
1       1016.15 ms      12424.99
2       856.81 ms       14673.91
3       815.44 ms       15397.44
4       779.69 ms       16087.83
5       773.50 ms       16209.43
6       769.64 ms       16292.87
7       761.90 ms       16451.80
8       750.70 ms       16691.78

nsys profile:

image.png

image.png

从这张具体执行的图中可以发现，基本上实现了多任务的流水线。

stream4Mapped

映射锁页内存，mapped memory 拥有两个地址：主机端地址（内存地址）和设备地址（显存地址），可以在 kernel 中直接访问 mapped memory 中的数据，而不必再在内存和显存间进行数据拷贝，即zero-copy 功能。如果内核程序只需要对 mapped memory 进行少量读写，这样做可以减少分配显存和数据拷贝的时间。优点是省去了H2D, D2H 的显示内存拷贝。 这样也可以减少GPU 操作的延时。

但是kernel 执行还是需要数据从HOST 传到 Device， 数据拷贝的时间会隐藏到 Kernel 执行当中。

Code：

    float *dptrOut = 0, *hptrOut = 0;
    float *dptrY = 0, *hptrY = 0;
    float *dptrX = 0, *hptrX = 0;
    cudaEvent_t evStart = 0;
    cudaEvent_t evStop = 0;

    cuda(HostAlloc( &hptrOut, N*sizeof(float), cudaHostAllocMapped ) );
    cuda(HostGetDevicePointer( &dptrOut, hptrOut, 0 ) );
    memset( hptrOut, 0, N*sizeof(float) );

    cuda(HostAlloc( &hptrY, N*sizeof(float), cudaHostAllocMapped ) );
    cuda(HostGetDevicePointer( &dptrY, hptrY, 0 ) );
    cuda(HostAlloc( &hptrX, N*sizeof(float), cudaHostAllocMapped ) );
    cuda(HostGetDevicePointer( &dptrX, hptrX, 0 ) );
    ......
    saxpyGPU<<<nBlocks, nThreads>>>( dptrOut, dptrX, dptrY, N, alpha );

Performance：

~/TensorRT/cudahandbook/streaming/build$ ./stream4Mapped -N 1024
Measuring times with 1024M floats
Total time (GPU event):  685.43 ms (18798.39 MB/s)
Total time (wall clock): 685.48 ms (18796.90 MB/s)

nsys profile:

这个 saxpy GPU kernel 执行时间是 685.245ms ，其中是包含了数据从 HOST -> Device 以及算完之后 Device->Host 的传输，我们可以与 stream2Async 的 Kernel 执行状态进行对比。 在stream4Mapped 中， 对于 Mapped memory ，saxpyGPU（Mapped memory） 相当于 MemcopyH2D + MemcopyH2D + 纯粹的 saxpyGPU 计算 + MemcopyD2H （pinned Mem with no mapped）

但是我们会发现， 在 Mapped memory 中 saxpyGPU 耗时 685.245ms 是小于 pinned Mem 中 （330.85+320.85+28.585+325.699）= 1005.98 ms，这表明 stream4Mapped 整体的 E2E 性能是优于 stream2Async。

image.png

image.png

avatar
文章
阅读
粉丝