NVVM IR Verifier (Deep Dive)

The NVVM IR Verifier (nvvm-verify) is NVIDIA's three-layer correctness gate that runs between optimization passes throughout the CICC pipeline. Unlike LLVM's generic Verifier pass, which validates structural IR invariants, this pass enforces the complete NVVM IR contract: valid target triples, legal address space usage, architecture-gated intrinsic availability, MMA dimension/type constraints, function attribute restrictions, and atomic operation rules. It is the single largest verification subsystem in CICC at approximately 230KB across three cooperating functions. The verifier is inserted at roughly a dozen points in every optimization tier, guarded only by NVVMPassOptions[600] (disable). Every NVVM intrinsic call, every address space cast, and every unsupported CPU-oriented feature triggers a check here; failure produces a diagnostic message and sets the module error flag, but compilation continues to collect as many errors as possible in a single run.

Key Facts

Three-Layer Verification Architecture

The pass operates as three nested verification functions. The module verifier is the entry point; it calls the function verifier once per function, and the function verifier dispatches to the intrinsic verifier for every intrinsic call instruction.

sub_2C80C90 (NVVM module verifier)
  |
  +-- Validate data layout string
  +-- Validate target triple against whitelist
  +-- sub_2C797D0() for each global variable
  +-- sub_2C7A130() for each function declaration
  +-- sub_2C7AA20() for each named metadata node
  |
  +-- For each function:
  |     |
  |     +-- sub_2C771D0 (NVVM function verifier)
  |     |     +-- Cluster dimension validation (Hopper+ gate)
  |     |     +-- Parameter width validation (>=32-bit or sext/zext)
  |     |     +-- Function attribute rejection (17 attributes)
  |     |     +-- Entry/exit handler constraints
  |     |
  |     +-- For each instruction in each basic block:
  |           |
  |           +-- Switch on opcode 0x1E..0x60
  |           +-- Opcode 0x55 (intrinsic call) --> sub_2C7B6A0
  |                 (NVVM intrinsic verifier, 143KB)
  |                 +-- Switch on intrinsic ID
  |                 +-- SM version gate checks
  |                 +-- Type, address space, constant arg validation
  |                 +-- MMA shape/type cross-validation

Context Object Layout

All three verifiers share a context object passed as the first argument:

Target Triple Whitelist

The module verifier validates the module's target triple against two whitelists depending on mode.

UnifiedNVVMIR Mode (mode == 1) -- Exact Match

Eight triples are accepted:

The nvsass triples confirm that CICC can compile directly to native GPU assembly (SASS) without the PTX intermediate step, and can do so for DirectX shader and SPIR-V/Vulkan shader pipelines. This reveals CICC's role in NVIDIA's shader compiler toolchain beyond CUDA.

Failure message: "Invalid target triple".

Standard Mode (mode != 1) -- Prefix + Suffix Match

The triple must begin with "nvptx-" or "nvptx64-" and end with "-cuda". The middle component is wildcarded.

Failure message: "Invalid target triple (<actual>), must be one of:" followed by "nvptx-*-cuda" and "nvptx64-*-cuda".

Data Layout Validation

If the module's data layout string is empty: "Empty target data layout, must exist".

Otherwise, sub_2C74F70 parses and validates the layout. On failure, the verifier prints "Example valid data layout:" with reference strings from:

Per-Instruction Validation (Module Verifier)

After calling sub_2C771D0 for function-level checks, the module verifier iterates every instruction in every basic block and dispatches on the LLVM IR opcode. The opcode range is 0x1E through 0x60:

The unsupported instructions -- indirectbr, invoke, resume, landingpad -- are CPU exception-handling features with no GPU equivalent. Their rejection at the IR level prevents downstream passes from encountering them.

Address Space Casting Rules

The addrspacecast validation enforces NVIDIA's GPU address space model:

Rule: At least one operand of addrspacecast must be AS 0 (generic).
      Non-generic-to-non-generic casts are illegal.

Legal:   addrspacecast i32* addrspace(0) to i32* addrspace(1)  ; generic -> global
Legal:   addrspacecast i32* addrspace(3) to i32* addrspace(0)  ; shared -> generic
Illegal: addrspacecast i32* addrspace(3) to i32* addrspace(1)  ; shared -> global

The valid address space range check uses the expression (AS + ~2) & 0xFFFFFF) > 2, which means AS values 0 (generic), 1 (global), and 3 (shared) are always valid for atomic and cast operations. AS 2 (constant) and higher values have restricted usage contexts.

Function Attribute Rejection

The function verifier (sub_2C771D0) rejects 17 LLVM function attributes that have no GPU meaning. Each is identified by its LLVM attribute kind ID:

These attributes fall into four categories: (1) CPU ABI (naked, alignstack, noredzone), (2) security hardening (ssp/sspreq/sspstrong, safestack, sanitizers), (3) EH-related (uwtable, returns_twice, personality), and (4) linker features (jumptable, nonlazybind, builtin, nobuiltin). None have GPU equivalents.

Additional Function-Level Checks

Architecture Gates (SM-Gated Features)

The intrinsic verifier (sub_2C7B6A0) uses the SM version stored at context offset +8 (encoded as SM*10) to gate feature availability. The threshold checks use <=, so e.g. <= 899 means "below sm_90".

Note the typo "instrinsic" in the Volta and Ada messages -- this is present in the binary. The Blackwell gate threshold of 1199 means the .offset.bindless intrinsics are available on sm_120 (value 1200) and above, covering all Blackwell-generation architectures including consumer (sm_120/121) and datacenter (sm_100/103).

Intrinsic Verification Categories

The intrinsic verifier is a single monolithic switch on the NVVM internal intrinsic ID (stored at function value offset +36). The 143KB function covers 26+ validation categories. The sections below group the major ones; reduction (nvvm.red), fence, setmaxnreg, tcgen05.cp, and arch-gated MMA variants each have dedicated check blocks documented in subsections G.1-G.3, O.1, P, and Q.

A. Constant Argument Validation

Many NVVM intrinsics require one or more arguments to be compile-time constants (typically mode selectors, masks, or task IDs):

• "arg0 of intrinsic not constant"
• "op0 of intrinsic not constant" / "op1 of intrinsic not constant"
• "Flag argument must be an immediate."
• "the task_id parameter must be constant"
• "the mask parameter must be constant"
• "Mode operand must be constant"

B. Rounding Mode Validation

Rounding mode encoding: bits[2:0] of the mode word
Valid range: 1..4 (round-to-nearest-even, round-down, round-up, round-to-zero)
Reject: value == 0 or value > 4
Message: "rounding mode not a valid value"

C. Subword Mode Validation

For conversion intrinsics that operate on sub-word portions:

Source subword mode:  bits[9:7], valid range 0..2
Dest subword mode:    bits[12:10], valid range 0..2
Messages: "src subword mode not a valid value"
          "dest subword mode not a valid value"

D. Reserved Bits Checking

Multiple locations verify that high/reserved bits in mode words are zero:

• "reserved flag bits used"

This prevents future-proofing conflicts if NVIDIA later assigns meaning to currently reserved fields.

E. Address Space Validation

Intrinsics that access memory enforce specific address space requirements:

F. Type Validation

Inline assembly type validation uses a bitmask check: valid bit widths are 1, 8, 16, 32, 64 (encoded as 0x1000000010001 for fast lookup).

G. Atomic Intrinsic Validation

G.1. NVVM Reduction (nvvm.red) Validation

The nvvm.red family (warp/CTA-wide reduction operators) has its own dedicated validation block, separate from generic atomics. The pointer operand identifies the destination memory; the operator selects the reduction kernel; the type-and-vector-length pair must match the hardware-supported reduction operator set.

G.2. Fence Intrinsic Validation

In addition to the IR-level fence instruction (opcode 0x40, see Per-Instruction Validation), there is a dedicated llvm.nvvm.fence intrinsic whose opcode/scope/ordering tuple is validated independently.

G.3. Register Count (nvvm.setmaxnreg) Validation

The Hopper register-reallocation intrinsic accepts only specific count values:

Combined, the legal reg_count set is {24, 32, 40, ..., 248, 256} -- 30 distinct values.

H. Texture/Surface Validation

H.1. Barrier Pointer Validation

Mbarrier / arrive-wait synchronization objects live exclusively in addrspace(3) (shared); a generic or global pointer is rejected.

I. SATF (Saturate-to-Float) Validation

For math intrinsics with saturation control (IDs 0x2281-0x229C, covering fma/mul/add variants):

Message: "satf operand must be a constant zero"

The satf parameter was deprecated but the intrinsic signatures retain it for ABI compatibility. The verifier enforces it must be zero.

J. Constant Load Validation

For ID 0x2310 (constant bank load):

K. TMA/Shared Memory Validation

For IDs 0x2319-0x231B:

L. Load Bounds Check

For ID 0x231C:

Validation: (value & 7) must be <= 2
Message: "invalid load bounds check type"
Also: "pointer address space not global"

M. Convergent Branch Result Validation

For IDs 8282 (llvm.nvvm.branch.if.all.convergent) and 8283 (llvm.nvvm.branch.if.convergent):

Message: "result of llvm.nvvm.branch.if.convergent and
          llvm.nvvm.branch.if.all.convergent can only be
          used by exactly one branch instruction"

This enforces that the convergent branch intrinsic's boolean result flows directly to a single terminator branch, preventing misuse that would break convergence guarantees.

N. MMA (Matrix Multiply-Accumulate) Validation

The most complex validation category (ID 0x2366 = 9062). Validates WMMA/MMA intrinsics against a multidimensional constraint space:

Opcode byte encoding:

MNK dimension decoding (selected cases):

Validation checks:

O. Type Conversion Validation

For IDs 0x2106 and 0x2107:

Conversion type: bits[3:1], must be 1..4
Messages: "conversion type not a valid value"
          "Invalid dst type" / "Invalid src type"
          "Src and dst type must be different types"
          "Src and dst type must be different bit widths"

O.1. Pack-Float Conversion (nvvm.cvt.packfloat) Validation

The pack-float family converts pairs of FP values to packed sub-word formats; the type pair must be one of a small whitelist (e.g., {f32, f32} -> v2bf16).

P. tcgen05 (Blackwell Tensor Core Gen-5) Validation

The Blackwell-introduced tcgen05.cp / tcgen05.mma intrinsics carry an encoded shape and multicast flag word; the verifier validates the tuple as a unit:

These checks fire only on sm_100+ targets; on earlier architectures the intrinsic is rejected by the SM gate first.

Q. Architecture-Gated MMA / Vector Intrinsic Variants

Beyond the eight monolithic SM gates in the Architecture Gates table, each MMA, vector-atomic, and TMA-2CTA variant has a private threshold check colocated with its lowering helper. The helpers live outside sub_2C7B6A0 proper but are reached from the same dispatch -- the verifier walks an arch-availability sub-table per feature rather than collapsing them into the main switch.

The encoding is uniform across all variants: *(_DWORD *)(ctx + 1136) resolves to the target descriptor, and field +344 carries the compute capability as major * 10 + minor (e.g. sm_70 = 70, sm_72 = 72, sm_75 = 75, sm_90 = 90, sm_100 = 100). The threshold uses <= so <= 0x45u rejects sm_69 and below, admitting sm_70+. Two helpers (2CTA-TMA, block-scale tcgen05) use a second field at +340 to encode an extended family token (values like 1101, 1102 select Blackwell-datacenter-family) and a third at +336 for the family-base SM.

Per-variant SM availability matrix (decoded from sub_36E5710, sub_36E72A0, sub_36E7580, sub_36E77C0, sub_36E7B50, sub_36E7EA0, sub_36E8630, sub_36E8BD0, sub_36E91F0, sub_36EC510, sub_304C610, sub_2179030):

SM tiers pinned by Section Q:

The IMMA family has a two-stage check: the base test v <= 0x47 admits sm_72, but several IMMA variants then add && opc > 1 or && a2 in {386,387,1595,1596} to forbid sm_72 specifically -- those subvariants need sm_75. The dispatch matches each sm_72/sm_75 split point to a different LLVM intrinsic ID, so a sm_72 user reaching a sm_75-only intrinsic ID gets the same "immaXXX is not supported" message rather than a distinct one. The BMMA threshold of 0x48 (72) rejects sm_72 cleanly because sm_73 and sm_74 do not exist in the NVIDIA ISA lineup -- the next valid value after 72 is 75, which is where BMMA debuted on Turing.

QUIRK: The 2CTA-TMA gate in sub_36EC510 is not a simple <= threshold but a three-clause family-token test on field +340 (extended SM token, e.g. 1101 = sm_110a, 1102 = sm_111a) joined to field +336 (family-base SM). The expression (__ROR4__(-858993459 * v38 + 1717986918, 1) > 0x19999999u || v39 <= 0x57) && v39 <= 0x55 is a compiler-folded modulo-5 test (-858993459 = -0x33333333 is the magic multiplier for n % 5), checking that the family token is not a multiple of 5 while the base SM is in the [85, 87] window -- this is the verifier's way of admitting sm_100a / sm_103a / sm_110a while rejecting sm_100 / sm_103 / sm_110 (non-a variants).

QUIRK: BMMA threshold 0x48u (72) versus the user-facing "BMMA requires sm_75". The verifier accepts any value >= 73, but no such SM exists in the NVIDIA range, so the effective minimum is sm_75. This single-byte threshold avoids encoding the sm_73/74 holes explicitly -- the gap is implicit in the SM numbering.

QUIRK: The IMMA ld-C gate sub_36E7EA0 uses a four-value whitelist {386, 387, 1595, 1596} on the intrinsic-ID argument to forbid sm_72 selectively. These four IDs correspond to s8/u8 packed-load intrinsics that landed in sm_75; the sm_72 IMMA path supports only the unpacked variants. A more conventional implementation would split into two separate intrinsic IDs with separate gates, but the verifier collapses both into one ID-range check, which is why a single helper raises "immaldc is not supported on this architecture" for two structurally different errors (pre-Volta+ vs. sm_72-extended).

Several "unexpected"-prefixed messages indicate the verifier detected an MMA variant that should have been lowered earlier in the pipeline; reaching the verifier in this state implies an upstream pass bug rather than user IR error: "unexpected imma_ld intrinsic!", "unexpected imma_mma intrinsic call!", "unexpected overloaded mma intrinsic call!", "unexpected overloaded mma load intrinsic call!", "unexpected overloaded mma store intrinsic call!", "unexpected WMMA intrinsic!".

R. Other Validation Categories

Cmpxchg Restrictions

The module verifier enforces strict constraints on cmpxchg:

Allowed types:  i32, i64, i128
Allowed spaces: generic (AS 0), global (AS 1), shared (AS 3)

Messages:
  "Atomic operations on non-i32/i64/i128 types are not supported"
  "cmpxchg pointer operand must point to generic, global, or shared address space"

This rules out i8/i16 atomics (hardware does not support sub-word CAS) and atomics on constant/local address spaces.

Tensor Memory Restrictions

Load and store instructions targeting address space 6 (tensor memory) are rejected at the IR level:

Message: "Tensor Memory loads/stores are not supported"

Tensor memory access is handled through dedicated intrinsics (TMA/cp.async) rather than generic load/store instructions. The verifier enforces this indirection.

Pipeline Placement

The NVVMVerifier is inserted repeatedly throughout the optimization pipeline, not just once. In the pipeline assembler (sub_12E54A0), it appears after nearly every major optimization pass, gated by !NVVMPassOptions[600]:

This aggressive re-verification catches bugs introduced by any optimization pass. In debug/development builds, this is the primary mechanism for detecting optimizer-introduced IR invalidity.

Configuration

Diagnostic Infrastructure

Error messages are produced through a chain of helper functions:

The error model is accumulate-and-continue: the verifier sets the error flag at context offset +16 and writes the diagnostic, but does not abort. This allows a single verification run to report all errors in the module.

Function Map

Role labels below are descriptive; only sub_<HEX> addresses are binary-confirmed.

Cross-References

• GPU Target Architecture -- SM table and architecture gating
• Hopper (sm_90) -- TMA, cluster operations, WGMMA
• Blackwell (sm_100) -- tcgen05, .offset.bindless
• Memory Space Optimization -- address space enforcement and resolution
• NVIDIA Custom Passes index -- pass inventory
• IP Memory Space Propagation -- inter-procedural address space analysis