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

elfLink Error Codes and Linker Diagnostic Catalog

The elfLink subsystem is nvlink's internal library for loading, compiling, and linking device code modules. It wraps libnvvm for NVVM IR compilation, manages cubin/fatbinary extraction, and interfaces with the linker library (libnvvm.so). When any of these operations fail, elfLink returns an integer error code in the range 0--13. The function sub_4BC270 at 0x4BC270 translates these codes into human-readable strings via a lookup table at off_1D489E0.

Beyond the elfLink subsystem's 14 return codes, nvlink emits over 100 distinct user-facing diagnostic messages through the diag_emit (sub_467460) pipeline. This page catalogs every error string, maps each to its source function and triggering conditions, and provides user-facing fix suggestions.

CodeString (addr)Symbolic Name
0(success -- no error)ELFLINK_OK
1elfLink internal error (0x1D4883F)ELFLINK_INTERNAL
2elfLink error on cubin (0x1D48856)ELFLINK_CUBIN_ERROR
3elfLink fatbinary error (0x1D4886D)ELFLINK_FATBIN_ERROR
4elfLink memory error (0x1D48885)ELFLINK_MEMORY_ERROR
5elfLink JIT compile failed (0x1D4889A)ELFLINK_JIT_COMPILE
6elfLink JIT link failed (0x1D488B5)ELFLINK_JIT_LINK
7elfLink nvvm error (0x1D488CD)ELFLINK_NVVM_ERROR
8elfLink error cubin not relocatable (0x1D488E0)ELFLINK_NOT_RELOCATABLE
9elfLink error cubin not compatible (0x1D48908)ELFLINK_NOT_COMPATIBLE
10elfLink cubin arch not compatible (0x1D48930)ELFLINK_ARCH_MISMATCH
11elfLink linker library load error (0x1D48958)ELFLINK_LIB_LOAD
12elfLink error incompatible formats (0x1D48980)ELFLINK_FORMAT_MISMATCH
13elfLink error finalization failed (0x1D489A8)ELFLINK_FINALIZE_FAILED
>13elfLink: unexpected error (0x1D48760)(default fallback)

The symbolic names above are inferred from the string content; the binary uses raw integer constants.

Lookup Function

// sub_4BC270 at 0x4BC270
// Table: off_1D489E0 -- 14 const char* entries (codes 0..13)
const char *elflink_error_string(unsigned int code)
{
    if (code <= 13)
        return error_table[code];   // off_1D489E0[code]
    return "elfLink: unexpected error";
}

The function is called from five sites in the binary:

CallerAddressContext
sub_4297B00x4297B0Generic error handler -- translates any non-zero elfLink code (except 4 and 7) into a fatal diagnostic
sub_42AF400x42AF40Input processing loop -- handles per-module load failures during LTO and cubin extraction
sub_42A2D00x42A2D0Host-ELF extraction path -- iterates embedded cubins and reports per-cubin elfLink errors
sub_427A100x427A10LTO add-module path -- loads NVVM IR into the LTO program object
main0x409800Top-level pipeline -- reports errors from library loading (sub_4BC470) and final LTO compilation (sub_4BC6F0)

Detailed Error Code Reference

Code 0 -- Success

No error. Returned when module loading, compilation, or linking completes without failure. Never passed to sub_4BC270 in practice.

Code 1 -- Internal Error

String: elfLink internal error (0x1D4883F)

Trigger: Returned by sub_4BDAC0 (host-ELF cubin extraction at 0x4BDAC0) and sub_4BDAF0 (host-ELF iterator at 0x4BDAF0) when the underlying ELF parser (sub_487C20, sub_487E10) returns an unrecognized status code (value > 2). Also returned by sub_4BC290 (library initialization at 0x4BC290) when the top-level context pointer is null or when nvvmCreateProgram fails. In the LTO compilation path (sub_4BC6F0), returned when nvvmGetProgramLog API calls fail.

Diagnosis: This is a catch-all for unexpected internal states. The input file may be corrupt, or an internal invariant was violated. Check that input object files are valid ELF. If the error persists with known-good inputs, it indicates a linker bug.

User fix: Verify input files with readelf -h to confirm valid ELF headers. Recompile inputs with the same toolkit version as nvlink.

Code 2 -- Cubin Error

String: elfLink error on cubin (0x1D48856)

Trigger: Produced when cubin extraction from a host ELF fails. The secondary table dword_1D48A50 maps parser status 1 to elfLink code 2. The parser returns status 1 when it encounters a malformed cubin section inside a host object file.

Diagnosis: The input .o file contains embedded device code (in .nv_fatbin or similar sections) that could not be parsed as a valid cubin. Recompile the source with the same toolkit version as nvlink.

User fix: Recompile the failing source file. If the file was produced by an older CUDA toolkit, upgrade the compilation to match the linker version.

Code 3 -- Fatbinary Error

String: elfLink fatbinary error (0x1D4886D)

Trigger: Produced when fatbinary extraction or decompression fails. The secondary table maps parser status 2 to elfLink code 3. This covers failures in the fatbinary header parsing, unsupported compression formats, or truncated fatbinary data.

Diagnosis: The fatbinary container embedded in the host object is invalid. Verify that the input was compiled with a compatible nvcc version. Truncated files (e.g., from interrupted builds) commonly produce this error.

User fix: Rebuild from scratch (make clean && make). Check for filesystem corruption or incomplete file copies.

Code 4 -- Incompatible Code (Memory Error)

String: elfLink memory error (0x1D48885)

Special handling: Code 4 receives unique treatment in sub_4297B0 -- instead of passing through sub_4BC270, the handler constructs a custom message by prepending a descriptive prefix (loaded from xmmword_1D34750 / xmmword_1D34760) and appending " code in <filename>". This produces a user-facing message like "elfLink found incompatible code in foo.o" rather than the generic table string. The same custom-message logic is duplicated in sub_42A2D0 for the host-ELF iteration path.

Trigger: Returned by sub_4BD0A0 (NVVM IR compilation driver at 0x4BD0A0) when the compilation pipeline fails at any stage: target architecture setup (sub_4CE2F0), debug mode configuration (sub_4CE380), 64-bit mode configuration (sub_4CE640), module addition (sub_4CE070), or final compilation (sub_4CE8C0). Also returned by sub_4BD240 (cubin post-processing at 0x4BD240) when ABI validation fails (-m32/-m64 mismatch) or when the cubin bytecode extractor (sub_4BE350) fails.

Diagnosis: Usually indicates a toolkit version mismatch. The cubin or NVVM IR module was compiled with options incompatible with the current link target. Check that all input objects target the same sm_ architecture and address size (32-bit vs 64-bit).

User fix: Ensure all .o files were compiled with the same -arch=sm_XX and -m64/-m32 flags. Check nvcc --version matches the nvlink version.

Code 5 -- JIT Compile Failed

String: elfLink JIT compile failed (0x1D4889A)

Trigger: Returned by sub_4BD0A0 when sub_4CE8C0 (the NVVM compilation call) returns a failure status other than 3 (which maps to code 7 instead). Also returned by sub_4BD240 when ABI checks fail (-m32/-m64 validation against sub_4CE3E0), or when a pass-through option string is rejected.

Diagnosis: The embedded NVVM IR or PTX could not be compiled to SASS for the target architecture. Check that the source was compiled for a compatible compute_ capability. If -dlto is in use, verify that all LTO objects were compiled with the same major CUDA toolkit version.

User fix: Verify the compute_ and sm_ targets are compatible. With LTO (-dlto), all objects must use the same CUDA toolkit major version.

String: elfLink JIT link failed (0x1D488B5)

Trigger: This code is only reachable through the error table. In the analyzed binary, no producer was found that explicitly returns 6. It is reserved for the case where the libnvvm link step (post-compilation module merging) fails, as distinct from compilation failure.

Diagnosis: If encountered, it means the NVVM linker was unable to merge compiled modules. This can happen when symbol visibility or linkage conflicts prevent merging.

User fix: Ensure that device-side extern declarations match across translation units. Check for conflicting function prototypes in device code headers.

Code 7 -- NVVM Error (Unresolved References)

String: elfLink nvvm error (0x1D488CD)

Special handling: Code 7 receives unique treatment parallel to code 4. In sub_4297B0, sub_42A2D0, and sub_42AF40, when the error code is 7 the handler checks byte_2A5F298 (the -lto-allow-unresolved flag at 0x2A5F298) and whether the filename contains "cudadevrt". If the flag is set or the input is libcudadevrt, the error is silently suppressed. Otherwise, a fatal diagnostic is emitted via descriptor unk_2A5B660:

nvvm IR for arch <target_arch> not found in <filename>

Trigger: Returned by sub_4BD0A0 when sub_4CE8C0 returns status 3, indicating that the NVVM compilation produced NVVM IR output (not SASS) -- meaning the module contains unresolved references that require a subsequent link step. This is expected for LTO intermediate objects and libcudadevrt.

Diagnosis: Typically not a true error. It indicates the module contains NVVM IR that cannot be finalized to SASS in isolation because it has unresolved references. When linking with -dlto, this is normal for all LTO input objects. It becomes a user-visible error only when encountered during non-LTO linking without -lto-allow-unresolved.

User fix: If you see this error, add -dlto to enable LTO linking. For libraries containing device runtime code, this is expected behavior.

Code 8 -- Cubin Not Relocatable

String: elfLink error cubin not relocatable (0x1D488E0)

Trigger: Returned by sub_4BD240 (cubin post-processing at 0x4BD240) when the cubin extraction via sub_4BE350 fails and the resulting context indicates no compilation output was produced (the compilation output pointer at offset +16 of the context is null). This means the cubin was compiled without relocatable device code (-rdc / -dc), making it unlinkable.

Diagnosis: The input cubin was compiled as a standalone (whole-program) object without separate compilation. Recompile with nvcc -dc (device code compilation) to produce relocatable device objects.

User fix: Add -dc (or equivalently -rdc=true) to the nvcc compile command. For static libraries, ensure they were built with -rdc=true.

Code 9 -- Cubin Not Compatible

String: elfLink error cubin not compatible (0x1D48908)

Trigger: Produced during cubin validation when the object's ELF metadata does not match the link-time target. This covers ABI version mismatches, size mismatches (32-bit vs 64-bit), and general compatibility flags that prevent linking.

Diagnosis: The input cubin was compiled with an incompatible ABI version or addressing mode. Check the nvcc version used to compile each input. The linker emits more specific diagnostics before this error:

  • "Input file '%s' abi does not match" (0x1D34C68)
  • "Input file '%s' size does not match target '%s'" (0x1D34C90)
  • "Input file '%s' ABI version '%u' is incompatible with target ABI version '%u'" (0x1D34CF0)

User fix: Recompile all objects with the same toolkit version and matching address-size flags.

Code 10 -- Architecture Not Compatible

String: elfLink cubin arch not compatible (0x1D48930)

Trigger: Returned by sub_4BC290 (library initialization at 0x4BC290) when libnvvm.so cannot be loaded (the dlopen handle is null) or when the __nvvmHandle or nvvmCreateProgram symbols cannot be resolved via dlsym. Also returned by sub_4BC4A0 when similar symbol resolution failures occur for __nvvmHandle during the add-module call. In the LTO compilation path (sub_4BC6F0), returned when nvvmGetCompiledResult symbol lookup fails.

Diagnosis: Despite the name suggesting architecture incompatibility, this code primarily means the linker library infrastructure is unavailable. The libnvvm.so library could not be loaded or is from an incompatible toolkit version.

User fix: Verify that the CUDA toolkit installation is complete and that libnvvm.so is on the library search path (set via --libnvvm-path or the toolkit's lib64/ directory). Check LD_LIBRARY_PATH if needed.

Code 11 -- Linker Library Load Error

String: elfLink linker library load error (0x1D48958)

Trigger: Returned by sub_4BC4A0 (the add-module-to-program call at 0x4BC4A0) when the __nvvmHandle callback returns a non-null function pointer but calling that function with the module data returns a non-zero status. This means libnvvm.so was loaded successfully but rejected the specific module being added.

Diagnosis: The libnvvm runtime rejected the NVVM IR module. The module may be in an unsupported bitcode format version, or contain constructs not supported by the installed libnvvm.

User fix: Ensure that all input objects were compiled with the same CUDA toolkit version as the nvlink binary. Check nvcc --version and nvlink --version output.

Code 12 -- Incompatible Formats

String: elfLink error incompatible formats (0x1D48980)

Trigger: This code is reachable through the error table but no direct producer was identified in the analyzed call paths. It is reserved for cases where the input module format (cubin vs PTX vs NVVM IR) is incompatible with the requested link operation.

Diagnosis: The linker encountered a module in a format it cannot process in the current mode. For example, attempting to link a Mercury (capmerc) object in a CUDA-only link, or mixing incompatible ELF object formats.

User fix: Ensure all inputs are of the same format type. Do not mix Mercury objects with standard CUDA objects unless the linker is configured for Mercury mode.

Code 13 -- Finalization Failed

String: elfLink error finalization failed (0x1D489A8)

Trigger: This code maps to failures in the finalization (off-target to on-target compilation) step of the pipeline.

Diagnosis: The finalizer could not convert the intermediate representation to native SASS for the target architecture. This typically occurs during forward-compatibility finalization when the input cubin targets a different SM architecture. See also the separate "Internal FNLZR error '%s'" diagnostic (0x1D34F74) which provides more detail.

User fix: Ensure the target architecture (-arch) matches the sm_ target that inputs were compiled for. If using forward compatibility, verify the toolkit supports the target.

Error Dispatch Logic

The generic error handler sub_4297B0 at 0x4297B0 implements a three-way dispatch:

sub_4297B0(error_code, filename):
    if error_code == 0:
        return                          // success, nothing to report
    if error_code == 7:
        if lto_allow_unresolved or "cudadevrt" in filename:
            return                      // silently suppress
        fatal("nvvm IR for arch %s not found in %s", target_arch, filename)
    if error_code == 4:
        msg = "elfLink found incompatible code in " + filename
        fatal(msg)
    else:
        msg = elflink_error_string(error_code)   // sub_4BC270
        fatal(msg)

The same dispatch logic is replicated in sub_42A2D0 (host-ELF extraction path), which additionally handles the host-ELF iterator loop (sub_4BDAF0) -- calling the dispatch for each extracted cubin.

All fatal diagnostics go through sub_467460, which formats the message with the "error " severity prefix and writes to stderr. When --Werror is active, warnings are promoted to the "error* " level.

Part 2: Broader Diagnostic Message Catalog

Beyond the 14 elfLink error codes, nvlink emits a comprehensive set of user-facing diagnostics during linking. These are organized by severity and category.

Severity Levels

The diagnostic system at 0x1D3C660 defines six severity levels:

LevelValuePrefixMeaning
note0(none)Silent; diag_emit returns immediately
info1(empty)Informational; suppressed by --disable-infos
info2info Informational with explicit label
warning3warning Non-fatal issue, link continues
error*4error* Warning promoted to error by --Werror
error5error Fatal error, link aborted
fatal6fatal Unrecoverable internal failure

Category 1: Input File Errors

MessageAddrSevSourceTriggerUser Fix
Cannot open file '%s'0x1D39250fatalsub_426570Input file does not exist or is not readableVerify file path and permissions
Bad file name '%s'0x1D39266fatalsub_427AE0Input filename is empty or contains invalid charactersCheck filename for special characters
Could not open input file '%s'0x1D34E10errorsub_426570Input file fopen returns nullVerify file exists and is not locked
Could not read file '%s'0x1D34FE0errormainI/O error during input readCheck disk and filesystem health
Could not find fatbinary in '%s'0x1D34C40errormainFatbinary not present in input fileVerify input was compiled with -rdc or -dc
Library file '%s' not found in paths0x1D34BF0errormain-l library not found in any -L pathAdd -L<dir> with the correct library directory
Unsupported file type '%s'0x1D34FABfatalsub_427AE0Input is not ELF, archive, PTX, or fatbinaryVerify file type with file <input>
No input files specified; use option --help for more information0x1D34ED0fatalmainNo object files provided on command lineProvide at least one input .o file
Too many options file opened '%s'0x1D357A0warningsub_42EC10Response-file nesting exceeds 14 levelsReduce @file nesting depth
Failed to read contents of options file0x1D357E8errorsub_42EC10I/O error reading @fileCheck response file path and permissions
Could not open options file '%s'0x1D35810errorsub_42EC10Response file fopen failedVerify @file path is accessible
Skipping incompatible '%s' when searching for -l%s0x1D34AB8warningmainLibrary found but wrong architectureProvide architecture-correct library

Category 2: Output File Errors

MessageAddrSevSourceTriggerUser Fix
Must specify output file with -o option0x1D34E30fatalmainMissing -o flagAdd -o <output> to command line
Could not open output file '%s'0x1D34DC8errormainOutput path is not writableCheck write permissions on target directory
Could not write file '%s'0x1D34FC6errormainI/O error during output writeCheck disk space and filesystem
Host linker script creation failed.0x1D34E58errormainInternal failure generating host linker scriptReport as bug; verify toolkit installation

Category 3: Architecture and Compatibility Errors

MessageAddrSevSourceTriggerUser Fix
Unknown arch name '%s'0x1D39279fatalsub_427AE0The -arch value is not a recognized sm_ targetUse a valid arch like sm_70, sm_80, sm_90, etc.
SM Arch ('%s') must be >= 200x1D34F8Efatalsub_427AE0Target architecture too oldUse sm_20 or newer
SM Arch ('%s') not found in '%s'0x1D34BC8errormainFatbinary does not contain code for the target archRecompile input for the correct -arch target
Cannot target %s when input '%s' is SASS0x1D34850errorsub_427AE0SASS object cannot be retargeted to a different architectureRecompile source for the desired target arch
Input file '%s' arch does not match target '%s'--errorsub_426570Object compiled for wrong sm_ targetRecompile with matching -arch flag
Input file '%s' abi does not match0x1D34C68errorsub_426570ELF ABI version mismatchRecompile with matching toolkit version
Input file '%s' size does not match target '%s'0x1D34C90errorsub_42657032-bit vs 64-bit object/target mismatchMatch -m32/-m64 across all objects
Input file '%s' ABI version '%u' is incompatible with target ABI version '%u'0x1D34CF0errorsub_426570Detailed ABI version incompatibilityRecompile with current toolkit
Target format and SM Arch ('%u') mismatch0x1D34D40errorsub_426570Cubin format does not match SM generationVerify -arch and -code flags are consistent
Input file '%s' must be recompiled with toolkit >= Cuda 12.00x1D34AF0errorsub_426570Object from pre-12.0 toolkitRecompile with CUDA 12.0 or later
Input file '%s' must be recompiled with toolkit >= Cuda 7.00x1D34B30errorsub_426570Object from pre-7.0 toolkitRecompile with CUDA 7.0 or later
Object '%s' cannot be linked due to version mismatch. Objects using tcgen05 in 12.x cannot be linked with 13.0 or later, they must be rebuilt with latest compiler0x1D39330errormaintcgen05 objects from CUDA 12.x linked with 13.0+Rebuild all objects with CUDA 13.0+ toolkit
Cannot link sanitized object '%s' from version %d with sanitized object from a different toolkit version (%d)0x1D393D8errormainSanitizer-instrumented objects from different toolkit versionsCompile all sanitized objects with same toolkit
For kernel functions with parameter size higher than 4k bytes on sm_7x and sm_8x, all objects must be compiled with 12.1 or later0x1D39488errormainLarge kernel params require recent toolkitRecompile all objects with CUDA 12.1+
%s '%s' is not compatible with driver version %s0x1D3ADE0warningsub_451D80Output binary not driver-compatibleUpdate GPU driver or target older sm_ arch

Category 4: Symbol Resolution Errors

MessageAddrSevSourceTriggerUser Fix
Multiple definition of '%s' in '%s', first defined in '%s'0x1D39A20errorsub_440BE0Symbol multiply defined across objectsRemove duplicate definitions; use static/__device__ scope
Undefined reference to '%s' in '%s'0x1D39A60errorsub_445000Unresolved symbol at link timeAdd the missing object file or library to the link
Prototype doesn't match for '%s' in '%s', first defined in '%s'0x1D398B0warningsub_450ED0Kernel parameter signature mismatch across TUsEnsure consistent kernel declarations in all TUs
Size doesn't match for '%s' in '%s', first specified in '%s'0x1D398F0warningsub_450ED0Symbol size conflict between definitionsCheck for struct layout differences between TUs
Duplicate weak parameter bank for '%s' is not the same size0x1D39690warningsub_437BB0Weak symbol parameter banks differEnsure weak symbols have consistent parameter sizes
Cannot link shared object '%s' with partially linked kernel '%s'0x1D396D0errormainShared object conflicts with partial linkRestructure the link to avoid mixing shared and partial objects
Could not replace weak symbol '%s'0x1D39888warningsub_450ED0Weak symbol replacement failureCheck for conflicting weak/global declarations
Section '%s' not found in object '%s'0x1D39610warningsub_451D80Expected ELF section missing from inputVerify input was compiled correctly for the link mode
Output is limited to run on SM Arch '%u' only, due to missing attribute '%s' in input file '%s'0x1D39550warningmainInput file lacks feature attributeRecompile input with current toolkit for full compatibility
Output is limited to run on SM Arch '%u' only due to missing attribute '%s' in input file '%s'0x1D395B0warningmainVariant of above with slightly different punctuationSame as above

Category 5: Resource Limit Errors

MessageAddrSevSourceTriggerUser Fix
File uses too much global %s data (0x%llx bytes, 0x%x max)0x1D39B88errorsub_445000Global/constant memory overflowReduce global variable sizes; split across modules
Entry function '%s' uses too much %s data (0x%llx bytes, 0x%x max)0x1D39BC8errorsub_445000Per-kernel resource limit exceededReduce per-kernel memory usage
More than %d %s used in entry function '%s'--errorsub_438DD0Register/barrier count exceededReduce register pressure; use __launch_bounds__
Entry function '%s' uses too much data for compiler-generated constants; please recompile with -Xptxas --disable-optimizer-constants0x1D39B00errorsub_445000Constant bank overflow from optimizationsAdd -Xptxas --disable-optimizer-constants
Stack size for entry function '%s' cannot be statically determined0x1D39A88warningsub_44AD40Indirect calls prevent stack analysisUse --suppress-stack-size-warning to silence
Function '%s' uses %d bytes stack but limited to %d--warningsub_44C030Stack usage exceeds budgetReduce function stack usage or increase limit

Category 6: CLI Option Errors

MessageAddrSevSourceTriggerUser Fix
Unknown option '%s'0x1D34F45fatalsub_427AE0Unrecognized CLI flagCheck nvlink --help for valid options
Invalid argument for option '%s'0x1D34DE8errorsub_427AE0Option value cannot be parsedProvide a valid value for the option
Conflicting options '%s' and '%s'0x1D34BA0errorsub_427AE0Mutually exclusive CLI flagsRemove one of the conflicting options
Incompatible options '%s' and '%s', using '%s'0x1D347C8warningsub_427AE0Options conflict but one takes precedenceReview option interactions
option '%s' has been deprecated0x1D357C8warningsub_427AE0Deprecated CLI option usedMigrate to the recommended replacement option
incompatible redefinition for option '%s', the last value of this option was used0x1D358E8warningsub_427AE0Option specified multiple times with different valuesSpecify each option only once
--cuda-api-version major number must be == toolkit version0x1D33DF0errorsub_429BA0API version mismatch with toolkitUse matching --cuda-api-version
--cuda-api-version value cannot be parsed0x1D33E30errorsub_429BA0Malformed version stringProvide version as major.minor
Option for ABI >= 8 is not compatible with -m32, hence ignored0x1D34788warningsub_427AE032-bit mode incompatible with ABI 8+Use -m64 for modern ABI levels

Category 7: LTO-Specific Errors

MessageAddrSevSourceTriggerUser Fix
Ignoring -dlto option because no LTO objects found0x1D34998warningmain-dlto specified but no LTO inputsRemove -dlto or ensure inputs contain NVVM IR
Some objects do not have '%s' specified but others do; will build everything with '%s=%d'--warningmainMixed compilation options across LTO objectsCompile all objects with the same options
error in LTO callback0x1D3425Dfatalmain__nvvmHandle callback returned errorInternal error; report bug with reproduction steps
could not find __nvvmHandle0x1D34241fatalmaindlsym for __nvvmHandle failedVerify libnvvm.so is from the correct toolkit version
could not find CALLBACK Handle0x1D345C0fatalmainSecond-level callback resolution failedSame as above; toolkit installation issue
Call to ptxjit failed in extended split compile mode0x1D345E0fatalmainThread pool work item failed during split compileReduce parallelism or report bug
Ptxjit compilation failed in extended split compile mode0x1D34618fatalmainPTX-to-SASS compilation failed in split modeCheck PTX compatibility with target arch
Cannot allocate pthread data0x1D342BEfatalmainmalloc failed for thread pool dataIncrease system memory or reduce parallelism
Unable to create thread pool0x1D342DBfatalmainThread pool initialization failedCheck system thread limits (ulimit -u)
unexpected object after cudadevrt0x1D346D8fatalmainInput ordering violation in LTO pipelineInternal error; report with input files

Category 8: Warning-Level Diagnostics

MessageAddrSevSourceTriggerUser Fix
Option '%s' is not fully implemented for gpu architecture '%s' and may not work as expected--warningsub_427AE0Feature partially implemented for targetCheck release notes for arch-specific feature support
Option '%s' not supported for gpu architecture '%s'0x1D348E0warningsub_427AE0Feature not available for targetRemove the unsupported option
Input file '%s' newer than toolkit (%d vs %d)0x1D34B70warningsub_426570Forward-compatibility concernUpdate nvlink to match input file toolkit version
Function '%s' and function '%s' have conflicting cache preference, falling back to use cache preference of entry '%s'0x1D39760warningsub_451D80Cache preference conflict across functionsUse consistent cache preference attributes

Category 9: Internal Fatal Errors

These indicate linker bugs or unexpected states. They are routed through descriptor unk_2A5B670 (severity 6, fatal) or unk_2A5B990 (severity 6, internal assertion):

MessageAddrSourceContext
Internal error: Aborting0x1D39290variousUnrecoverable internal assertion failure
Internal error: %s0x1D39C80variousInternal assertion with detail string
Internal FNLZR error '%s'0x1D34F74sub_4275C0Finalizer returned an error string
Bailing out due to earlier errors0x1D3B978mainAccumulated error count exceeded threshold
merge_elf failed0x1D34360mainELF merge phase set fatal-error flag
expected libcudadevrt object--mainLTO pipeline expected cudadevrt but found different input
linking with -ewp objects requires using current toolkit--sub_426570EWP objects from different toolkit version
cubin not an elf?--mainInput claimed to be cubin but failed ELF validation
cubin not a device elf?--mainELF is valid but not a device ELF
fatbin wrong format?--mainFatbinary header validation failed
should never see bc files--mainUnexpected bitcode file in input stream
unexpected cpuArch0x1D34002sub_42A2D0Host CPU architecture not recognized

Category 10: External Tool Errors

MessageAddrSevSourceTriggerUser Fix
exec() failed with '%s'0x1D3B902fatalsub_42FA70External tool launch failedVerify tool path and permissions
wait() failed with '%s'0x1D3B91Afatalsub_42FA70waitpid returned errorSystem-level error; check ulimit and process limits
fork() failed with '%s'0x1D3B932fatalsub_42FA70Process fork failedReduce system load or increase process limits
Finalization failed '%s'0x1DFCBC4errorfinalizerExternal finalizer returned errorCheck target arch compatibility

Warning Suppression Flags

Several CLI options control diagnostic emission:

FlagEffect
--disable-warnings (0x1D325F0)Suppress all warning diagnostics
--Werror / --warning-as-error (0x1D3261E / 0x1D32625)Promote all warnings to errors
--suppress-stack-size-warning (0x1D32450)Suppress stack-size-related warnings
--suppress-arch-warning (0x1D3246C)Suppress architecture mismatch warnings
--extra-warnings (0x1D32735)Enable additional advisory warnings
--allow-undefined-globals (0x1D325C3)Allow undefined globals and their relocations in linked executable

Error Flow Architecture

Input Processing (sub_42AF40 / sub_42A2D0)
    |
    +-- sub_4BD0A0  (NVVM IR compilation)    --> codes 0, 5, 7
    +-- sub_4BD240  (cubin post-processing)  --> codes 0, 1, 5, 8
    +-- sub_4BDAC0  (host-ELF extraction)    --> codes 0, 1, 2, 3
    +-- sub_4BDAF0  (host-ELF iteration)     --> codes 0, 1, 2, 3
    |
    v
sub_4297B0  (generic error dispatch)
    |
    +-- code 0:  no-op
    +-- code 4:  custom "elfLink found incompatible code in <file>" message
    +-- code 7:  conditional suppress (cudadevrt or -lto-allow-unresolved) or arch-specific fatal
    +-- other:   sub_4BC270(code) --> error string --> sub_467460 (fatal via unk_2A5B670)

Library Loading (main / sub_427A10)
    |
    +-- sub_4BC470  (dlopen libnvvm.so)      --> sub_4BC290 --> codes 0, 1, 10
    +-- sub_4BC4A0  (add module to program)  --> codes 0, 10, 11
    |
    v
    sub_4BC270(code) --> error string --> sub_467460 (fatal via unk_2A5B670)

LTO Compilation (sub_4BC6F0)
    |
    +-- nvvmCompileProgram via dlsym         --> code 8 (log + compile error)
    +-- nvvmGetProgramLog                    --> code 1 (API failure)
    +-- nvvmGetCompiledResult                --> code 10 (symbol lookup failure)
    |
    v
    Returns raw code to caller for sub_4BC270 translation

Split Compilation (main, thread pool)
    |
    +-- sub_43FDB0 / sub_43FF50             --> "Call to ptxjit failed"
    +-- sub_4264B0  (worker function)        --> per-item error codes
    +-- sub_4297B0  (per-result dispatch)    --> <lto ptx> as filename
    |
    v
    Per-item errors aggregated; fatal on first failure

ELF Merge + Layout + Relocate + Finalize (main phases)
    |
    +-- sub_439830  (merge)                  --> via unk_2A5B990 internal assertions
    +-- sub_469D60  (layout)                 --> resource limits (unk_2A5BA60..BA80)
    +-- sub_445000  (relocate)               --> undef refs, resource limits, section errors
    +-- sub_4475B0  (finalize / output)      --> write errors (unk_2A5B710)

Key Implementation Details

Secondary mapping table (dword_1D48A50): Functions sub_4BDAC0, sub_4BDAF0, and sub_4BDB30 use a 3-entry mapping table at 0x1D48A50 to convert their internal parser return codes (0, 1, 2) into elfLink codes (0, 2, 3). Values > 2 map to code 1 (internal error).

setjmp/longjmp error recovery: Several elfLink functions (sub_4BC290, sub_4BC4A0, sub_4BD240) use setjmp/longjmp for non-local error recovery. This allows deep call stacks within libnvvm to unwind cleanly back to the elfLink entry point. The longjmp path always returns code 1 (internal error), protecting the linker from crashes inside the dynamically loaded libnvvm.

Code 7 suppression logic: The NVVM error code (7) is the only error code with conditional suppression. The flag byte_2A5F298 (-lto-allow-unresolved) and the filename check for "cudadevrt" together determine whether the error is fatal or silently ignored. This allows libcudadevrt's NVVM IR to pass through the non-LTO path without aborting. The check is performed identically in three functions: sub_4297B0, sub_42A2D0, and sub_42AF40.

Code 4 custom message construction: Rather than using the generic table string, code 4 constructs a message from two SSE-loaded 16-byte constants (xmmword_1D34750, xmmword_1D34760) forming the prefix "elfLink found incompatible", followed by " code in " and the filename. This provides file-specific context that the generic string lacks. The 42-byte allocation (strlen(filename) + 42) accounts for the 33-byte prefix plus " code in " (9 bytes). This construction is duplicated in sub_42A2D0's iterator loop.

Multiple errors aggregation: The string "Multiple errors:\n" at 0x1D486CE indicates nvlink can aggregate multiple NVVM errors into a single diagnostic. The nvvmGetErrorString function (0x1D487F1) is resolved via dlsym to retrieve libnvvm-specific error details.

Descriptor unk_2A5B990 dominance: The internal-assertion descriptor unk_2A5B990 accounts for 230 of the 380+ total diag_emit call sites. It serves as the universal internal-error descriptor (always fatal, severity 6). Every call site passes a literal string describing the violated invariant. 40+ unique assertion messages have been cataloged; see Error Reporting for the complete list.

Complete String Address Index

All error-related strings confirmed in nvlink_strings.json, organized by address:

AddressString
0x1D321C0specified arch exceeds buffer length
0x1D323AEInternal error
0x1D33DF0--cuda-api-version major number must be == toolkit version
0x1D33E30--cuda-api-version value cannot be parsed
0x1D33E5Funexpected data in module_ids
0x1D34002unexpected cpuArch
0x1D34241could not find __nvvmHandle
0x1D3425Derror in LTO callback
0x1D342BECannot allocate pthread data
0x1D342DBUnable to create thread pool
0x1D34360merge_elf failed
0x1D345C0could not find CALLBACK Handle
0x1D345E0Call to ptxjit failed in extended split compile mode
0x1D34618Ptxjit compilation failed in extended split compile mode
0x1D346D8unexpected object after cudadevrt
0x1D34788Option for ABI >= 8 is not compatible with -m32, hence ignored
0x1D347C8Incompatible options '%s' and '%s', using '%s'
0x1D34850Cannot target %s when input '%s' is SASS
0x1D348E0Option '%s' not supported for gpu architecture '%s'
0x1D34998Ignoring -dlto option because no LTO objects found
0x1D34AB8Skipping incompatible '%s' when searching for -l%s
0x1D34AF0Input file '%s' must be recompiled with toolkit >= Cuda 12.0
0x1D34B30Input file '%s' must be recompiled with toolkit >= Cuda 7.0
0x1D34B70Input file '%s' newer than toolkit (%d vs %d)
0x1D34BA0Conflicting options '%s' and '%s'
0x1D34BC8SM Arch ('%s') not found in '%s'
0x1D34BF0Library file '%s' not found in paths
0x1D34C40Could not find fatbinary in '%s'
0x1D34C68Input file '%s' abi does not match
0x1D34C90Input file '%s' size does not match target '%s'
0x1D34CF0Input file '%s' ABI version '%u' is incompatible with target ABI version '%u'
0x1D34D40Target format and SM Arch ('%u') mismatch
0x1D34DC8Could not open output file '%s'
0x1D34DE8Invalid argument for option '%s'
0x1D34E10Could not open input file '%s'
0x1D34E30Must specify output file with -o option
0x1D34E58Host linker script creation failed.
0x1D34ED0No input files specified; use option --help for more information
0x1D34F45Unknown option '%s'
0x1D34F74Internal FNLZR error '%s'
0x1D34F8ESM Arch ('%s') must be >= 20
0x1D34FABUnsupported file type '%s'
0x1D34FC6Could not write file '%s'
0x1D34FE0Could not read file '%s'
0x1D357A0Too many options file opened '%s'
0x1D357C8option '%s' has been deprecated
0x1D357E8Failed to read contents of options file
0x1D35810Could not open options file '%s'
0x1D358E8incompatible redefinition for option '%s', the last value of this option was used
0x1D39250Cannot open file '%s'
0x1D39266Bad file name '%s'
0x1D39279Unknown arch name '%s'
0x1D39290Internal error: Aborting
0x1D39330Object '%s' cannot be linked due to version mismatch...
0x1D393D8Cannot link sanitized object '%s' from version %d...
0x1D39488For kernel functions with parameter size higher than 4k bytes...
0x1D39550Output is limited to run on SM Arch '%u' only, due to missing attribute...
0x1D395B0Output is limited to run on SM Arch '%u' only due to missing attribute...
0x1D39610Section '%s' not found in object '%s'
0x1D39690Duplicate weak parameter bank for '%s' is not the same size
0x1D396D0Cannot link shared object '%s' with partially linked kernel '%s'
0x1D39760Function '%s' and function '%s' have conflicting cache preference...
0x1D39888Could not replace weak symbol '%s'
0x1D398B0Prototype doesn't match for '%s' in '%s', first defined in '%s'
0x1D398F0Size doesn't match for '%s' in '%s', first specified in '%s'
0x1D39A20Multiple definition of '%s' in '%s', first defined in '%s'
0x1D39A60Undefined reference to '%s' in '%s'
0x1D39A88Stack size for entry function '%s' cannot be statically determined
0x1D39B00Entry function '%s' uses too much data for compiler-generated constants...
0x1D39B88File uses too much global %s data (0x%llx bytes, 0x%x max)
0x1D39BC8Entry function '%s' uses too much %s data (0x%llx bytes, 0x%x max)
0x1D39C80Internal error: %s
0x1D3ADE0%s '%s' is not compatible with driver version %s
0x1D3B902exec() failed with '%s'
0x1D3B91Await() failed with '%s'
0x1D3B932fork() failed with '%s'
0x1D3B978Bailing out due to earlier errors
0x1D486CEMultiple errors:
0x1D48760elfLink: unexpected error
0x1D487F1nvvmGetErrorString
0x1D4883FelfLink internal error
0x1D48856elfLink error on cubin
0x1D4886DelfLink fatbinary error
0x1D48885elfLink memory error
0x1D4889AelfLink JIT compile failed
0x1D488B5elfLink JIT link failed
0x1D488CDelfLink nvvm error
0x1D488E0elfLink error cubin not relocatable
0x1D48908elfLink error cubin not compatible
0x1D48930elfLink cubin arch not compatible
0x1D48958elfLink linker library load error
0x1D48980elfLink error incompatible formats
0x1D489A8elfLink error finalization failed

Cross-References

Internal (nvlink wiki):

  • Error Reporting -- diag_emit (sub_467460) diagnostic pipeline with the complete 88-descriptor catalog and 40+ internal assertion messages
  • LTO Overview -- LTO pipeline context where elfLink codes 5--7 arise from NVVM IR compilation
  • libnvvm Integration -- dlsym/dlopen lifecycle for libnvvm.so loading (codes 10--11)
  • Split Compilation -- LTO split-compile path where elfLink errors propagate from parallel PTX-to-SASS workers
  • Fatbin Extraction -- Fatbinary parsing path where code 3 (fatbin error) originates
  • Cubin Loading -- Input cubin validation where codes 8 (not relocatable) and 9 (not compatible) originate
  • Architecture Compatibility -- Arch mismatch checks that produce code 10 (arch not compatible)
  • Pipeline Entry -- main() top-level error handling that translates elfLink codes to user diagnostics
  • Symbol Resolution -- Symbol merge/resolve logic that produces multiple-definition and undefined-reference errors
  • Dead Code Elimination -- DCE callgraph logic that can trigger "reference to deleted symbol" assertions

Confidence Assessment

AspectConfidenceBasis
elfLink error strings (all 14 codes + fallback)HIGHAll 14 strings found verbatim in nvlink_strings.json at addresses 0x1D4883F--0x1D489A8 plus fallback at 0x1D48760; exact match confirmed
Lookup function (sub_4BC270 at 0x4BC270)HIGHDecompiled function confirms: if (a1 <= 0xD) return off_1D489E0[a1]; return "elfLink: unexpected error";
Error table address (off_1D489E0)HIGHConfirmed in decompiled sub_4BC270; 14 const char* entries
Caller sites (5 callers)HIGH5 callers confirmed via grep of decompiled/ for sub_4BC270: sub_4297B0, sub_42AF40, sub_42A2D0, sub_427A10, and main (at two sites within main)
Code 4 / Code 7 special handlingHIGHConditional branches in decompiled sub_4297B0 and sub_42A2D0 explicitly check error_code == 4 and error_code == 7 with distinct handling paths
Secondary mapping table (dword_1D48A50)HIGHReferenced in decompiled sub_4BDAC0/sub_4BDAF0/sub_4BDB30 as a 3-entry index mapping
User-facing diagnostic catalog (90+ messages)HIGHAll strings found in nvlink_strings.json; source functions confirmed via grep of 380+ sub_467460 call sites across decompiled code
Severity level prefixesHIGHPrefix strings (warning , info , error , error* , fatal ) confirmed at 0x1D3C660--0x1D3C690
Descriptor-to-severity mappingHIGH88 descriptors in BSS range 0x2A5B530--0x2A5BB70 confirmed; severity inferred from diag_emit dispatch logic and caller behavior
Version-mismatch errors (tcgen05, sanitizer, 4k params)HIGHStrings at 0x1D39330, 0x1D393D8, 0x1D39488 confirmed; these are CUDA 13.0 additions
Symbolic names (ELFLINK_OK, etc.)LOWInferred from string content; no symbol table or debug info in the binary contains these names
Error flow architecture diagramMEDIUMCall chain reconstructed from decompiled code; covers primary paths but may miss indirect callers through function pointers
Split-compile error pathsMEDIUMThread pool error propagation confirmed via sub_43FF50 call sites; per-worker error codes routed through sub_4297B0