Adding New Intrinsics#

So you want to teach cuda-oxide a new GPU trick. Maybe NVIDIA just shipped a new instruction, or you need an existing PTX operation that nobody has wired up yet. Good news: the process is mechanical. Five crates, five steps, and roughly thirty minutes once you have done it before.

This chapter walks through the full pipeline using two real examples – one trivially simple, one with a few twists – so you can see exactly what happens at each stage.

The Five-Stage Pipeline#

Every GPU intrinsic follows the same path through the compiler:

cuda-device          User writes:  thread::threadIdx_x()
    │
    ▼
mir-importer       Compiler sees the call, emits a `dialect-nvvm` op
    │
    ▼
dialect-nvvm       The op lives here as a verified IR node
    │
    ▼
mir-lower          Converts the `dialect-nvvm` op into an LLVM dialect op
    │
    ▼
llvm-export        Exports textual LLVM IR  →  llc turns it into PTX

At the mir-lower stage you pick one of two strategies:

Strategy	When to use	Examples
LLVM intrinsic call	LLVM already has a built-in for it	`threadIdx_x`, warp shuffles, barriers
Inline PTX assembly	No LLVM intrinsic exists, or you need exact control over the PTX instruction	`trap`, `wgmma`, `tcgen05`, `mbarrier`

Both strategies are demonstrated below.

Example 1: `threadIdx_x` (the Simple Case)#

threadIdx_x() is the “Hello, World” of GPU intrinsics: zero arguments, one u32 result, maps directly to a single LLVM NVVM intrinsic. If you can follow this example, you can add any simple intrinsic.

Stage 1 – Declare in `cuda-device`#

File: crates/cuda-device/src/thread.rs

#[inline(never)]
pub fn threadIdx_x() -> u32 {
    unreachable!("threadIdx_x called outside CUDA kernel context")
}

Two rules that might look odd until you understand the trick:

#[inline(never)] keeps the function visible as a distinct call in MIR. If rustc inlined it, the compiler would see unreachable!() instead of a call it can intercept. That would be… less than helpful.
The body is unreachable!() because this function never actually runs. The compiler replaces the entire call with a dialect-nvvm operation before any GPU code executes. Think of the function as a placeholder – a “dear compiler, please insert a GPU instruction here” note.

Tip

Document what LLVM IR or PTX the intrinsic maps to in a comment above the function. Future you (or future contributors) will thank present you.

Stage 2 – Define the `dialect-nvvm` Op#

File: one of the hand-written op modules in crates/dialect-nvvm/src/ops/ – asm.rs, atomic.rs, cluster.rs, debug.rs, grid.rs, memory.rs, or wgmma.rs.

Important

Two kinds of op live in dialect-nvvm, and only one of them is written by hand.

Anything described by intrinsics/catalog.json is generated by cuda-intrinsics-gen into crates/dialect-nvvm/src/ops/generated/, where every file opens with // @generated by cuda-intrinsics-gen ... DO NOT EDIT. The op used as the example just below is one of them: tid.x is a catalog intrinsic and its op lives in ops/generated/sreg.rs today. Adding an op there by hand is undone by the next generator run.

The modules named above are the hand-written ones – ops with bespoke verification or lowering that the catalog does not describe. The code in this section shows the shape an op takes either way; write it yourself only for that second kind.

#[pliron_op(name = "nvvm.read_ptx_sreg_tid_x", dialect = "nvvm", format)]
pub struct ReadPtxSregTidXOp;

impl Verify for ReadPtxSregTidXOp {
    fn verify(&self, ctx: &Context) -> Result<(), Error> {
        let op = &*self.get_operation().deref(ctx);
        if op.get_num_operands() != 0 {
            return verify_err!(op.loc(), "expected 0 operands");
        }
        if op.get_num_results() != 1 {
            return verify_err!(op.loc(), "expected 1 result");
        }
        Ok(())
    }
}

Then register it so pliron knows the op exists:

pub(super) fn register(ctx: &mut Context) {
    ReadPtxSregTidXOp::register(ctx, ReadPtxSregTidXOp::parser_fn);
}

The Verify trait catches structural bugs early. If something accidentally creates this op with two operands, verification fails with a clear message instead of producing garbage PTX three stages later.

Stage 3 – Recognize in `mir-importer`#

File: crates/mir-importer/src/translator/terminator/intrinsics/generated.rs

When the translator processes MIR, every function call passes through try_dispatch_intrinsic() in terminator/mod.rs. The callee’s fully qualified domain name (FQDN) comes from CrateDef::name() via extract_func_info(), producing paths like cuda_device::thread::threadIdx_x. For a catalog intrinsic like tid.x, try_dispatch_intrinsic() defers to the generated dispatcher, whose match arm checks the FQDN:

"cuda_intrinsics::__cuda_oxide_intrinsic_abi_v1::i0001"
| "cuda_device::thread::threadIdx_x"
| "cuda_device::threadIdx_x" => {
    require_arity(name, args.len(), 0, &loc)?;
    Ok(Some(helpers::emit_generated_nvvm_intrinsic(
        ctx,
        ReadPtxSregTidXOp::get_concrete_op_info(),
        "v1:i0001",
        destination, target, block_ptr, prev_op,
        value_map, block_map, loc,
    )?))
}

The arm matches the full module path (cuda_device::thread::threadIdx_x), the re-exported name (cuda_device::threadIdx_x), and a versioned ABI alias, so the call is intercepted regardless of how the user imports the function. The FQDN is used as-is for matching – no :: to __ conversion happens before the intrinsic check.

The emit_generated_nvvm_intrinsic() helper works for any zero-argument, single-result NVVM intrinsic. It creates the operation, tags it with its catalog ABI marker, stores the result in the value map, and emits a branch to the next basic block. You never write this arm yourself: describing the intrinsic in intrinsics/catalog.json makes cuda-intrinsics-gen emit it. Hand-written arms in terminator/mod.rs exist only for what the catalog cannot describe.

Stage 4 – Lower to the LLVM dialect#

File: crates/mir-lower/src/convert/generated_intrinsics.rs

The op implements the MirToLlvmConversion op interface. For a catalog intrinsic the impl is generated alongside a shared converter:

#[op_interface_impl]
impl MirToLlvmConversion for ReadPtxSregTidXOp {
    fn convert(
        &self,
        ctx: &mut Context,
        rewriter: &mut DialectConversionRewriter,
        _operands_info: &OperandsInfo,
    ) -> Result<()> {
        convert_zero_operand_scalar_direct(
            ctx,
            rewriter,
            self.get_operation(),
            32,
            "llvm_nvvm_read_ptx_sreg_tid_x",
        )
    }
}

The shared converter emits the LLVM intrinsic call:

fn convert_zero_operand_scalar_direct(
    ctx: &mut Context,
    rewriter: &mut DialectConversionRewriter,
    op: Ptr<Operation>,
    width: u32,
    intrinsic_name: &str,
) -> Result<()> {
    let result_ty = IntegerType::get(ctx, width, Signedness::Signless);
    let function_ty = llvm_types::FuncType::get(ctx, result_ty.into(), vec![], false);
    let call = call_intrinsic(ctx, rewriter, op, intrinsic_name, function_ty, vec![])?;
    rewriter.replace_operation(ctx, op, call);
    Ok(())
}

Hand-written conversions for non-catalog ops follow the same pattern from convert/interface_impls.rs, delegating into the convert/intrinsics/ modules.

Note

LLVM intrinsic names use dots (llvm.nvvm.read.ptx.sreg.tid.x), but pliron identifiers cannot contain dots. Internally we use underscores (llvm_nvvm_read_ptx_sreg_tid_x). The export stage converts them back.

Stage 5 – Export (Nothing to Change)#

File: crates/llvm-export/src/export/ (the export module)

The CallOp exporter already handles the underscore-to-dot conversion:

let fixed_name = if name.starts_with("llvm_nvvm") {
    name.replace('_', ".")
} else {
    strip_device_prefix(&name)
};

Since threadIdx_x is not convergent (it is a per-thread register read, not a collective operation), there is nothing to add to is_convergent_intrinsic().

Final output:

declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
%v5 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()

After llc:

mov.u32  %r1, %tid.x;

Five stages, one PTX instruction. Not bad.

Example 2: `shuffle_xor` (the Complex Case)#

Warp shuffles are more interesting. The user’s shorthand passes two arguments, but the underlying LLVM intrinsic expects four. One extra (the member mask) is filled in at the source level, the other (the clamp) by the lowering. The operation is also convergent, meaning LLVM must not move, duplicate, or speculate it across control flow.

Stage 1 – Declare in `cuda-device`#

File: crates/cuda-device/src/warp.rs

#[inline(never)]
pub fn shuffle_xor_sync(mask: u32, var: u32, lane_mask: u32) -> u32 {
    let _ = (mask, var, lane_mask);
    unreachable!("shuffle_xor_sync called outside CUDA kernel context")
}

#[inline(always)]
pub fn shuffle_xor(var: u32, lane_mask: u32) -> u32 {
    shuffle_xor_sync(u32::MAX, var, lane_mask)
}

The intrinsic stub is the masked _sync form, same pattern as before; the let _ = (mask, var, lane_mask); suppresses unused-variable warnings – a small courtesy that costs nothing. The bare shuffle_xor shorthand is not a stub at all: it is an #[inline(always)] wrapper that supplies the full-warp mask (u32::MAX) and calls the stub. Only the _sync name is ever matched by the compiler; the wrapper is ordinary Rust code.

Stage 2 – Define the `dialect-nvvm` Op#

File: crates/dialect-nvvm/src/ops/generated/warp_shuffle.rs. The warp shuffles are catalog intrinsics like tid.x, so this op is generated rather than typed out; here is the real definition, condensed:

/// Exchange a 32-bit value with a named XOR partner lane.
///
/// Operands are `[member_mask, value, lane_or_delta]`.
/// Generated lowering inserts the fixed clamp required by the selected shuffle mode.
#[pliron_op(
    name = "nvvm.shfl_sync_bfly_i32",
    format,
    interfaces = [NOpdsInterface<3>, NResultsInterface<1>],
)]
pub struct ShflSyncBflyI32Op;

impl Verify for ShflSyncBflyI32Op {
    fn verify(&self, ctx: &Context) -> Result<(), Error> {
        let op = self.get_operation().deref(ctx);
        if op.get_num_operands() != 3 || op.get_num_results() != 1 {
            return verify_err!(
                op.loc(),
                "nvvm.shfl_sync_bfly_i32 requires exactly [member_mask, value, lane_or_delta] and one result"
            );
        }
        // ... i32 type checks on all three operands and the result ...
        Ok(())
    }
}

Three operands this time (member mask, value, and lane mask), one result.

Stage 3 – Recognize in `mir-importer`#

File: crates/mir-importer/src/translator/terminator/intrinsics/generated.rs

Important

The warp shuffles are catalog intrinsics, so their dispatch arms and their argument-translating emitters are generated into generated.rs, next to the ops from Stage 2. Nothing about a shuffle is written by hand. Hand-written dispatch lives in try_dispatch_intrinsic() in terminator/mod.rs and covers what the catalog cannot describe. For example, core::intrinsics::typed_swap_nonoverlapping (the primitive behind core::mem::swap) is matched there and lowered by the hand-written emit_typed_swap() in the same file as a load/load/store/store crossover over two non-overlapping pointers. No catalog entry could say that.

An intrinsic with arguments translates each of the user’s operands before building the op. Here is the generated arm for shuffle_xor_sync, condensed:

"cuda_device::warp::shuffle_xor_sync" => {
    require_arity(name, args.len(), 3, &loc)?;
    let (member_mask, last_op) =
        rvalue::translate_operand(ctx, body, &args[0], /* ... */)?;
    let (value, last_op) =
        rvalue::translate_operand(ctx, body, &args[1], /* ... */)?;
    let (lane_or_delta, last_op) =
        rvalue::translate_operand(ctx, body, &args[2], /* ... */)?;
    let shuffle = ShflSyncBflyI32Op::build(ctx, member_mask, value, lane_or_delta);
    helpers::set_generated_intrinsic_marker(ctx, shuffle, "v1:i0051");
    helpers::insert_op(ctx, shuffle, block_ptr, last_op);
    let result = shuffle.deref(ctx).get_result(0);
    Ok(Some(helpers::emit_store_result_and_goto(
        ctx, destination, result, target, /* ... */
        "shuffle_xor_sync call without target block",
    )?))
}

The arm translates the user’s MIR operands into pliron values – typically the results of mir.loads from each operand’s alloca slot, or mir.constants for literal arguments – and wires them into the NVVM operation. (These are not SSA values yet; pliron::opts::mem2reg will collapse the load/store chains once translation is complete.) For a catalog intrinsic you never write any of this: cuda-intrinsics-gen derives it from the entry in intrinsics/catalog.json.

Stage 4 – Lower to the LLVM dialect#

File: crates/mir-lower/src/convert/intrinsics/warp.rs

Here is where the fourth argument appears. The dialect-nvvm op carries three operands; the LLVM intrinsic for butterfly shuffle takes four:

User's shorthand: shuffle_xor(value, lane_mask)                     →  2 args
Intrinsic stub:   shuffle_xor_sync(mask, value, lane_mask)          →  3 args
                                   ^^^^
                                   u32::MAX from the Stage 1 wrapper
LLVM intrinsic:   shfl.sync.bfly.i32(mask, value, lane_mask, clamp) →  4 args
                                                             ^^^^^
                                                             always 31

The converter forwards the three operands and adds only the clamp constant:

pub(crate) fn convert_shuffle_i32(
    ctx: &mut Context,
    rewriter: &mut DialectConversionRewriter,
    op: Ptr<Operation>,
    _operands_info: &OperandsInfo,
    intrinsic_name: &str,
    clamp: i32,
) -> Result<()> {
    let i32_ty = IntegerType::get(ctx, 32, Signedness::Signless);

    let operands: Vec<_> = op.deref(ctx).operands().collect();
    if operands.len() != 3 {
        return pliron::input_err_noloc!(
            "Warp shuffle i32 requires 3 operands [mask, value, lane_or_delta]"
        );
    }
    let (mask, val, lane_or_delta) = (operands[0], operands[1], operands[2]);

    let clamp_val = create_i32_const(ctx, rewriter, clamp);

    let func_ty = llvm_types::FuncType::get(
        ctx,
        i32_ty.into(),
        vec![i32_ty.into(), i32_ty.into(), i32_ty.into(), i32_ty.into()],
        false,
    );

    let call_op = call_intrinsic(
        ctx, rewriter, op, intrinsic_name, func_ty,
        vec![mask, val, lane_or_delta, clamp_val],
    )?;
    rewriter.replace_operation(ctx, op, call_op);
    Ok(())
}

The mask is an ordinary operand here, not a converter-made constant: the Stage 1 wrapper baked in u32::MAX (all 32 lanes participate), and the masked _sync forms let the user pass any mask at all. The clamp value of 31, supplied per shuffle mode by the generated MirToLlvmConversion impl that calls this converter, means the shuffle spans the full warp width. It is a fixed property of the instruction mode, and exposing it to the user would just be noise.

Stage 5 – Export (Add Convergent)#

Warp shuffles are convergent – LLVM must not reorder them relative to control flow. The export step checks is_convergent_intrinsic():

fn is_convergent_intrinsic(name: &str) -> bool {
    name == "llvm.nvvm.barrier0"
        || name.starts_with("llvm.nvvm.shfl")      // shuffles
        || name.starts_with("llvm.nvvm.vote")       // votes
        || name.starts_with("llvm.nvvm.mbarrier")   // async barriers
        || name.starts_with("llvm.nvvm.cp.async.bulk")
        // ...
}

If your new intrinsic is convergent, add it here. If you forget, LLVM might hoist it out of an if block, and your warp-level code will produce wrong results or deadlock. Not great.

Final output:

declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)

%v8 = call i32 @llvm.nvvm.shfl.sync.bfly.i32(
    i32 -1, i32 %v3, i32 %v4, i32 31) #0

attributes #0 = { convergent }

After llc:

shfl.sync.bfly.b32  %r3, %r1, %r2, 31;

The Inline PTX Path#

Some operations do not have LLVM intrinsics. For those, we emit inline PTX assembly directly. The helper inline_asm_convergent() handles the boilerplate:

// wgmma fence: no inputs, no outputs, just a side effect
inline_asm_convergent(
    ctx, void_ty.into(), vec![],
    "wgmma.fence.sync.aligned;", ""
);

// mbarrier arrive: one input (pointer), one output (token)
inline_asm_convergent(
    ctx, i64_ty.into(), vec![ptr_val],
    "mbarrier.arrive.shared.b64 $0, [$1];", "=l,r"
);

The constraint string follows LLVM inline assembly syntax:

Constraint	Meaning
`=l`	Output: 64-bit register
`=r`	Output: 32-bit register
`r`	Input: 32-bit register
`l`	Input: 64-bit register
(empty)	No inputs or outputs (side-effect only)

The sideeffect convergent markers on the inline assembly tell LLVM to leave it alone – do not move it, do not delete it, do not duplicate it.

End-to-End: The Full Journey#

Here is every representation threadIdx_x passes through, top to bottom:

Rust:         thread::threadIdx_x()
MIR:          _3 = threadIdx_x() -> bb1
Pliron MIR:   %v = nvvm.read_ptx_sreg_tid_x : i32
Pliron LLVM:  %v = call i32 @llvm_nvvm_read_ptx_sreg_tid_x()
LLVM IR:      %v5 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
PTX:          mov.u32 %r1, %tid.x;

And shuffle_xor_sync, the stub behind the shuffle_xor shorthand:

Rust:         warp::shuffle_xor_sync(u32::MAX, val, lane)
MIR:          _5 = shuffle_xor_sync(const u32::MAX, _3, _4) -> bb2
Pliron MIR:   %v = nvvm.shfl_sync_bfly_i32 %mask, %val, %lane : u32
Pliron LLVM:  %v = call i32 @llvm_nvvm_shfl_sync_bfly_i32(i32 -1, %val, %lane, i32 31)
LLVM IR:      %v8 = call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 -1, %v3, %v4, i32 31) #0
PTX:          shfl.sync.bfly.b32 %r3, %r1, %r2, 31;

Six representations. One Rust function call becomes one PTX instruction. The intermediate steps exist so that each transformation is small, verifiable, and independently testable.

Quick-Reference Checklist#

For a hand-written intrinsic (one intrinsics/catalog.json cannot describe), every file you need to touch, in order:

cuda-device/src/<module>.rs – pub fn with #[inline(never)] and unreachable!() body. This is the user-facing API.
dialect-nvvm/src/ops/<module>.rs – #[pliron_op(name = "nvvm.<name>", ...)] struct, Verify impl (check operand/result counts), register() call.
mir-importer/src/translator/terminator/mod.rs – match arm in try_dispatch_intrinsic(), modelled on the hand-written typed_swap arm. Zero-operand intrinsics are catalog territory: the generated dispatcher emits them through helpers::emit_generated_nvvm_intrinsic(), which also stamps the ABI marker a hand-written arm would have to stamp itself.
mir-lower/src/convert/interface_impls.rs – MirToLlvmConversion impl for the new op, dispatching to the converter function.
mir-lower/src/convert/intrinsics/<module>.rs – conversion logic. Use call_intrinsic() for LLVM intrinsics, or inline_asm_convergent() for inline PTX.
llvm-export/src/export/ – only if convergent: add to is_convergent_intrinsic().

For a catalog intrinsic, steps 2 through 4 are not yours to write: add the entry to intrinsics/catalog.json and cuda-intrinsics-gen regenerates the op, the dispatch arm, and the conversion impl.

Side-by-Side Comparison#

Stage	`threadIdx_x` (simple)	`shuffle_xor_sync` (complex)
cuda-device	`fn() -> u32`	`fn(u32, u32, u32) -> u32`
dialect-nvvm	0 operands, 1 result	3 operands, 1 result
mir-importer	Generated arm, generic helper	Generated arm, operand translation
mir-lower	`call @intrinsic()`	`call @intrinsic(4 args)`
Convergent?	No	Yes (`#0`)
PTX	`mov.u32 %r1, %tid.x`	`shfl.sync.bfly.b32 ...`

That is the entire process. Five files, each with a clear and narrow responsibility. The pattern is mechanical enough that adding a new intrinsic should take about thirty minutes once you have done it once – most of that time spent reading the PTX ISA spec to figure out exactly what instruction you want.

Adding New Intrinsics#

The Five-Stage Pipeline#

Example 1: threadIdx_x (the Simple Case)#

Stage 1 – Declare in cuda-device#

Stage 2 – Define the dialect-nvvm Op#

Stage 3 – Recognize in mir-importer#

Stage 4 – Lower to the LLVM dialect#

Stage 5 – Export (Nothing to Change)#

Example 2: shuffle_xor (the Complex Case)#

Stage 1 – Declare in cuda-device#

Stage 2 – Define the dialect-nvvm Op#

Stage 3 – Recognize in mir-importer#

Stage 4 – Lower to the LLVM dialect#

Stage 5 – Export (Add Convergent)#

The Inline PTX Path#

End-to-End: The Full Journey#

Quick-Reference Checklist#

Side-by-Side Comparison#

Example 1: `threadIdx_x` (the Simple Case)#

Stage 1 – Declare in `cuda-device`#

Stage 2 – Define the `dialect-nvvm` Op#

Stage 3 – Recognize in `mir-importer`#

Example 2: `shuffle_xor` (the Complex Case)#

Stage 1 – Declare in `cuda-device`#

Stage 2 – Define the `dialect-nvvm` Op#

Stage 3 – Recognize in `mir-importer`#