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.

    Source: §I · p. 1

  • Method

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

    Source: §I · p. 1, §V · p. 6

  • Results

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

    Source: §X · p. 12

  • Takeaways & Limitations

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

    Source: §X · p. 12

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

    Source: §IX · p. 12

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.Source: §I · p. 1
  • vDNN virtualizes DNN memory by exposing CPU and GPU memory concurrently, improving the ability to train larger and deeper networks.Source: §I · p. 1
  • 31% average performance loss, with a 52% worst case, can result when PCIe data movement is bottlenecked relative to DNN computation.Source: §I · p. 1
  • cDMA reduces copied data sizes by compressing sparse activation maps during transfers between GPU and CPU memory.Source: §I · p. 1
  • CNNs use convolutional, activation, pooling, and fully-connected layers for feature extraction and classification.Source: A. Deep Neural Networks · p. 2

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.Source: B. Training versus Inference · p. 2
  • Forward propagation processes layers sequentially from input to output, generating output activation maps from input maps and layer operations.Source: B. Training versus Inference · p. 2
  • Backward propagation reverses the layer order, uses incoming gradients to produce gradients for preceding layers, and updates layer weights.Source: B. Training versus Inference · p. 2
  • CNN activation maps are four-dimensional arrays organized by batch size, channels, height, and width.Source: C. Data Layout for Activation Maps · p. 2
  • vDNN offloads inter-layer activation maps during forward propagation and prefetches them before reuse during backward propagation.Source: Figure 2 · p. 3

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.Source: §III · p. 3
  • vDNN can stall when CPU–GPU data movement takes longer than backpropagation computation.Source: §III · p. 3
  • 2.2× average speedup from cuDNN v1 to v5 contrasts with unchanged 16 GB/sec PCIe gen3 bandwidth, increasing vDNN overheads.Source: §III · p. 3
  • Activation maps contain substantial zero-valued sparsity, making them highly compressible for PCIe transfer.Source: §III · p. 3
  • The evaluation primarily focuses on CNNs because realistic, publicly available datasets support analysis of input-dependent compression effectiveness.Source: §III · p. 4

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.

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.Source: §V · p. 6
  • The approach targets activation-map movement because activation sparsity makes these offloaded data structures highly compressible.Source: §V · p. 6

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.

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.

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.Source: C. Design Overheads · p. 8

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.Source: C. Design Overheads · p. 9

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.Source: D. Software Interface · p. 9
  • 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.Source: §VI · p. 9
  • vDNN offloads all layer activation maps, with allocation sizes determined by cDMA’s observed compression ratio.Offload and prefetch use cudaMemcpyAsync().Source: §VI · p. 9
  • 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.Source: §VI · p. 9
  • 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.Source: §VI · p. 9
  • 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.Source: §VI · pp. 9–10
  • 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.Source: §VI · p. 10
  • 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.Source: §VI · p. 10

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.Source: §VII · p. 10

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.Source: A. Compression Efficiency · p. 10
  • 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.Source: A. Compression Efficiency · p. 10
  • 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.Source: A. Compression Efficiency · p. 10

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.Source: B. Performance · p. 11
  • 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.Source: B. Performance · p. 11

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.Source: C. Energy Efficiency · p. 11
  • 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.Source: C. Energy Efficiency · p. 11
  • 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.Source: C. Energy Efficiency · p. 11

B. Performance

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.Source: §IX · p. 12

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.Source: §IX · p. 12

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.Source: §X · p. 12
  • 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.Source: §X · p. 12

Related papers

Loading 1705.01626v1…