NVVM Builtin Table Structure

770 builtins mapped to integer IDs (1--770) in a wyhash open-addressing hash table. Dual tables exist: pre-optimization (sub_90AEE0) and post-optimization (sub_126A910), both with identical content but separate address spaces.

Architecture

sub_913450 (public API: name -> builtin ID)
  |
  +-- Guard: context+492 == 0?
  |    +-- sub_90AEE0 (lazy init: populate all 770 entries, once)
  |
  +-- strlen(name)
  +-- sub_C92610(name, len)         -> compute wyhash
  +-- sub_C92860(context+480, ...)  -> quadratic probe find
  |
  +-- return *(uint32*)(entry + 8)  -> the builtin ID

Hash Table Infrastructure

The builtin name table uses a specialized 20-byte hash table header at context+480 with a parallel hash cache array and wyhash-v4 string hashing. The table employs quadratic probing with triangular-number increments and grows at 75% load factor. For 770 entries the capacity sequence is 16 -> 32 -> 64 -> 128 -> 256 -> 512 -> 1024.

Full structural details -- table layout, bucket format, string entry format, wyhash length-dispatch table with pseudocode, probing algorithm, triple-gated comparison guard, rehash procedure, and sentinel values -- are documented in Hash Table and Collection Infrastructure. The "wyhash v4 String Hasher" and "Probing Strategy" sections on that page are the canonical references.

Complete Builtin ID Inventory

Synchronization & Compiler Intrinsics (IDs 1–7)

Cluster Operations — SM 90+ (IDs 8–14)

Barrier Extensions (IDs 15–20)

Bit Manipulation (IDs 21–26)

__nvvm_clz_{i,ll}, __nvvm_popc_{i,ll}, __nvvm_brev{32,64}

Math — Rounding/Abs/Saturate (IDs 27–56)

__nvvm_{floor,ceil,abs,fabs,round,trunc,saturate}_{ftz_f,f,d}, __nvvm_{ex2,lg2,sin,cos}_approx_{ftz_f,f,d}

Reciprocal / Sqrt / Rsqrt (IDs 57–87)

__nvvm_rcp_{rn,rz,rm,rp}_{ftz_f,f,d}, __nvvm_sqrt_{f,rn,rz,rm,rp}_{ftz_f,f,d}, __nvvm_rsqrt_approx_{ftz_f,f,d}

Type Conversions (IDs 88–184)

97 entries covering all float↔int, double↔int, float↔half, bitcast combinations with all four rounding modes and FTZ variants.

Address Space & Memory Queries (IDs 185–204)

Atomic Operations — Legacy NVVM (IDs 207–275)

69 entries: __nvvm_atom_{,cta_,sys_}{add,xchg,min,max,inc,dec,and,or,xor}_gen_{i,ll,f,d,ui,ull,128}

FP Arithmetic (IDs 276–349)

__nvvm_{min,max}_{i,ui,ll,ull}, __nvvm_f{min,max}_{f,ftz_f,d}, __nvvm_mulhi_{i,ui,ll,ull}, __nvvm_mul_{rn,rz,rm,rp}_{ftz_f,f,d}, __nvvm_div_*, __nvvm_add_*

Vote Operations (IDs 351–358)

__nvvm_vote_{all,any,uni,ballot} + _sync variants

Match Operations (IDs 361–364)

__match{32,64}_{any,all}_sync

FMA (IDs 383–403)

__nvvm_fma_{rn,rz,rm,rp}_{ftz_f,f,d,ftz_f2,f2}

C++11 Atomics (IDs 417–473)

Sized variants: __nv_atomic_{load,store,fetch_add,fetch_sub,fetch_and,fetch_or,fetch_xor,fetch_max,fetch_min,exchange,compare_exchange}_{1,2,4,8,16}_{u,s,f}

Surface Stores — sust (IDs 474–638)

165 entries covering __nvvm_sust_b_{1d,1d_array,2d,2d_array,3d}_{i8,...,v4i32}_{clamp,trap,zero}.

Pattern: sust_b_<dim>_<type>_<oob_mode> across 5 dimensions × 11 types × 3 OOB modes.

CUDA Varargs (IDs 639–642)

__cu_va_{start,end,arg,copy}

Tex/Surf Handler (ID 647)

__nv_tex_surf_handler — generic dispatch for texture/surface reads (surface stores use the dedicated sust builtins above).

C++ ABI (IDs 648–677)

__cxa_vec_{ctor,cctor,dtor,new2,new,new3,delete2,delete,delete3}, __gen_nvvm_mem{cpy,set}_*, _Znw{j,m,y}, _Zna{j,m,y}, _ZdlPv{,m,y}, _ZdaPv{,m,y}

WMMA Tensor Core — SM 70+ (IDs 678–707)

30 entries: __hmma_m{16n16k16,32n8k16,8n32k16}_{ld_a,ld_b,ld_c_f16,ld_c_f32,st_c_f16,st_c_f32,mma_f16f16,mma_f32f16,mma_f16f32,mma_f32f32}

Integer/Binary Tensor Core — SM 75+ (IDs 708–745)

38 entries: __imma_m{16n16k16,32n8k16,8n32k16}_{ld_a,ld_b,ld_c,st_c,mma}_{s8,u8}, __imma_m8n8k32_{s4,u4}, __bmma_m8n8k128_{b1}

Extended Tensor Core — SM 80+ (IDs 746–764)

__dmma_m8n8k4_mma_f64, __mma_tf32_m16n16k8_mma_f32, __mma_bf16_m*_mma_f32 + load/store variants

WGMMA — SM 90+ (IDs 765–768)

__wgmma_mma_async_{f16,bf16,tf32,f8}

Alloca (IDs 769–770)

_alloca, __builtin_alloca

Category Summary

SM Generation Coverage

All 770 builtins are registered regardless of target SM. Architecture gating happens in the lowering layer that consumes the builtin IDs.

Key Observations

• Lazy initialization: The entire table is built on first lookup. Guard: context+492 != 0.
• No texture reads (suld): Only surface store builtins are registered. Texture/surface reads go through __nv_tex_surf_handler (ID 647).
• Write-once table: Tombstone mechanics exist but deletions never occur for the builtin table.
• Duplicate prefix optimization: IDA shows SSE xmmword constant loads for long common prefixes (__nvvm_sust_b_2d_array_*) — this is compiler optimization of string literal loads, not a different code path.