Introducing hipThreads: A C++ - Style Concurrency Library for AMD GPUs

2026-02-19 · Source: AMD ROCm Blogs · Field: Technology & Digital — Software Development & Engineering, Artificial Intelligence & Machine Learning · Depth: Intermediate, long

Summary

AMD has introduced hipThreads, a C++-style concurrency library designed to accelerate C++ code on AMD GPUs by enabling incremental porting of `std::thread` patterns to GPU-resident `hip::thread` code. This library aims to lower the barrier to GPU programming by providing familiar C++ concurrency primitives like `hip::mutex` and `hip::condition_variable`. hipThreads operates by launching a persistent GPU manager that handles a work queue, eliminating the overhead of constant kernel launches. It supports multi-fiber execution, where each `hip::thread` can utilize up to 32 fibers for SIMD parallelism. Benchmarks demonstrate significant performance gains, with a SAXPY operation showing a 6.4x speedup (from 271.88ms on a Ryzen™ 9 9900X to 42.60ms on an AMD Radeon™ AI PRO R9700) and other workloads achieving 2.9-3.6x speedups, often with minimal code changes (e.g., 16 lines for SAXPY).

Key takeaway

For AI Engineers and Software Engineers porting C++ applications to AMD GPUs, hipThreads offers a streamlined path to performance gains. You can incrementally adapt existing `std::thread` code to `hip::thread` with minimal changes, leveraging familiar concurrency primitives and achieving significant speedups (e.9., 2.9-6.4x) without a complete mental model shift. Consider adopting hipThreads to accelerate your C++ workloads on AMD GPUs, especially if you prioritize maintainability and a reduced learning curve.

Key insights

hipThreads enables incremental GPU acceleration of C++ `std::thread` code using familiar concurrency patterns and persistent multi-fiber execution.

Principles

• Familiar APIs reduce cognitive load.
• Persistent GPU schedulers minimize overhead.
• Multi-fiber execution leverages SIMD parallelism.

Method

Port CPU `std::thread` code to GPU by replacing `std::thread` with `hip::thread`, managing device memory, and optimizing with `hip::thread::max_width()` for multi-fiber parallelism.

In practice

• Use `thrust::unique_ptr` for GPU memory allocation.
• Annotate GPU functions with `__device__`.
• Distribute work across fibers using `get_fiber_id()` and `get_width()`.

Topics

• hipThreads
• GPU Programming
• C++ Concurrency
• AMD ROCm
• Multi-fiber Execution

Code references

• ROCm/hipThreads
• ROCm/libhipcxx
• ROCm/hipThreads

Best for: Software Engineer, AI Engineer, AI Architect

Related on AIssential

• Unlocking Sparse Acceleration on AMD GPUs with hipSPARSELt — hipSPARSELt accelerates sparse matrix operations on AMD GPUs using 2:4 structured sparsity for improved AI model efficiency.
• GPU Scaling from 1 to 261120 threads Part 1 — GPU performance scales non-linearly with threads due to memory latency and inefficient warp utilization.
• Taking on CUDA With ROCm: ‘One Step After Another’ — AMD's ROCm software stack is rapidly maturing, leveraging open-source tools and community engagement to challenge Nvidia's CUDA dominance.
• Unlocking asynchronicity in continuous batching — Asynchronous batching separates CPU and GPU workloads to maximize parallel execution and GPU utilization during LLM inference.
• Building and Deploying Custom hipBLASLt Libraries on AMD Instinct GPUs — Custom "hipBLASLt" builds offer precise GPU optimization and flexible deployment for AMD Instinct GPUs, avoiding system-wide conflicts.
• Using Accelerated Computing to Live-Steer Scientific Experiments at Massive Research Facilities — GPU-accelerated computing enables real-time analysis and steering of exascale scientific experiments, drastically reducing data processing times.

Open in AIssential →

Editorial summary, takeaway, and curation by AIssential. Original article published by AMD ROCm Blogs.