Minor NVIDIA Passes

This page indexes NVIDIA-proprietary passes that are too small or insufficiently decompiled for dedicated pages. For the ten passes that were previously documented here and now have full pages, see the links below.

Passes with Dedicated Pages

alloca-hoisting -- Entry Block Alloca Consolidation

PTX requires all stack allocations to reside in the entry block. This pass moves alloca instructions inserted by inlining or loop transforms into the entry block, preserving order and alignment. Without it, non-entry-block allocas produce invalid PTX.

image-optimizer -- Texture/Surface Access Optimization

Groups related texture loads for cache utilization and merges redundant surface operations. Works in coordination with Replace Image Handles (below). See also Machine-Level Passes.

nvptx-peephole -- Machine-Level Peephole

PTX-specific peephole that folds redundant cvta address space conversions, optimizes predicate patterns, and simplifies PTX-specific instruction sequences. Distinct from the IR-level NVVM Peephole. See Machine-Level Passes for pipeline position.

proxy-reg-erasure -- Redundant cvta.to.local Removal

Removes redundant cvta.to.local instructions left by address space lowering. Runs late in the pipeline after register allocation. See Machine-Level Passes.

valid-global-names -- PTX Identifier Sanitization

Rewrites global symbol names to comply with PTX naming rules, removing characters illegal in PTX identifiers (@, $, etc.). Runs immediately before PTX emission.

replace-image-handles -- Texture/Surface Handle Substitution

Replaces IR-level texture/surface handle references with PTX-level .tex / .surf declarations. Paired with image-optimizer above. See Machine-Level Passes.

extra-mi-printer -- Register Pressure Diagnostics

Prints per-function register pressure statistics. Used for tuning pressure heuristics during development. Not active in release builds.

nvvm-intr-range -- Intrinsic Range Metadata

Attaches !range metadata to NVVM intrinsics that return hardware-bounded values (threadIdx.x, blockIdx.x, etc.), enabling downstream known-bits analysis and range-based dead code elimination. Tightens ranges when __launch_bounds__ metadata is present. Documented in detail in KnownBits & DemandedBits.

GenericToNVVM -- Global Address Space Migration

Moves global variables from generic address space (AS 0) to global address space (AS 1), inserting addrspacecast at use sites. Required because PTX globals must reside in .global memory. Documented in detail in PTX Emission.

Other Passes Documented Elsewhere

These passes appear in the NVPTX backend but have primary documentation on other pages: