第11章 流式负载

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

鉴于GPU对于高计算密度的负载具有最大效力，我们将讨论几种情形，以助于流式负载移植到GPU。

·如果输入和输出均在设备内存上，让数据传回CPU却仅执行一次计算操作是得不偿失的。
·如果GPU有比CPU更好的指令级操作支持（例如，频繁使用特殊功能单位指令的Black-Scholes期权计算任务），尽管存在额外的内存传输开销，GPU仍可以超过CPU。
·GPU操作与CPU并发执行可以增加约一倍的性能（即使假设它们的速度相同）。
·对于给定负载，CUDA代码相对于高度优化的CPU代码可能会更可读或更可维护。

• 对于包含集成显卡的系统（即CPU与支持CUDA的GPU共存于同一芯片，并作用于同一内存），不存在传输开销。CUDA技术可以使用“零复制”方法，完全避免复制操作。

本章涵盖流式负载的方方面面。针对同一负载，采用不同实现策略，以暴露可能出现的多个问题。这里关注的负载是来自BLAS库的SAXPY操作。它在一个操作中同时执行一个标量乘法和矢量加法。

代码清单11-1给出了一个实现SAXPY的简单C程序。每次取出两个输入数组相对应的两个数，对其中一个施行常倍数的缩放，并加到另一个数上，然后写回输出数组的对应位置。两个输入数组和输出数组均包含N个元素。由于GPU包含原生乘加指令，SAXPY的最内层循环在每次内存访问时，仅包含较少的指令数。

代码清单11-1 saxpyCPU函数

void   
saxpyCPU( float \*out, const float \*x, const float \*y, size_t N, float alpha) { for (size_t i = 0; i < N; i++) { out[i] += alpha\*x[i] + y[i]; }

代码清单11-2给出了SAXPY的一个简单CUDA实现版本。这个版本可以运行于任何网格或各种大小的线程块，并且其执行性能能满足大多

数应用的需要。由于这个内核严重受制于带宽，对多数应用程序而言，与其在它的基础上进行优化，不如重构该程序以增加计算密度。

代码清单11-2 saxpyGPU内核函数

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];
            }
}

本章的主旨在于讨论如何将数据高效地传入和传出主机内存。但是，首先我们要花些时间研究如何从操作设备内存的角度改善这个内核的性能。