`vllm.models.muse_spark.nvidia.specialized_decode.cutedsl.common.primitives` ¶

PTX building blocks and launch helpers of the Muse Spark specialized decode.

Every Lamport kernel of the package (the boundary gathers, the attention boundary, the MoE tail and the unpooler's zero-copy epilogue) is built from these. BF16 negative zero (0x8000) marks a slot that has not been written yet; a three-generation stage state (NEXT_STAGE written by the consumer, ACTIVE_STAGE by the publisher) lets a fast rank publish the next call while a slow one still reads the current one.

Functions:

add_ftz_f32 –

A + b, round to nearest, subnormal inputs and result flushed to a zero
add_rn –

A + b, round to nearest, subnormals kept.
atomic_add1_acq_rel_u32 –

Add 1 to the global word at pointer and return its value before
atomic_add1_acq_rel_u32_address –

atomic_add1_acq_rel_u32 of the global word at address.
atomic_add1_u32_address –

atomicAdd(address, 1) (relaxed, GPU scope): the word before.
atomic_exch_zero_u32_address –

atomicExch(address, 0) (relaxed, GPU scope).
bf16x2_to_f32x2 –

The BF16 pair in word widened to FP32 (exact): (bits 0-15, 16-31).
cp_async_cg_16 –

Copy 16 bytes from global source to this CTA's shared
cp_async_wait –

Wait until at most pending of this thread's cp.async groups are in
div_full_f32 –

A / b as div.full.f32 (Triton's FP32 division).
ex2_approx_ftz –

2 ** a by ex2.approx, subnormals flushed.
f32_bits –

The bit pattern of an FP32 value.
fence_mbarrier_init_cluster –

Release this CTA's mbarrier initializations to the cluster.
fma_ftz_f32 –

A * b + c with one rounding, subnormals flushed.
fma_rn –

A * b + c with one rounding.
fms_rn –

A * b - c with one rounding (the negation is exact, -0 included).
global_timer_ns –

The GPU's global nanosecond timer (bounds the Lamport spins).
ldmc_bf16x8 –

Reduce-load eight BF16s; an omitted predicate loads, false returns zeros.
load_f16x8_as_f32 –

8 FP16 elements from element offset of base (one 16-byte
load_global_cg_f32_address –

An FP32 load cached in L2 only (another CTA's store, after a fence).
load_global_cg_f32x2_address –

load_global_cg_f32_address of two FP32.
load_global_f32x4_address –

Four FP32 at the 16-byte aligned address in one load; ordered
load_global_nc_f32_address –

A read-only FP32 load through the non-coherent path, in program order.
load_global_nc_s32_or –

The read-only word at address if predicate, else other.
load_global_nc_u32x4_address –

Four words of data no thread of the grid writes (the read-only path),
load_global_u32x4 –

Load four words; an omitted predicate loads normally, false returns zeros.
load_shared_f32_or –

The FP32 at shared address if predicate, else other.
load_shared_u32x4_address –

load_shared_u32x4 by address.
load_shared_u32x4_unpacked –

load_shared_u32x4_address as a tuple of four Uint32.
lowest_set_bit –

The index of the lowest set bit of a nonzero mask (CUDA's
map_shared_address_to_peer –

The shared::cluster address of this CTA's shared address in the
max_ftz_f32 –

The larger of a and b (a NaN operand yields the other), subnormals
mbarrier_arrive_expect_tx_address –

Arrive on the shared mbarrier at bar expecting num_bytes
mbarrier_init_address –

Initialize the shared mbarrier at bar for count arrivals (a
mbarrier_try_wait_phase0_cluster –

1 once phase 0 of the shared mbarrier at address has completed
mbarrier_wait_address –

Spin until the phase of parity completes (acquire, CTA scope).
mbarrier_wait_cluster_address –

mbarrier_wait_address with cluster-scope acquire: for a phase that
membar_gl –

The fence of __threadfence().
mma_m16n8k16_bf16 –

mma.sync m16n8k16 row.col with BF16 operands and FP32 accumulation: the
movmatrix_trans_b16 –

This lane's word of the warp's 8x8 b16 matrix, transposed.
mul_ftz_f32 –

A * b, round to nearest, subnormals flushed; never contracted.
mul_rn –

A * b, round to nearest, subnormals kept; never contracted.
pack_bf16x2_rn –

Two FP32 rounded to nearest-even BF16, lo in bits 0-15, hi in bits 16-31.
pack_bf16x8_rn –

8 FP32 scalars rounded to nearest-even BF16, as 4 packed words (an
rcp_approx_ftz –

1 / a by rcp.approx, subnormals flushed (the fast-math 1 / a).
regs –

An 8-wide FP32 TensorSSA (load_f32x8, load_bf16x8_as_f32) as
round_bf16 –

A rounded to BF16 (nearest even), widened back.
rsqrt_approx_f32 –

rsqrt.approx with subnormal inputs kept (rsqrtf without fast math).
rsqrt_approx_ftz –

rsqrt.approx with subnormals flushed (libdevice's rsqrt under FTZ).
scale_or_neg_inf –

scale * x (round to nearest, subnormals flushed) for a key in
sigmoid_f32 –

tl.sigmoid(x): x * -log2(e), ex2.approx, + 1, then 1 / that with
stmc_bf16x8 –

Multicast-store eight BF16s; an omitted predicate stores, false skips.
store_async_cluster_f32x2 –

Store (lo, hi) at cluster address remote
store_async_cluster_f32x4 –

Four FP32 into a cluster CTA's shared memory
store_bf16x8_rn –

8 FP32 scalars rounded to nearest-even BF16 and stored as one 16-byte
store_global_f32_address –

Store value at global address if predicate.
store_shared_f32_address –

Store value at shared address if predicate.
sub_ftz_f32 –

A - b, round to nearest, subnormals flushed.
tma_load_2d –

The box of tensor_map at (c0, c1) (inner coordinate first) into
tma_prefetch_2d_l2 –

Prefetch the box of tensor_map at (c0, c1) into L2.
trap –

Abort the kernel: the launch fails instead of hanging.