Getting started

Install

pkg> add PTX, CUDA

PTX.jl depends on LLVM.jl for the inline-asm machinery (@asmcall). Kernel launch and memory go through CUDA.jl. PTX.jl itself carries no hard GPU dependency.

First kernel

Every PTX instruction is one @ptx_str call. Special registers (%tid.x, %ctaid.y, …) are spelled the same way, ptx"%tid.x". Inputs flow in as Julia arguments, the result is the call's return value:

using PTX, CUDA

function add_kernel!(c, a, b)
    tid = ptx"mov.u32"(ptx"%tid.x") # access instructions and special registers,
    i = tid + one(tid)               # while also writing familiar julia code.
    c[i] = ptx"add.f32"(a[i], b[i])  # same as a[i] + b[i], assuming Float32
    return
end

n = 128;
a = CUDA.randn(n);
b = CUDA.randn(n);
c = similar(a);
@cuda threads=n add_kernel!(c, a, b);
@assert c == a + b

# or simply use an instruction in a broadcast:
@assert c == ptx"add.f32".(a, b)

ptx"add.f32" constructs an Operation singleton; the call site is @generated, picks constraint letters from the argument types, infers the return type from the trailing .f32 modifier, and lowers to a single @asmcall("add.f32 $0, $1, $2;", "=f,f,f", false, Float32, Tuple{Float32, Float32}, a, b).

Composing with CUDA.jl

PTX.jl is strictly additive. CUDA.jl owns everything around the instruction; PTX.jl owns the instruction itself.

Rule of thumb: write what you'd write in normal CUDA C++ in Julia + CUDA.jl; drop into PTX.jl for what you'd write as a <header> library call that nvcc lowers to asm volatile-equivalent inline PTX. Real kernels mix both.

Device code only

PTX instructions run only in device code. Called from the host, they throw a PTX.DeviceOnlyError rather than reaching the host's LLVM backend, which would abort the Julia process on PTX constraint letters and NVVM intrinsics. For the same reason, host reflection (code_typed, code_llvm) on a function using PTX.jl shows only that error path. Inspect device code with CUDATools.code_llvm and @device_code_llvm, or PTX.device_code_typed.

The device/host split is the PTX.on_device overlay, which PTX.jl's CUDACore extension installs for CUDA.jl's compiler. A GPUCompiler client other than CUDA.jl that targets NVPTX with its own method table must add the same overlay there.

Address spaces

ptx"ld.global.f32" / ptx"cp.async.ca.shared.global" and friends take typed pointers. PTX.jl provides a small AS module (not exported — spell it PTX.AS) of LLVM address-space numbers that match CUDACore.AS.*, plus two NVPTX-only spaces CUDACore doesn't name:

PTX.AS.Generic        # 0
PTX.AS.Global         # 1
PTX.AS.Shared         # 3
PTX.AS.Const          # 4
PTX.AS.Local          # 5
PTX.AS.Param          # 101
PTX.AS.Tmem           # 6 — Blackwell tensor memory (tcgen05 taddr operands)
PTX.AS.SharedCluster  # 7 — distributed shared memory; `mapa.shared::cluster` returns these

A pointer reaches a wrapper as Core.LLVMPtr{T, AS.X}. NVPTX lowers non-Generic address-space pointers as i64 at the LLVM IR level even when the underlying PTX address is 32-bit (shared, param, local). The chain emits l (i64) for any LLVMPtr; hand-written shared-memory wrappers (cp.async, ldmatrix, stmatrix, mbarrier) override to r (i32) where ptxas requires the 32-bit form.

PTX.TMADescriptorPtr uses AS.Const as the TMA wrapper carrier convention. Descriptors uploaded by PTX.upload_tma_descriptor live in device global memory; their pointer conversion happens on the host.

Where to next

• Chain DSL — the full set of conventions the chain encodes.
• Wrappers — when to reach for a hand-written wrapper.
• Transpiler — turn existing .ptx files into Julia.