PaperlayerView exact revision on arXiv

Source-linked AI summary

Compressing DMA Engine: Leveraging Activation Sparsity for Training Deep Neural Networks

Minsoo Rhu, Mike O'Connor, Niladrish Chatterjee, Jeff Pool, Stephen W. Keckler

arXiv:1705.01626v1cs.LGcs.AR

TL;DR

DNN training can exceed physical GPU memory, while virtualizing memory through CPU–GPU transfers can become a performance bottleneck. The paper introduces cDMA, which compresses sparse activation maps during transfer, reporting 2.6× average compression and 32% average performance improvement for virtualized DNNs.

  • Problem

    GPU memory virtualization enables larger DNN training but can incur substantial overhead when PCIe transfers take longer than forward and backward propagation.

  • Method

    cDMA compresses sparse activation maps on the fly before transferring them between GPU and CPU memory over PCIe.

  • Results

    2.6× average compression and 32% average performance improvement are reported for virtualized DNNs.

  • Takeaways & Limitations

    cDMA reduces CPU–GPU communication traffic while providing a general-purpose virtualization strategy that can be adopted into existing ML frameworks.

  • Takeaways & Limitations

    cDMA does not reduce GPU DRAM capacity or bandwidth requirements because activation maps remain uncompressed inside GPU DRAM.

Abstract

from arXiv · show

Popular deep learning frameworks require users to fine-tune their memory usage so that the training data of a deep neural network (DNN) fits within the GPU physical memory. Prior work tries to address this restriction by virtualizing the memory usage of DNNs, enabling both CPU and GPU memory to be utilized for memory allocations. Despite its merits, virtualizing memory can incur significant performance overheads when the time needed to copy data back and forth from CPU memory is higher than the latency to perform the computations required for DNN forward and backward propagation. We introduce a high-performance virtualization strategy based on a "compressing DMA engine" (cDMA) that drastically reduces the size of the data structures that are targeted for CPU-side allocations. The cDMA engine offers an average 2.6x (maximum 13.8x) compression ratio by exploiting the sparsity inherent in offloaded data, improving the performance of virtualized DNNs by an average 32% (maximum 61%).

I. INTRODUCTION

DNN frameworks rely on GPU acceleration but are constrained by limited physical GPU memory. Virtualized memory expands capacity using CPU memory, yet PCIe transfers can impose substantial performance overheads, motivating cDMA's compression-based approach.

  • GPU acceleration trains DNNs much faster than CPUs, but frameworks are constrained by limited physical GPU memory.
  • vDNN virtualizes DNN memory by exposing CPU and GPU memory concurrently, improving the ability to train larger and deeper networks.
  • 31% average performance loss, with a 52% worst case, can result when PCIe data movement is bottlenecked relative to DNN computation.
  • cDMA reduces copied data sizes by compressing sparse activation maps during transfers between GPU and CPU memory.
  • CNNs use convolutional, activation, pooling, and fully-connected layers for feature extraction and classification.

B. Training versus Inference

DNN training repeatedly performs forward propagation, loss calculation, and backward propagation while managing layer activation maps. vDNN coordinates activation offloads and prefetches so these maps can move between GPU and CPU memory during training.

  • Each training pass combines forward propagation, error calculation against ground truth, and backward propagation.
  • Forward propagation processes layers sequentially from input to output, generating output activation maps from input maps and layer operations.
  • Backward propagation reverses the layer order, uses incoming gradients to produce gradients for preceding layers, and updates layer weights.
  • CNN activation maps are four-dimensional arrays organized by batch size, channels, height, and width.
  • vDNN offloads inter-layer activation maps during forward propagation and prefetches them before reuse during backward propagation.

III. MOTIVATION

Faster GPU computation has widened the gap with fixed PCIe transfer bandwidth, increasing vDNN overheads. The paper motivates cDMA by exploiting sparse, compressible activation maps to reduce transferred data.

  • vDNN offloads activation maps that occupy more than 90% of GPU-side training allocations, reducing average GPU memory usage.
  • vDNN can stall when CPU–GPU data movement takes longer than backpropagation computation.
  • 2.2× average speedup from cuDNN v1 to v5 contrasts with unchanged 16 GB/sec PCIe gen3 bandwidth, increasing vDNN overheads.
  • Activation maps contain substantial zero-valued sparsity, making them highly compressible for PCIe transfer.
  • The evaluation primarily focuses on CNNs because realistic, publicly available datasets support analysis of input-dependent compression effectiveness.

IV. SPARSITY OF DNN ACTIVATIONS

The study characterizes how activation sparsity changes across DNN layers and training, finding substantial and structured sparsity across CNNs. Sparsity generally follows a U-shaped training trajectory, varies by layer type and depth, and provides a basis for compression.

  • Except for conv0, average activation density follows a U-shaped training pattern: it decreases early and increases later as accuracy improves.
  • Conv0 remains near 50% average activation sparsity throughout training, varying within ±2%.
  • Pooling layers increase activation density, while later layers—especially fully connected layers—are generally sparser than earlier layers.
  • 49.4% average activation sparsity is observed across AlexNet, implying that perfect zero removal could reduce activation size by about half.
  • 62% average network-wide activation sparsity, with a 93% maximum, is observed across six CNNs throughout training.
  • Early training density drops alongside rapid loss improvement, later density rises first rapidly and then more slowly, and final fine-tuning changes average sparsity minimally.

V. COMPRESSING DMA ENGINE

The cDMA approach compresses sparse activation maps before transferring them between GPU and CPU memory. It targets PCIe transfer overhead by exploiting zero-valued activations in offloaded data.

  • cDMA compresses activation maps before PCIe transfer to reduce the size of data moved between GPU and CPU memory.
  • The approach targets activation-map movement because activation sparsity makes these offloaded data structures highly compressible.

A. Compression Algorithm

The compression algorithms considered for cDMA prioritize high throughput, with zero-value compression selected to exploit the prevalence of zero activations. ZVC encodes non-zero positions using masks and appends the corresponding values.

  • Compression must sustain rates of hundreds of GB/sec while saturating PCIe bandwidth with compressed data.
  • RLE exploits clustered zeros and non-zeros but its effectiveness depends strongly on consecutive-value sparsity patterns.
  • Approximately 50% to 90% of activations are zero-valued, motivating zero-value compression for removing zero elements.
  • ZVC processes 32 activation values with a 32-bit zero/non-zero mask, then appends the non-zero elements.
  • A window of 32 consecutive zeros compresses to one all-zero 32-bit mask, yielding a 32× compression ratio.
  • ZVC hardware complexity is dominated by MUXes and mask pop-count/prefix-sum operations for gathering, scattering, and offset calculation.

B. Compressing DMA Engine Architecture

The cDMA architecture places compression and decompression near GPU memory controllers while adding buffering at the PCIe-facing DMA engine. This placement reduces crossbar bandwidth pressure while supporting high-throughput transfers.

  • cDMA adds compression/decompression units beside GPU memory controllers and additional buffering in the DMA engine at the PCIe interface.
  • An alternative design places compression directly inside the DMA unit, but it raises concerns about GPU on-chip crossbar bandwidth requirements.
  • 13.8× maximum per-layer compression is observed, with the design objective of generating compressed activation maps at PCIe-link throughput.
  • Placing compression in GPU memory controllers reduces compressed-DMA crossbar demand to levels similar to baseline non-compressing DMA.
  • The compression engine processes 32B per cycle, comparing eight words to zero and using prefix sums to determine compressed offsets.
  • The decompression engine expands compressed data into 128-byte outputs while producing 32B of decompressed data per cycle.
  • The decompression pipeline adds only two cycles of latency for a 128-byte line.

C. Design Overheads

The cDMA design adds compression and decompression support around the DMA engine while sizing buffering to sustain data delivery without output-stream bubbles.

  • Decompression units: cDMA supplements existing GPU compression hardware for a conservative area estimate while allowing DRAM compression schemes to be reused.The design leverages compression units already employed in GPU memory controllers.
  • Buffer sizing: 70KB of buffering is required for 200 GB/sec compression-read bandwidth and 350 ns memory-request latency.Experiments found that provisioning 200 GB/sec captures most benefits of sparse compression.

D. Software Interface

The evaluation integrates cDMA with vDNN and Caffe-based compression processing, while modeling bandwidth constraints and using a defined DNN training workload.

  • D. Software Interface: The software interface could expose compressed memory copies through a cudaMemcpyCompressed() call that returns the compressed region size.The proposed API enables DMA-engine compression or decompression for ML frameworks and other applications.
  • D. Software Interface: Cycle-level simulation is difficult because GPU simulators cannot execute precompiled cuDNN APIs and training iterations can take tens of seconds.These constraints make cycle-level simulation of the workloads impractical within a reasonable timeframe.
  • D. Software Interface: vDNN offloads all layer activation maps, with allocation sizes determined by cDMA’s observed compression ratio.Offload and prefetch use cudaMemcpyAsync().
  • D. Software Interface: The Caffe compression pipeline checkpoints activations, compresses each layer’s offloaded data, and measures CPU-side copy latency.The compressed maps are returned to the vDNN memory manager.
  • D. Software Interface: cDMA performance depends on reduced PCIe traffic and generating compressed activations fast enough to saturate PCIe bandwidth.The reduced transfer latency improves vDNN, but compression throughput must match the transfer path.
  • D. Software Interface: cDMA bandwidth is restricted to 236 GB/sec of leftover Titan X DRAM bandwidth, and layers exceeding that limit receive increased modeled offload latency.The platform provides 336 GB/sec DRAM bandwidth and 16 GB/sec PCIe bandwidth.
  • D. Software Interface: The six evaluated ImageNet networks are AlexNet, OverFeat, NiN, VGG, SqueezeNet, and GoogLeNet.Training uses SGD with an initial learning rate of 0.01, with dropout applied to fully connected layers.
  • D. Software Interface: The average network-wide compression ratio of 2.6 implies average cDMA memory-bandwidth usage of 41.3 GB/sec when PCIe bandwidth is 16 GB/sec.Peak cDMA bandwidth can approach 200 GB/sec, but average usage is lower.

VII. RESULTS

The results evaluate compression efficiency and PCIe traffic reduction across three algorithms and three activation layouts, showing ZVC as the most robust option overall.

  • VII. RESULTS: The evaluation compares run-length encoding, zero-value compression, and zlib while offloading all activation maps under a memory-scalable vDNN policy.An oracle baseline was also established for performance comparison.
  • A. Compression Efficiency: The study reports maximum per-layer and average network-wide compression ratios using a 4KB compression window, with larger windows producing similar results.The maximum ratio determines required DRAM bandwidth, while the average ratio reflects PCIe traffic reduction.
  • A. Compression Efficiency: 2.6× average compression ratio makes ZVC the best average compressor across networks and activation layouts.ZVC’s compression ratio is decoupled from activation arrangement, unlike the layout-sensitive RLE and zlib results.
  • A. Compression Efficiency: 3% average traffic reduction separates zlib from ZVC across six networks, despite zlib’s 30% additional reduction for GoogLeNet.The sophisticated compressor therefore offers little average PCIe advantage over ZVC.

B. Performance

cDMA improves virtualized-DNN performance by reducing PCIe traffic, while its energy-efficiency evaluation remains qualitative because the CUDA stack cannot vary key bandwidths.

  • B. Performance: 0.7% average speedup over ZVC, with a maximum 2.2% for GoogLeNet, results from zlib’s higher compression ratios.The gains remain marginal because offloading is often hidden by DNN computation and overlaps occur in already-hidden layers.
  • B. Performance: ZVC is selected as the best DNN-virtualization option because of its simple compression and robustness across activation layouts.Its practical advantage is not solely the highest compression ratio for every network.
  • C. Energy Efficiency: 2.6× average DRAM-read bandwidth is cDMA’s primary additional memory-side energy overhead, alongside compression units and buffers.The design primarily leverages compression hardware already present in GPU memory controllers.
  • C. Energy Efficiency: Energy-efficiency effects cannot be measured precisely because CUDA does not let users change DRAM-read or PCIe-transfer bandwidth.The paper therefore provides a qualitative comparison with vDNN.
  • C. Energy Efficiency: 1–7% power overhead from vDNN offload and prefetch is contrasted with cDMA’s reduced PCIe traffic and CPU-memory energy use.The paper expects overall energy consumption to improve when accounting for cDMA’s performance gains.
  • B. Performance: Figure 13 compares cDMA performance with vDNN and an oracle baseline.The comparison summarizes overall performance normalized to the oracle baseline.

IX. DISCUSSION AND FUTURE WORK

cDMA reduces offloading traffic and CPU-side memory footprint, while future work addresses interconnect-specific bandwidth constraints and GPU DRAM footprint reduction. Its reported benefits include lower data movement and improved virtualized-DNN performance with modest implementation overhead.

  • Future CPU-GPU interconnects: 10–20 GB/sec per-GPU bandwidth on multi-GPU NVLINK systems remains similar to PCIe, preserving the importance of reducing CPU–GPU offloading traffic.NVLINK can provide up to 80 GB/sec for CPU–GPU communication in coupled IBM Power systems, but shared multi-GPU channels reduce bandwidth per GPU.
  • Compression for GPU footprint reduction: cDMA leaves GPU memory allocation unchanged, so compressing activation maps inside GPU DRAM remains future work.That optimization would require efficient memory addressing to retrieve compressed data without disrupting performance and energy efficiency.
  • Discussion: 2.6× average and 13.8× maximum savings in data movement reduce traffic on the CPU–GPU communication link.The savings exploit sparsity in offloaded DNN data structures.
  • Discussion: 32% average and 61% maximum performance improvements are reported for virtualized DNNs with modest implementation overhead.The design can be adopted into existing machine-learning frameworks.
Loading 1705.01626v1…