Logo
ScienceStack
Table of Contents
Fetching ...
Sign In

Optimizing sDTW for AMD GPUs

Daniel Latta-Lin, Sofia Isadora Padilla Munoz

TL;DR

This work presents a ROCm/HIP-based, AMD GPU-optimized implementation of subsequence Dynamic Time Warping (sDTW). It introduces two kernel modules: a normalizer for z-score standardization and a specialized sDTW kernel that uses per-wavefront processing, warp shuffles for inter-segment communication, and two buffers to transfer state across wavefronts, all operating on float16 with __half2 packing. The approach builds on prior GPU DTW techniques, adapting them to ROCm with careful attention to memory access, synchronization, and segment-width tuning; results show correct outputs against CPU references and substantial throughput improvements, achieving 9.26e-4 Gsps for sDTW and 4.81 Gsps for normalization under a batch of 512 queries of length 2,000 against a 100,000-length reference. The work demonstrates portability of DTW optimizations to AMD hardware and outlines potential future gains through quantization and pruning strategies to further boost performance. The practical impact lies in enabling real-time, large-scale sequence alignment on AMD GPUs for applications like finance, genomics, and signal processing.

Abstract

Subsequence Dynamic Time Warping (sDTW) is the metric of choice when performing many sequence matching and alignment tasks. While sDTW is flexible and accurate, it is neither simple nor fast to compute; significant research effort has been spent devising parallel implementations on the GPU that leverage efficient memory access and computation patterns, as well as features offered by specific vendors and architectures (notably NVIDIA's). We present an implementation of sDTW on AMD hardware using HIP and ROCm. Our implementation employs well-known parallel patterns, as well as lower-level features offered by ROCm. We use shuffling for intra-wavefront communication and shared memory to transfer data between consecutive wavefronts. By constraining the input data to batches of 512 queries of length 2,000, we optimized for peak performance the width of reference elements operated on by a single thread.

Optimizing sDTW for AMD GPUs

TL;DR

This work presents a ROCm/HIP-based, AMD GPU-optimized implementation of subsequence Dynamic Time Warping (sDTW). It introduces two kernel modules: a normalizer for z-score standardization and a specialized sDTW kernel that uses per-wavefront processing, warp shuffles for inter-segment communication, and two buffers to transfer state across wavefronts, all operating on float16 with __half2 packing. The approach builds on prior GPU DTW techniques, adapting them to ROCm with careful attention to memory access, synchronization, and segment-width tuning; results show correct outputs against CPU references and substantial throughput improvements, achieving 9.26e-4 Gsps for sDTW and 4.81 Gsps for normalization under a batch of 512 queries of length 2,000 against a 100,000-length reference. The work demonstrates portability of DTW optimizations to AMD hardware and outlines potential future gains through quantization and pruning strategies to further boost performance. The practical impact lies in enabling real-time, large-scale sequence alignment on AMD GPUs for applications like finance, genomics, and signal processing.

Abstract

Subsequence Dynamic Time Warping (sDTW) is the metric of choice when performing many sequence matching and alignment tasks. While sDTW is flexible and accurate, it is neither simple nor fast to compute; significant research effort has been spent devising parallel implementations on the GPU that leverage efficient memory access and computation patterns, as well as features offered by specific vendors and architectures (notably NVIDIA's). We present an implementation of sDTW on AMD hardware using HIP and ROCm. Our implementation employs well-known parallel patterns, as well as lower-level features offered by ROCm. We use shuffling for intra-wavefront communication and shared memory to transfer data between consecutive wavefronts. By constraining the input data to batches of 512 queries of length 2,000, we optimized for peak performance the width of reference elements operated on by a single thread.
Paper Structure (10 sections, 3 equations, 3 figures, 1 table)

This paper contains 10 sections, 3 equations, 3 figures, 1 table.

Table of Contents

  1. Introduction
  2. Background
  3. Prior Works
  4. Methodology
  5. Implementation
  6. Normalizer implementation
  7. sDTW implementation
  8. Results
  9. Conclusion
  10. Discussion

Figures (3)

  • Figure 1: On each iteration threads calculate the cost values for the next row in their segment. The right most value is then propagated up to be used by the next thread in the following iteration.
  • Figure 2: Once a thread is finished with it's segment the found minimum result for the segment is propagated up for further enrichment. The shared memory buffer allows for communication between a completed and newly started wavefront.
  • Figure 3: Initially increasing segment width improves performance but any gains are soon lost once segment width grows past 14 values wide.