API reference guide

API reference guide#

Threading#

The threading groups cover threads, mutexes, and condition variables.

group threading: Thread objects, synchronization and waiting primitives adapted for HIP execution.

Thread management#

group thread

Creation, identification and control of GPU threads.

Functions

template<class _CharT, class _Traits> ::std::basic_ostream<_CharT, _Traits> &operator<<(::std::basic_ostream<_CharT, _Traits> &__os, hip::__thread_id __id)#

Stream insertion for cuda::__thread_id.

Produces a textual representation; distinct ids yield distinct text, equal ids yield identical text (except default constructed which prints 0).

class __thread_id#

#include <id.h>

Opaque handle identifying a logical GPU thread (work node lane).

Acts similarly to std::thread::id but for the hipThreads runtime. A default constructed id (value 0) represents “no thread” and is always ordered before any non‑default id. Instances are obtained via:

cuda::this_thread::get_id()
cuda::wthread::get_id()

Ordering:

All non‑zero ids are ordered by underlying integral value.
Zero (sentinel) < any non‑zero.

Lifetime:

An id may outlive the associated thread; equality remains valid for comparison but does not retain resources.

Public Functions

__host__ __device__ inline void __reset()#: Resets this id back to the sentinel (value 0).

__host__ __device__ inline __thread_id() noexcept#: Constructs a sentinel (non-joinable / no-thread) id (value 0).

Friends

template<class _CharT, class _Traits> friend ::std::basic_ostream<_CharT, _Traits> &operator<<(::std::basic_ostream<_CharT, _Traits> &__os, __thread_id __id)#

Stream insertion for cuda::__thread_id.

Produces a textual representation; distinct ids yield distinct text, equal ids yield identical text (except default constructed which prints 0).

class wthread#

#include <thread.h>

Handle owning (at most) one scheduled GPU work node (wave of execution).

Constructing schedules a callable for execution (host or device path). Non-copyable; movable to transfer ownership. Must be joined or detached before destruction if joinable.

Memory / transfer constraints:

Captured callable and argument types must be trivially copyable (host path) or trivially destructible (device path) as enforced by static assertions.

Concurrency model:

join() waits for completion (synchronizing-with the end of the callable).
detach() releases association; work may continue without the handle.

Width:

Optional first constructor parameter width ( <= max_width() ) describes how many lanes participate; implementation may map this to a warp subset.

Deleted copy / move operations

These special members are intentionally disabled (handle is non-copyable / non-movable).

Public Types

using id = __thread_id #: Alias for thread identifier type (may be refined later).

Public Functions

__host__ __device__ void detach()#

Releases ownership allowing work to proceed independently.

After detach() the handle becomes not joinable.

__host__ __device__ wthread::id get_id(uint32_t index = 0) const#

Returns the id of the (possibly width-partitioned) logical lane.

Parameters:: index – Lane index (default 0).

__host__ __device__ void join()#

Waits for completion of the associated work.

Undefined behavior if !joinable().

__host__ __device__ inline bool joinable() const noexcept#: Returns true if an execution context is owned and not yet joined/detached.

__host__ __device__ wthread &operator=(wthread &&other) noexcept#: Move assignment transfers ownership; source becomes not joinable.

__host__ __device__ inline void swap(wthread &__t) noexcept#: Swaps underlying work node ownership with another wthread.

__host__ wthread() noexcept#: Default constructs a non-joinable wthread (no associated work node).

__device__ inline wthread() noexcept: Device-side default constructor (no work node).

template<class Fn_t, class ...Args_t, ::std::enable_if_t<!::std::is_arithmetic_v<::std::remove_reference_t<Fn_t>>, bool> = true> __device__ inline explicit wthread(Fn_t &&typed_fn, Args_t&&... args)#: Device-side convenience constructor (width = 1) for initial drop-in replacement of std::thread.

template<class Fn_t, class ...Args_t, ::std::enable_if_t<!::std::is_arithmetic_v<::std::remove_reference_t<Fn_t>>, bool> = true> __host__ inline explicit wthread(Fn_t &&typed_fn, Args_t&&... args): Host-side convenience constructor (width = 1) for initial drop-in replacement of std::thread.

template<class Fn_t, class ...Args_t> __device__ inline explicit wthread(uint32_t width, Fn_t &&typed_fn, Args_t&&... args)#

Construct with explicit width and callable (device path).

Parameters:

width – Logical participation width (1..max_width()).
typed_fn – Callable object.
args – Argument pack forwarded to callable.

template<class Fn_t, class ...Args_t> __host__ inline explicit wthread(uint32_t width, Fn_t &&typed_fn, Args_t&&... args)

Construct with explicit width and callable (host path).

Schedules work node for device execution.

__host__ __device__ inline wthread(wthread &&other) noexcept#: Move construction transfers ownership; source becomes not joinable.

__host__ __device__ ~wthread()#: Destructor: wthread must be not joinable (joined or detached).

Public Static Functions

__device__ static unsigned int hardware_concurrency() noexcept#: Number of concurrent hardware slots usable (device variant).

__host__ static unsigned int hardware_concurrency() noexcept: Number of concurrent hardware slots usable (host reflection / cached).

__host__ __device__ static inline constexpr unsigned int max_width() noexcept#: Maximum supported width (implementation constant).

Mutexes#

group mutex

Mutual exclusion types preventing simultaneous access to shared data.

Includes a spinning mutex (spin_mutex) and RAII lock wrappers (lock_guard, unique_lock).

template<class _Mutex> class lock_guard#

#include <lock_guard.h>

Scoped non-copyable guard that owns a mutex for its lifetime.

Guarantees:

Acquires the mutex in the locking constructor.
Releases the mutex in the destructor.
Non-copyable and non-assignable to prevent multiple owners.

Template Parameters:: _Mutex – Mutex type meeting BasicLockable (lock(), unlock()).

Public Types

typedef _Mutex mutex_type#: Alias for the underlying mutex type.

Public Functions

__device__ inline explicit lock_guard(mutex_type &__m)#

Locks the given mutex.

Parameters:: __m – Mutex to lock (must not already be locked by this wthread).

__device__ inline lock_guard(mutex_type &__m, ::std::adopt_lock_t)#

Adopts ownership of an already-locked mutex.

Parameters:

__m – Mutex already locked by the calling wthread.
(adopt_lock) – Tag indicating adoption (no lock attempt made).

Pre:

The calling wthread holds the mutex.

__device__ inline ~lock_guard()#: Unlocks the mutex.

class pseudo_mutex#

#include <pseudo_mutex>

Busy-wait (spin) mutex with periodic pseudo_yield.

Characteristics:

Exclusive, non‑recursive.
lock(): loops atomicCAS; on repeated failure yields every 0x10000 iterations.
try_lock(): single CAS attempt.
unlock(): __threadfence() then atomicExch to release (establishes release/acquire ordering with next successful lock).

Ownership model:

Tracks owning block via computed block id (not per-thread recursion aware).

When to choose:

Prefer spin_mutex for short sections with minimal contention.
Use pseudo_mutex only when you can guarantee no yield-loops will occur.

Warning

Avoid pseudo_mutex if yield-loops are possible. A deadlock occurs when thread A holds the lock and yields to thread B, then thread B attempts to acquire the same lock. Since pseudo_yield() prevents resumption until the yieldee completes, thread B cannot yield back to thread A, causing permanent deadlock. Only use pseudo_mutex when you can guarantee this scenario won’t occur.

Deleted copy / move operations

Instances are neither copyable nor movable.

Public Functions

__device__ inline void lock()#

Acquire (spins with periodic pseudo_yield).

Yields every 0x10000 failed attempts. Debug assert triggers on recursive attempt or attempting to acquire a lock already owned by another fiber in the same SIMD wave/hip::wthread.

__device__ constexpr pseudo_mutex() = default#: Constructs unlocked.

__device__ inline bool try_lock() noexcept#

Try to acquire without spinning.

Provides acquire semantics on success; no synchronizes-with edge on failure.

Returns:: true on success, false if already owned.

__device__ inline void unlock() noexcept#

Release ownership.

Issues device fence to publish prior writes, then clears owner with atomicExch. Debug assert validates current block was the owner.

class spin_mutex#

#include <spin_mutex.h>

Busy‑wait (spin) non‑recursive mutex.

Characteristics:

Exclusive, non‑recursive ownership (re‑entering is UB / debug assert).
lock() spins with atomicCAS until acquired.
try_lock() single non‑blocking attempt.
unlock() issues a release fence (__threadfence) then clears ownership.

GPU specifics:

It is illegal for more than one fiber to be active while acquiring the lock (debug assert or livelock). This results in the extra fibers looping, trying to acquire the lock when it’s already owned by a fiber in the same SIMD wave/hip::wthread.

Ownership is tracked using the block id so that we can attempt to detect this invalid use in debug and assert instead of livelocking.

Prolonged contention wastes execution resources (avoid long critical sections).

Use RAII helpers (lock_guard / unique_lock) for safer scope management.

Not copyable or movable.

Deleted copy / move operations

Instances are neither copyable nor movable.

Public Functions

__device__ inline void lock()#

Acquires the mutex (spins until success).

Uses atomicCAS with acquire semantics to claim the owner slot. Undefined behavior (asserts in debug) if the same block attempts to re‑enter (non‑recursive) or more than one fiber is active while acquiring..

__device__ constexpr spin_mutex() = default#: Constructs an unlocked spin_mutex.

__device__ inline bool try_lock() noexcept#

Attempts to acquire without spinning.

No memory ordering relationship is created with prior failed try_lock attempts or successful lock() calls that did not transfer ownership to this block (mirrors std::mutex semantics).

Returns:: true if ownership obtained; false if already locked.

__device__ inline void unlock() noexcept#

Releases ownership.

Performs a release fence ensuring all writes inside the critical section become visible to a subsequent acquiring block before clearing owner. Debug builds assert the current block owns the mutex.

template<class _Mutex> class unique_lock#

#include <unique_lock.h>

Movable (not copyable) lock owning at most one mutex.

Features:

Multiple constructor tags: (default), defer_lock, try_to_lock, adopt_lock.
Manual lock()/unlock()/try_lock().
Release without unlocking (release()).
Ownership query (owns_lock() / operator bool()).
Move construct / assign transfers ownership; source is left empty.

Invariants:

If owns_lock() is true then mutex() is non-null and locked by this wthread.
Destruction unlocks only if ownership is held.

Timed locking hooks are commented out (TODO) — will forward to underlying mutex once TimedLockable support is implemented.

Template Parameters:: _Mutex – Mutex type meeting BasicLockable (lock(), unlock()).

Deleted copy operations

Instances are movable but not copyable.

Public Types

typedef _Mutex mutex_type#: Underlying mutex type.

Public Functions

__device__ void lock()#

Locks the associated mutex (must be associated & not owned).

Throws:: (asserts – in device build on misuse).
Post:: owns_lock() == true

__device__ inline mutex_type *mutex() const noexcept#: Pointer to the associated mutex (may be null if default-constructed or released).

__device__ inline explicit operator bool() const noexcept#: Same as owns_lock().

__device__ inline unique_lock &operator=(unique_lock &&__u) noexcept#

Move assigns, releasing any owned mutex then taking __u’s.

Parameters:: __u – Source; becomes empty.
Returns:: *this

__device__ inline bool owns_lock() const noexcept#: True if this object currently owns the mutex.

__device__ inline mutex_type *release() noexcept#

Releases association without unlocking.

Returns:: Previously associated mutex pointer (may be null).
Post:: mutex()==nullptr && owns_lock()==false

__device__ inline void swap(unique_lock &__u) noexcept#

Exchanges state with another unique_lock.

Parameters:: __u – Other lock.
Post:: Ownership flags and mutex pointers swapped.

__device__ bool try_lock()#

Attempts to lock without blocking.

Returns:: true if the lock was obtained.
Post:: owns_lock() reflects result.

__device__ inline unique_lock() noexcept#: Constructs an empty non-owning lock (mutex() == nullptr).

__device__ inline explicit unique_lock(mutex_type &__m)#

Locks the supplied mutex immediately.

Parameters:: __m – Target mutex (must not already be locked by this wthread).
Post:: owns_lock() == true

__device__ inline unique_lock(mutex_type &__m, ::std::adopt_lock_t)#

Assumes caller already holds the mutex.

Parameters:

__m – Target mutex already locked by this wthread.
(adopt_lock) – Tag asserting prior lock ownership.

Pre:

Current wthread owns __m.

__device__ inline unique_lock(mutex_type &__m, ::std::defer_lock_t) noexcept#

Associates with a mutex without locking it yet.

Parameters:

__m – Target mutex.
(defer_lock) – Tag indicating deferred acquisition.

Post:

owns_lock() == false

__device__ inline unique_lock(mutex_type &__m, ::std::try_to_lock_t)#

Attempts to lock without blocking.

Parameters:

__m – Target mutex.
(try_to_lock) – Tag requesting a non-blocking attempt.

Post:

owns_lock() reflects success.

__device__ inline unique_lock(unique_lock &&__u) noexcept#

Move constructs, transferring ownership.

Parameters:: __u – Source; becomes empty afterward.

__device__ void unlock()#

Unlocks the mutex.

Pre:: owns_lock() == true
Post:: owns_lock() == false

__device__ inline ~unique_lock()#: Destructor unlocks if ownership is held.

Condition variables#

group condition_variable

Blocking/waiting mechanisms allowing one or more threads to wait until they are notified.

Includes standard semantics (condition_variable_any) and a lightweight spinning variant (spin_condition_variable).

class condition_variable_any#

#include <condition_variable_any.h>

Condition variable that works with any BasicLockable lock.

Maintains separate counters for waiters and notifications. Each wait() obtains an arrival index; notify operations advance a notification counter.

Note

Not copyable or movable.

Warning

Implementation spins while waiting; avoid long waits to prevent excessive resource usage.

Subclassed by hip::spin_condition_variable

Deleted copy / move operations

Condition variable objects are neither copyable nor movable.

Public Functions

__device__ constexpr condition_variable_any() noexcept = default#: Constructs an empty condition variable (no waiters).

__device__ inline void notify_all() noexcept#

Wakes all current waiters.

Advances the notification counter to match the waiter counter so every waiter observing the value will proceed.

__device__ inline void notify_one() noexcept#

Wakes at most one waiting wthread/work item (if any).

If no waiter is currently blocked, the call is a no-op (no “stored” wake token).

template<class _Lock> __device__ void wait(_Lock &__lock)#

Blocks (spins) until notified.

Atomically:

Records arrival index.
Releases the supplied lock (lock.unlock()).
Spins until its index is < current notify counter.
Reacquires the lock before returning.

Warning

Spurious wakeups are possible—prefer the predicate form.

Template Parameters:: _Lock – Lock type supporting unlock() / lock().
Parameters:: __lock – Acquired lock protecting the predicate.

template<class _Lock, class _Predicate> __device__ inline void wait(_Lock &__lock, _Predicate __pred)#

Waits until predicate returns true.

Repeatedly invokes the simple wait() and rechecks __pred() under the lock. Returns only when __pred() evaluates to true.

Template Parameters:

_Lock – Lock type.
_Predicate – Callable returning bool, evaluated with the lock held.

Parameters:

__lock – Acquired lock.
__pred – Predicate defining the wake condition.

class pseudo_condition_variable#

#include <pseudo_condition_variable>

Spinning condition variable with periodic pseudo_yield.

Mechanism:

Each waiter atomically takes a ticket (wait_counter).
Notifiers advance notify_counter.
A waiter proceeds when its ticket < notify_counter.
During the spin loop, every K iterations (default 0x10000) a hip::this_thread::pseudo_yield() is issued to reduce wasted GPU resources.

Guarantees / caveats:

No true sleep; still busy-waiting between yields.
The “unlock and sleep” operation in wait() isn’t technically atomic. A notify that occurs between taking the ticket and unlocking may “wake” the wait-er before it unlocked.
Predicate form recommended to guard against spurious wakeups.

Deleted copy / move operations

Instances are neither copyable nor movable.

Public Functions

__device__ inline void notify_all() noexcept#

Wakes all current waiters.

Sets notify_counter to wait_counter so every waiter observes progress.

__device__ inline void notify_one() noexcept#

Wakes at most one waiter (if any present).

Increments notify_counter iff there is at least one unmatched wait ticket. No effect if no current waiters.

__device__ constexpr pseudo_condition_variable() noexcept = default#: Constructs an empty pseudo_condition_variable.

template<class _Lock> __device__ inline void wait(_Lock &__lock)#

Waits until notified (may spuriously wake).

Steps:

Atomically obtain ticket (wait_counter++).
Unlock supplied lock.
Spin until ticket < notify_counter; pseudo_yield periodically.
Re-lock before returning.

Template Parameters:: _Lock – BasicLockable providing unlock()/lock().
Parameters:: __lock – Held lock protecting the predicate.

template<class _Lock, class _Predicate> __device__ inline void wait(_Lock &__lock, _Predicate __pred)#

Waits until predicate returns true.

Repeats wait(lock) while !pred().

Template Parameters:

_Lock – BasicLockable.
_Predicate – Callable returning bool.

Parameters:

__lock – Held lock.
__pred – Predicate to satisfy.

class spin_condition_variable : private hip::condition_variable_any #

#include <spin_condition_variable.h>

Spin-based condition variable bound to spin_mutex.

Privately inherits condition_variable_any to reuse its counter logic, re-exposing notify and tailored wait overloads for unique_lock<spin_mutex>.

Note

Not copyable.

Warning

Pure spinning; avoid long waits to reduce resource burn.

Deleted copy / move operations

Instances are neither copyable nor movable.

Public Functions

__device__ inline void notify_all() noexcept#

Wakes all current waiters.

Advances the notification counter to match the waiter counter so every waiter observing the value will proceed.

__device__ inline void notify_one() noexcept#

Wakes at most one waiting wthread/work item (if any).

If no waiter is currently blocked, the call is a no-op (no “stored” wake token).

__device__ constexpr spin_condition_variable() noexcept = default#: Constructs an empty spin condition variable.

__device__ inline void wait(unique_lock<spin_mutex> &__lk) noexcept#

Waits (spins) until notified.

Releases the lock, spins polling the internal counters, then reacquires before returning.

Parameters:: __lk – Acquired unique_lock guarding the predicate.

template<class _Predicate> __device__ inline void wait(unique_lock<spin_mutex> &__lk, _Predicate __pred)#

Waits until predicate returns true.

Repeatedly performs wait() then rechecks __pred() under the lock.

Template Parameters:

_Predicate – Callable returning bool.

Parameters:

__lk – Acquired unique_lock.
__pred – Predicate tested after each wake/spin cycle.

C library utilities#

The C library groups cover memory allocation and byte and string manipulation.

group c_library: Wrappers mirroring portions of the C library with HIP/GPU semantics.

Memory allocation#

group c_memory

Malloc/free style GPU allocation helpers.

Functions

__host__ inline void free(void *ptr)#

Asynchronously frees memory obtained from hip::malloc.

Enqueues hipFreeAsync on the internal stream and returns immediately. Passing nullptr is a no-op (matches free semantics).

Parameters:: ptr – Pointer previously returned by hip::malloc or nullptr.

__host__ inline void *malloc(::std::size_t size)#

Allocates size bytes of uninitialized GPU-accessible memory.

Uses hipMallocAsync on an internal non-blocking stream, then synchronizes that stream so the returned pointer is ready for immediate use.

Differences vs C malloc:

Throws on HIP failure instead of returning nullptr.
Guarantees stream sync before return (blocking semantics).

Note

Unlike hipMalloc, this function is safe to call while hip::wthread objects are alive. It uses hipMallocAsync on a non-blocking stream, avoiding the deadlock that would occur with synchronous HIP APIs when the hipThreads persistent scheduler holds the GPU context.

Parameters:: size – Number of bytes to allocate.
Throws:: HIP – error via __LIBHIPTHREADS_HIP_CHECK__ on failure.
Returns:: Pointer to at least size bytes.

Byte and string manipulation#

group c_bytestring

Low-level raw memory / byte sequence utilities.

Functions

__host__ __device__ inline void *memcpy(void *dest, const void *src, ::std::size_t count)#

Copies a block of memory from a source address to a destination address.

Copies count bytes from the object pointed to by src to the object pointed to by dest. Both objects are reinterpreted as arrays of unsigned char.

This version is available for use in both __host__ and __device__ code.

Warning

The behavior is undefined if the memory areas of src and dest overlap.

Parameters:

dest – Pointer to the destination array where the content is to be copied.
src – Pointer to the source of data to be copied.
count – Number of bytes to copy.

Returns:

A copy of dest.

API reference guide

Contents

API reference guide#

Threading#

Thread management#

Mutexes#

Condition variables#

C library utilities#

Memory allocation#

Byte and string manipulation#