Skip to main content

Module cp_async

Module cp_async 

Source
Expand description

cp.async pipeline macro shim.

cp.async (sm_80+) and the Hopper-introduced cp.async.bulk variants live entirely on the device side. This module provides Rust constants for the macro names defined in atomr_hopper.cuh plus host-side helpers that compute the right mbarrier arrival count for a given pipeline stage count.

Modules§

macro_names
Macro names exposed by atomr_hopper.cuh for callers to reference in their NVRTC sources.

Enums§

PipelineStages
Pipeline stage policy for cp.async-driven shared-memory double/triple/quad buffering.

Functions§

mbarrier_arrival_count
Number of mbarrier arrival slots needed to fence a pipelined producer/consumer with the given stages. Matches the formula 2 * stages used by stage-balanced double-buffer kernels.