Streaming workloads are among the simplest that can be ported to CUDA: computations where each data element can be computed independently of the others, often with such low computational density that the workload is bandwidth-bound. Streaming workloads do not make use of many of the hardware resources of the GPU, such as caches and shared memory, that are designed to optimize reuse of data.
Since GPUs give the biggest benefits on workloads with high computational density, it might be useful to review the cases when streaming workloads still make sense to port to GPUs:
If the input and output are in device memory, it doesn’t make sense to transfer the data back to the CPU just to perform one operation;
If the GPU has much better instruction-level support than the CPU for the operation (e.g., Black-Scholes options computation, which uses Special Function Unit instructions intensively), the GPU can outperform the CPU despite memory transfer overhead;
The GPU operates concurrently with the CPU – even if they are the same speed, that can approximately double performance;
The CUDA code for a given workload may be more readable or maintainable than highly optimized CPU code for the same computation;
On integrated systems (i.e., systems-on-a-chip with CPU and CUDA-capable GPU operating on the same memory), there is no transfer overhead – CUDA can use “zero-copy” methods and avoid the copy entirely.
This chapter will cover every aspect of streaming workloads, giving different formulations of the same workload to highlight the different issues that arise. The workload in question, the SAXPY operation from the BLAS library, performs a scalar multiplication and vector addition together in a single operation.
Listing 11-1 gives a trivial C implementation of SAXPY: for corresponding elements in the two input arrays, one element is scaled by a constant, added to the other, and written to the output array. Both input arrays and the output arrays consist of N elements. Since GPUs have a native multiply-add instruction, the innermost loop of SAXPY has an extremely modest number of instructions per memory access.
voidsaxpyCPU( 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]; }}saxpyCPU() (source on GitHub)__global__ voidsaxpyGPU( 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]; }}saxpyGPU()Listing 11-2 gives a trivial CUDA implementation of SAXPY. This version works for any grid or block size, and performs adequately for most applications. This kernel is so bandwidth-bound that most applications would benefit more from restructuring the application to increase the computational density than from optimizing this tiny kernel.
The bulk of this chapter discusses how to move data to and from host memory efficiently, but we’ll spend a moment examining how to improve this kernel’s performance when operating on device memory.