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 (NVVMModuleVerifier)
  |
  +-- 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 (NVVMFunctionVerifier)
  |     |     +-- 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
  |                 (NVVMIntrinsicVerifier, 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:

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

H. Texture/Surface Validation

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"

P. 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

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