a PTX JIT compilation failed: causes and fix

The driver understood the PTX version but could not assemble its contents. A bad inline register reproduced CUDA error 218 at launch on a T4.

The driver accepted the PTX version, tried to assemble it for the GPU and found invalid or unsupported code.


Enum	`cudaErrorInvalidPtx`
Code	218
Runtime string	`a PTX JIT compilation failed`
Source	https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html (CUDA 13.3, checked 2026-08-29)

Measured reproduction

The repro embeds inline PTX that names a nonexistent special register, then ships PTX only so offline ptxas cannot reject it during the build. On a Tesla T4 with driver 595.84 and CUDA 12.6 (V12.6.85), CUDA_CACHE_DISABLE=1 produced this transcript in code/errors/ptx-jit-compilation-failed/evidence/run-2026-09-01.txt:

launch:  cudaErrorInvalidPtx: a PTX JIT compilation failed
sync:    cudaSuccess: no error
exit: 0

Like error 209 and 222, this is a launch-time failure. The following sync is clean because no kernel executed. Check cudaGetLastError() immediately after every launch.

Cause 1: invalid hand-written or generated PTX

This is the measured case. Extract the PTX and run ptxas directly so the one-line runtime error becomes a compiler diagnostic with a line number:

cuobjdump -ptx ./a.out > kernel.ptx
ptxas -arch=sm_75 kernel.ptx -o /dev/null

Fix the named instruction, operand type, register or target, then rebuild.

Cause 2: an instruction is unavailable on the card

Inline PTX may be valid for one architecture and illegal for another. Guard architecture-specific code and compile a real target that supports it. Features such as wgmma require an architecture-specific target, so a generic fallback is part of the program rather than a driver setting.

Cause 3: stale or unwritable JIT cache state

Retry once with CUDA_CACHE_DISABLE=1. If that changes the result, clear the user's compute cache and check its permissions. This is different from code 222: there the driver compiler is too old for the PTX version; here it understood the container and rejected what was inside.

Related errors

unsupported PTX toolchain (222): PTX version newer than the driver JIT.
no kernel image is available (209): no compatible image can be loaded.
invalid device function (98): a related image or symbol failure.

The lesson

Day 46 teaches PTX, SASS and direct assembly. Day 69 measures JIT compilation and fat-binary fallback. Setup starts at /setup/install-cuda, and the full code table is on the error hub.

Written 2026-09-02. Transcript from code/errors/ptx-jit-compilation-failed/evidence/run-2026-09-01.txt, captured on a Tesla T4, driver 595.84, CUDA 12.6. Author and reviewer: not yet assigned.