Keyboard shortcuts

Press or to navigate between chapters

Press S or / to search in the book

Press ? to show this help

Press Esc to hide this help

Function Map

Binary: ptxas v13.0.88, 37.7 MB stripped ELF, ~40,000 functions
Documented: 2,063 unique functions across 70 wiki pages
This page: Top ~100 most cross-referenced functions, plus routing tables
Complete listings: Each wiki page has its own Function Map section with full details

This page is the central lookup index for identified functions in ptxas. It lists the functions that appear most frequently across the wiki (cross-cutting infrastructure and major entry points), and provides routing tables to find any function by address range or subsystem.

Confidence levels: CERTAIN = named in symbols or strings. HIGH = strong evidence from strings and call patterns (>90%). MEDIUM = structural analysis with partial string evidence (70-90%).


Core Infrastructure

These functions appear in 10+ wiki pages -- they are the universal building blocks called by nearly every subsystem.

AddressIdentityPagesCallersNotes
0x424070pool_alloc(pool, size)193,809Custom slab allocator, 8-byte aligned
0x4248B0pool_free(ptr)81,215Coalescing free, boundary tags
0x4280C0get_thread_local_context103,928Most-called function in ptxas; 280-byte TLS struct
0x42BDB0fatal_OOM_handler83,825Called on every allocation failure
0x426150hashmap_put(map, key, value)112,800Open-addressing + chaining, auto-resize
0x426D60hashmap_get(map, key)11422Returns value or 0
0x425CA0hashmap_create(hash_fn, cmp_fn, cap)7127Integer/pointer/custom hash modes
0x427630murmurhash3_x86_32(str)573Constants: 0xcc9e2d51, 0x1b873593
0x42D850hashset_insert(set, key)4282Hash set variant
0x42FBA0diagnostic_emit(desc, loc, fmt...)72,350Central error/warning reporter
0x42F590fatal_internal_error(desc, ...)83,825Assertion handler
0x4279D0starts_with(str, prefix)4185Returns suffix pointer or 0
0x42CA60list_push_front(node, head_ptr)4298Pool-allocated linked list
0xBDBA60bitvector_allocate8many(bits+31)>>5 word count
0xBDCDE0bitvector_or_assign (SSE2)5many_mm_or_si128 on 128-bit chunks

Details: Memory Pools, Hash & Bitvector, Threading


Compilation Driver & CLI

AddressIdentityPagesCallersNotes
0x409460main51Delegates to 0x446240
0x446240real_main (top-level driver)131Orchestrates entire pipeline
0x4428E0ptx_input_setup61Version/target validation
0x43CC70per_entry_compile_unit51Processes each entry through pipeline
0x43F400function_abi_config41Parameter regs, return addr, scratch
0x43A400compilation_target_config71SM-specific defaults
0x43B660register_constraint_calculator51Balances .maxnreg, occupancy
0x432A00option_registration91CLI option definitions
0x434320option_parser91Validates combinations, applies state

Details: Pipeline Entry, Pipeline Overview, CLI Options


PTX Front End

AddressIdentityPagesCallersNotes
0x46E000instruction_table_builder9193 KB, 1168 callees, one per PTX opcode
0x451730parser_setup (special register init)91%ntid, %laneid, %clock, etc.
0x4CE6B0bison_parser (directive/decl)71.local_maxnreg, .alias, .pragma
0x720F00flex_lexer (ptxlex / yylex)82~550 Flex rules, DFA scanner
0x4B2F20ptx_validator_general41Validates texture, surface, cvt, call
0x4C5FB0ptx_validator_mma_wmma_tcgen0541MMA, WMMA, tensor core validation
0x71F630preprocessor_dispatch41.MACRO, .ELSE, .INCLUDE
0x489050ptx_to_ori_converter51PTX AST to ORI IR translation

Details: PTX Parser, PTX Directives, PTX to ORI


Static Initialization

AddressIdentityPagesCallersNotes
0x4094C0ctor_001 -- thread infra init40pthread_key_create, mutex
0x4095D0ctor_003 -- PTX opcode name table60~900 ROT13-encoded PTX mnemonics
0x40D860ctor_005 -- tuning knob registry6080 KB, 2000+ ROT13 knob names
0x421290ctor_007 -- scheduler knob registry4098 ROT13 scheduler knobs

Details: Pipeline Entry, Binary Layout


Phase Manager & Optimization Framework

AddressIdentityPagesCallersNotes
0xC60D30phase_factory (159-case switch)121Allocates phase objects
0xC62720PhaseManager_ctor102159-entry phase table
0xC64F70phase_dispatch_loop52Executes phases, reports timing
0xC64310per_phase_timing_reporter51"[Total N KB] [Freeable N KB]"
0xC641D0phase_name_to_index_lookup53Binary search, case-insensitive
0x7DDB50phase_run_dispatch14manyVtable-based phase execution
0x9F4040NamedPhases_parse_and_build61"shuffle", "OriCopyProp", etc.
0x798B60NamedPhases_parser42PTXAS_DISABLE env var parsing
0x799250IsPassDisabled54Checks knob index 185
0xA36360pass_sequence_builder61Constructs NvOptRecipe pass list

Details: Phase Manager, Pass Inventory, Optimizer Pipeline


ORI IR & Instruction Access

AddressIdentityPagesCallersNotes
0x9253C0instruction_operand_get11manyOperand accessor on ORI instructions
0x7E6090instruction_modifier_set10manyIR modification helper
0x781F80instruction_iterator12manyDoubly-linked list traversal
0x7DF3A0instruction_property_query5manyInstruction flag/attribute checker
0x91BF30register_type_query8manyRegister class/type inspection
0x9314F0register_class_id_query71,547Most-called non-trivial regalloc fn
0x931920register_class_compat_checker6328Pair register class handling
0x934630register_id_packer9856Packs reg#/class/type into 32-bit
0xB28E00ir_node_type_query5manyNode kind discrimination
0xB28E90ir_node_field_accessor6manyGeneric field getter
0xA50650CodeObject_EmitRecords1874 KB, ORI record serializer (56 section types)
0xA53840EmitRecords_wrapper11Thin wrapper, adds type-44 header

Details: Instructions, Registers, Data Structures, CFG


Intrinsic Infrastructure

AddressIdentityPagesCallersNotes
0x5D1660intrinsic_table_register (608 entries)71Master name-to-ID table
0x5D4190intrinsic_dispatch_builder131PTX opcode -> codegen handler mapping
0x5FF700intrinsic_prototype_emitter51354 KB -- largest function in binary
0x5C7A50wmma_mma_codegen41173 KB, all shapes/types/layouts
0x5C10A0mma_codegen (mma.sync)41120 KB, m8n8k4 through m16n8k256
0x5BBC30tcgen05_mma_codegen (Blackwell)5190 KB, 5th-gen tensor core
0x70FA00ocg_intrinsic_handler81OCG-level intrinsic routing
0x6A97B0intrinsic_lowering_main4126 KB, switch-based lowering
0x6C9EB0ocg_builtin_name_lookup51Blackwell+ OCG name table

Details: Intrinsics Index, Math Intrinsics, Tensor Intrinsics, Sync & Warp


Register Allocator

AddressIdentityPagesCallersNotes
0x9721C0regalloc_entry ("REGALLOC GUIDANCE")61Top-level allocator entry
0x957160fatpoint_allocator_core71Core fatpoint graph coloring
0x96D940spill_guidance_engine51Determines spill strategy
0x971A90full_alloc_with_spill_retry41"NOSPILL REGALLOC" path
0x9714E0regalloc_failure_reporter61"Register allocation failed..."
0x926A30interference_graph_builder9722 KB, SSE bitvectors
0x92C240liveness_bitvector_ops587Set/clear/query with aliasing
0x917A60opcode_to_regclass_mapping4221Massive switch
0x910840ConvertMemoryToRegisterOrUniform51Pass driver

Details: RegAlloc Overview, RegAlloc Algorithm, Spilling, ABI


Instruction Scheduling

AddressIdentityPagesCallersNotes
0x8D0640ScheduleInstructions (top-level)71String: "ScheduleInstructions"
0x688DD0scheduler_engine (main BB loop)51ReduceReg / DynBatch selection
0x8C9320scheduling_priority_function40~300 locals, core heuristic
0x68B9C0dependency_graph_builder41RAW/WAR/WAW hazard analysis
0x6820B0build_ready_list51Zero-dependency instructions
0x8CD6E0reverse_scheduling_driver41Reverse post-order iteration
0x8CEE80register_budget_with_occupancy41Pressure coeff default 0.045
0x8E4400hw_profile_table_init63Encoding/latency property tables
0xA9CDE0scheduling_metadata_builder61Per-instruction sched metadata
0xA9CF90scheduling_metadata_accessor5manySched metadata field queries
0xAED3C0scheduling_optimization_mega_pass40137 KB, ~560 locals, largest vtable pass

Details: Scheduling Overview, Scheduling Algorithm, Latency Model, Scoreboards


Codegen & ISel

AddressIdentityPagesCallersNotes
0x169B190isel_pattern_dispatch (master)51280 KB, 65,999 insns -- largest function
0x143C440sm120_peephole_dispatch41SM120 (RTX 50), 373-case switch
0x198BCD0sm100_peephole_dispatch41SM100 (Blackwell), 1336 callees
0x83EF00main_peephole_pass6029 KB, 392 callees
0x6D9690master_instruction_encoder7194 KB, opcode switch
0x6E4110sass_codegen_main41EmitSASSForFunction, FNV-1a BB hash
0x6F52F0SASS_pipeline_run_stages51Mercury SASS compilation pipeline
0x9ED2D0MercConverter_entry61ORI to Mercury IR conversion
0x9F1A90MercConverter_builder61Mercury instruction construction

Details: ISel, Encoding, Peephole, Mercury, Templates


Bitfield Encoding

AddressIdentityPagesCallersNotes
0x7B9B80bitfield_insert(insn, off, wid, val)918,347Most-called by caller count
0x7BC030encode_register_operand46,1471-bit + 4-bit type + 10-bit reg
0x7B9D60encode_reuse_flags_predicate42,4081-bit reuse + 5-bit predicate
0x7BC5C0encode_immediate_const_operand41,449Const buffer index or immediate
0x7BCF00encode_predicate_register41,657PT=14, 2-bit type + 3-bit condition
0x10B61801_bit_boolean_encoder38,091.S/.U, .STRONG, etc.

Details: Encoding, SASS Printing


ELF / CUBIN Output

AddressIdentityPagesCallersNotes
0x612DE0section_attr_builder11176 KB, ELF section/attribute config
0x1C9F280master_elf_emitter91Complete CUBIN assembly
0x1CB53A0elf_world_init71672-byte ELFW context
0x1CB68D0symbol_table_builder51.symtab from internal symbols
0x1CABD60master_section_allocator51Shared/const/local memory
0x1CB3570add_function_section544Creates .text.FUNCNAME + .rela
0x1CD48C0relocation_processor51Relocation section emission
0x1C9B110mercury_capsule_builder41Creates embedded .nv.merc ELF

Details: ELF Emitter, Sections, Relocations, Debug Info, Capsule Mercury


Knobs System

AddressIdentityPagesCallersNotes
0x79B240GetKnobIndex62ROT13 name lookup, case-insensitive
0x79D070ReadKnobsFile51Parses [knobs] section from file
0x79F540ParseKnobValue4112-type switch: bool/int/float/string/...
0x79D990ProcessKnobs (top-level)41File + pragma + numbered config
0xA0F020knob_conditional_evaluator5many[WHEN condition] handler

Details: Knobs, Opt Levels


Target-Specific Code

AddressIdentityPagesCallersNotes
0x6765E0target_profile_selector71SM-dependent profile dispatch
0x607DB0target_feature_query7manySM feature capability checks
0x896D50sass_mnemonic_table_init (ROT13)41~400+ SASS instruction names
0x89FBA0instruction_latency_init43Encoding/latency property tables

Details: Targets Index, Turing-Ampere, Ada-Hopper, Blackwell, tcgen05


Subsystem Routing Table

To find a specific function, locate it by address range or subsystem topic in this table. Each page contains a detailed Function Map section with complete listings.

By Subsystem Topic

SubsystemPrimary PagesFunctions
Memory allocator, poolsmemory-pools.md30
Hash maps, bitvectors, setshash-bitvector.md51
Threading, TLS, jobserverthreading.md41
CLI parsing, option handlingcli-options.md17
Tuning knobs (2000+ knobs)knobs.md56
Optimization levelsopt-levels.md14
DumpIR debug outputdumpir.md14
Compilation pipelineoverview.md, entry.md56+25
PTX lexer & parserptx-parser.md75
PTX directivesptx-directives.md41
PTX-to-ORI translationptx-to-ori.md41
Optimizer pipelineoptimizer.md28
ORI instruction IRinstructions.md80
CFG constructioncfg.md18
Register representationregisters.md40
IR data structuresdata-structures.md74
Phase manager (159 phases)phase-manager.md26
Copy propagation, CSE, GVNcopy-prop-cse.md65
General optimization passesgeneral-optimize.md71
Loop optimization (unroll, LICM, SWP)loop-passes.md92
Branch/switch optimizationbranch-switch.md24
Strength reductionstrength-reduction.md25
Predicationpredication.md28
Rematerializationrematerialization.md55
Liveness analysisliveness.md42
Sync barrierssync-barriers.md66
Late legalizationlate-legalization.md59
Hot/cold splittinghot-cold.md10
GMMA pipelininggmma-pipeline.md47
Uniform registersuniform-regs.md22
Register allocator corealgorithm.md50
Spillingspilling.md54
ABI handlingabi.md87
Scheduling overviewoverview.md112
Scheduling algorithmalgorithm.md121
Latency model & HW profileslatency-model.md78
Scoreboards & barriersscoreboards.md56
ISel pattern matchingisel.md182
SASS encodingencoding.md92
Peephole optimizationpeephole.md67
Mercury IR conversionmercury.md79
SASS templatestemplates.md46
SASS printing / renderersass-printing.md96
Capsule Mercurycapmerc.md20
Intrinsic infrastructureindex.md159
Math intrinsicsmath.md42
Tensor core intrinsicstensor.md45
Sync & warp intrinsicssync-warp.md65
SM targets & featuresindex.md70
ELF emitterelf-emitter.md29
ELF sectionssections.md33
Debug info (DWARF)debug-info.md33
Relocationsrelocations.md19

By Address Range

Functions in the binary are clustered by subsystem. This table maps address ranges to the pages that document them.

Address RangePrimary SubsystemKey Pages
0x400000-0x424000Entry, static init, mainentry.md, binary-layout.md
0x424000-0x42E000Memory pools, hash maps, listsmemory-pools.md, hash-bitvector.md
0x42E000-0x446000Diagnostics, CLI parsingcli-options.md, entry.md
0x446000-0x452000Compilation driveroverview.md, entry.md
0x452000-0x4D5000PTX parser & validatorptx-parser.md, ptx-directives.md
0x4D5000-0x5AA000PTX-to-ORI, early IRptx-to-ori.md, instructions.md
0x5AA000-0x612000Intrinsic infrastructureindex.md, math.md, tensor.md
0x612000-0x67F000Section builder, target configsections.md, index.md
0x67F000-0x6E4000Scheduling engine, OCG lowering, encodingoverview.md, encoding.md
0x6E4000-0x754000SASS codegen, SASS pipelinemercury.md, overview.md
0x754000-0x7C0000Liveness, knobs, bitfield encodingliveness.md, knobs.md, encoding.md
0x7C0000-0x8FE000Peephole, SASS mnemonics, scheduling upperpeephole.md, algorithm.md
0x8FE000-0x9D3000Register allocatoroverview.md, algorithm.md, abi.md
0x9D3000-0xAA8000Post-regalloc, named phases, rematrematerialization.md, phase-manager.md
0xAA8000-0xC52000Mega-passes, sync barriers, dataflowsync-barriers.md, general-optimize.md
0xC52000-0xD27000Phase manager, phase factoryphase-manager.md, optimizer.md
0xD27000-0x10B7000592 SASS encoder bodiesencoding.md, isel.md
0x10B7000-0x1225000Field encoders, ISel helpersencoding.md, isel.md
0x1225000-0x13CF000Bitvector, ISel coordinatorshash-bitvector.md, isel.md
0x13CF000-0x17F8000SM-specific ISel, pattern matchers, templatesisel.md, templates.md
0x17F8000-0x1C21000SASS printing, peephole mega-dispatcherssass-printing.md, peephole.md
0x1C21000-0x1CE3000ELF emitter, capsule mercury, relocationself-emitter.md, capmerc.md

Statistics

Top 10 Most-Called Functions

RankAddressIdentityCallers
10x7B9B80bitfield_insert18,347
20x10B61801-bit boolean encoder8,091
30x7BC030encode_register_operand6,147
40x4280C0get_thread_local_context3,928
50x42BDB0fatal_OOM_handler3,825
60x424070pool_alloc3,809
70x426150hashmap_put2,800
80x7B9D30clear_const_buffer_slots2,408
90x7B9D60encode_reuse_flags_predicate2,408
100x42FBA0diagnostic_emit2,350

Top 5 Largest Functions

RankAddressIdentitySize
10x5FF700intrinsic_prototype_emitter354 KB
20x169B190isel_pattern_dispatch280 KB
30x198BCD0sm100_peephole_dispatch233 KB
40x143C440sm120_peephole_dispatch233 KB
50x5C7A50wmma_mma_codegen173 KB

Top 10 Most Cross-Referenced (by wiki page count)

RankAddressIdentityPages
10x424070pool_alloc19
20x7DDB50phase_run_dispatch14
30x446240real_main13
30x5D4190intrinsic_dispatch_builder13
50x781F80instruction_iterator12
50xC60D30phase_factory12
70x9253C0instruction_operand_get11
70x612DE0section_attr_builder11
70x426150hashmap_put11
70x426D60hashmap_get11

Documentation Coverage

MetricCount
Total unique functions documented2,063
Wiki pages with function maps70
Functions in 5+ pages (high cross-reference)89
Functions in 1 page only (subsystem-internal)1,324
Confidence CERTAIN~40
Confidence HIGH~1,400
Confidence MEDIUM~620