GPU Compilers and Kernel Languages

Your vector-add kernel passed every correctness check yesterday. Today the same source takes 300 milliseconds on its first call, or fails on a newer GPU with "no kernel image is available." Reading the source again won't explain either symptom. The GPU never executes CUDA C++ or Python directly.

The accelerator field guide mapped execution and memory owners across hardware. The kernel performance lab then treated memory traffic, launch geometry, occupancy, and numerical checks as evidence. Keep those habits. Here we follow the compiler artifacts between source and hardware, because each boundary owns a different class of bugs.

One addition, several representations

Start with one operation: for every valid index $i$, write $out_i = x_i + y_i$. CUDA C++ expresses the work from one thread's point of view. The launch supplies many threads.

At source level, you can inspect the bounds check and the index formula. You can't yet prove which machine instructions were emitted, how many registers they require, or whether a deployed binary contains code for the target GPU.

NVIDIA's offline compiler driver, nvcc, coordinates separate host and device compilation. For device code, the documented path reaches Parallel Thread Execution (PTX), a versioned virtual instruction set architecture (ISA). ptxas then assembles PTX into a CUDA binary, called a cubin, for a particular streaming multiprocessor target such as sm_90. A cubin is an ELF file containing encoded GPU instructions plus symbols, relocations, and resource metadata.^{[1]Reference 1NVIDIA CUDA Compiler Driver NVCChttps://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/[2]Reference 2CUDA Binary Utilitieshttps://docs.nvidia.com/cuda/cuda-binary-utilities/}

Engineers commonly call the human-readable native instruction listing SASS. SASS isn't another portable source format in this workflow. It's a disassembly of target-specific instructions stored in a cubin.

There can be an earlier compiler layer. NVVM intermediate representation (IR) is NVIDIA's GPU-oriented subset and extension of LLVM IR. Language frontends can generate NVVM IR, and libNVVM compiles compatible NVVM IR into PTX.^{[3]Reference 3NVVM IR Specificationhttps://docs.nvidia.com/cuda/nvvm-ir-spec/[4]Reference 4libNVVM APIhttps://docs.nvidia.com/cuda/libnvvm-api/} That makes this conceptual path useful:

$$\text{source} \rightarrow \text{LLVM/NVVM IR} \rightarrow \text{PTX} \rightarrow \text{cubin with native code}$$

But don't turn a conceptual compiler layer into a deployment promise. nvcc publicly documents source-to-PTX and PTX-to-cubin stages; it doesn't promise a stable, human-readable NVVM file for every build. Another frontend may use different IRs before it reaches PTX, and a non-NVIDIA backend won't target PTX at all.

You find an FFMA instruction in nvdisasm output. Which layer has supplied the strongest evidence that this native operation exists?

Answer

The target cubin. nvdisasm reads encoded machine instructions from a cubin and renders them as SASS. A source expression or PTX instruction can suggest an FFMA, but a later compiler stage may still fuse, split, remove, or replace it.

nvcc builds a package, not one universal binary

A CUDA translation unit may contain CPU host code and GPU device code. nvcc separates those paths, sends the host path to a supported C++ compiler, compiles device code for requested virtual and real architectures, and embeds device images into the host object. The container of one or more device images is a fatbinary, usually shortened to fatbin.^{[1]Reference 1NVIDIA CUDA Compiler Driver NVCChttps://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/}

The distinction between compute_90 and sm_90 records intent:

PTX offers forward compatibility, not backward compatibility. PTX generated for compute_90 may be JIT-compiled for a later target, but it can't run on an older GPU that lacks the required feature set. Cubin compatibility is narrower: NVIDIA documents compatibility within some major compute-capability families, and not across major versions.^{[1]Reference 1NVIDIA CUDA Compiler Driver NVCChttps://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/}

The runtime flow branches after packaging:

This branch explains two otherwise confusing observations. A cold process may pay compilation or lazy-loading cost before the first kernel. A package may run on the build machine yet fail on another GPU because neither a compatible cubin nor usable PTX was embedded.

Produce artifacts you can inspect

Compile the CUDA source for one native Hopper target and keep PTX for future targets. --keep asks nvcc to retain supported intermediate files, while the two --generate-code entries package both sm_90 cubin and compute_90 PTX.

Create standalone PTX and cubin files as well. Keeping both makes the boundary explicit and gives nvdisasm a cubin it can read directly.

Now inspect the package before judging performance. cuobjdump accepts a standalone cubin or a host executable, object, library, or external fatbin. nvdisasm accepts standalone cubins and adds richer control-flow and register-liveness views.^{[2]Reference 2CUDA Binary Utilitieshttps://docs.nvidia.com/cuda/cuda-binary-utilities/}

Read these outputs as a chain of evidence:

1. PTX header: Check .version, .target, address size, and kernel entry name. A mismatched .target can explain a deployment failure before launch.
2. PTX body: Look for global loads, bounds predication, arithmetic, and the store. This confirms virtual operations, not final scheduling.
3. Resource report: Record registers, static shared memory, local memory, and stack per kernel. Local memory can indicate spills, although not every local-memory use is a spill.
4. SASS listing: Confirm native loads, arithmetic, stores, predicates, and architecture-specific instructions. Compare against compiler flags and source line info.
5. Profiler trace: Measure executed instructions, memory behavior, and time on the target GPU. Static disassembly can't tell you which path dominates wall time.

⚠️ Common mistake: A shorter SASS listing isn't automatically faster. Instruction latency, issue rate, dependency chains, memory transactions, occupancy, and input-dependent control flow still decide runtime behavior.

Keep an artifact manifest

A benchmark result becomes much easier to reproduce when its compiler state travels beside it. Record at least these fields:

Hash the source, compile options, target, compiler version, and any compile-time constants into a JIT cache key. Shape-specialized kernel systems also need shape, stride, data type, alignment, and algorithmic mode when those values change generated code. Caching only by function name can silently load a valid cubin for the wrong contract.

JIT is part of request latency

Just-in-time (JIT) compilation moves work from build time to runtime. NVIDIA's runtime compiler, NVRTC, accepts CUDA C++ strings and can emit PTX, cubin, CUDA Tile IR, or link-time optimization IR depending on target and options. The CUDA Driver API can load the result, while nvJitLink can combine PTX or link-time IR and produce a linked cubin.^{[5]Reference 5NVRTChttps://docs.nvidia.com/cuda/nvrtc/}

PTX loaded at runtime is compiled by the device driver's JIT compiler. Generated cubins enter a compute cache, and a driver upgrade invalidates that cache so the new JIT compiler can regenerate code.^{[6]Reference 6CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

That runtime behavior creates an operational contract:

Use the driver's switches to test both deployment branches. The first command forces embedded PTX through JIT and disables its disk cache for the run. The second disables PTX JIT, so success requires a compatible embedded cubin.^{[7]Reference 7CUDA Environment Variableshttps://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/environment-variables.html}

Run each mode in a fresh process. Reset the environment afterward. These are packaging tests, not kernel benchmarks.

CUDA module loading is lazy by default in current documentation. First use can therefore include module or kernel loading even when a compatible cubin exists. Warm the exact kernel or use eager loading when a latency-sensitive service needs that work outside the request path.^{[7]Reference 7CUDA Environment Variableshttps://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/environment-variables.html}

A service packages sm_90 cubin plus compute_90 PTX. Its first request on a future compatible GPU is slow, but later requests are fast. What should you separate in the trace?

Answer

Separate process initialization, module loading, PTX-to-cubin JIT, cache write, first kernel launch, and steady-state execution. A fast warm kernel doesn't disprove expensive compilation or loading on the first request.

Same kernel, different unit of thought

CUDA C++ asks you to describe one thread, then organize threads into blocks. Triton asks you to describe one blocked program instance operating on vectors of offsets. Both versions below compute the same masked vector add.

The Triton kernel forms a block of 256 indices, loads two blocks under a mask, adds them, and stores the result. The compiler maps that blocked program to GPU threads and memory operations.^{[8]Reference 8Triton Programming Guide: Introductionhttps://triton-lang.org/main/programming-guide/chapter-1/introduction.html}

The arithmetic stayed fixed. Ownership changed:

Triton removes some thread-level bookkeeping. It doesn't remove the need to reason about coalescing, block size, occupancy, specialization, or numerical behavior. A blocked expression is an optimization hypothesis until generated artifacts and measurements support it.

Six authoring surfaces, six control contracts

A kernel language isn't a ranking from easy to powerful. Each surface chooses which decisions belong to author, compiler, library, and runtime.

CUDA C++

CUDA C++ uses single instruction, multiple threads (SIMT) semantics. You control thread and block geometry, address calculations, synchronization, memory spaces, and low-level intrinsics. That directness fits irregular algorithms, architecture experiments, and cases where a higher layer hides a critical decision. Its portability surface is NVIDIA GPUs, with source, PTX, and fatbins handling generations inside that ecosystem.^{[6]Reference 6CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Evidence should reach the cubin. Keep ptxas resource output, disassemble native code, profile target hardware, and compare results against a trusted implementation.

Triton block programs

Triton is a Python-embedded domain-specific language (DSL) whose program instances operate on blocks. Its compiler performs block-level data-flow analysis and can automate coalescing, vectorization, shared-memory management, synchronization, and instruction selection.^{[8]Reference 8Triton Programming Guide: Introductionhttps://triton-lang.org/main/programming-guide/chapter-1/introduction.html} Current project documentation lists NVIDIA and AMD GPU support, but backend support and generated code still vary by hardware and Triton release.^{[9]Reference 9Triton Language and Compilerhttps://github.com/triton-lang/triton}

Triton fits custom fusion and dense or structured kernels where block ownership is natural. Inspect specialized IR and target assembly when available, because one Python function can compile into many shape, type, and meta-parameter variants.

CUTLASS and CuTe

CUTLASS supplies CUDA C++ templates for high-performance linear algebra. CuTe, used inside modern CUTLASS, models hierarchical tensor layouts, tensors, copy atoms, and matrix multiply-accumulate atoms. Its C++ surface keeps detailed control over thread-data mapping and architecture features. CuTe DSL brings the same low-level concepts into a Python JIT stack while preserving explicit memory, thread, and data hierarchy.^{[10]Reference 10CUTLASS and CuTe Documentationhttps://docs.nvidia.com/cutlass/latest/overview.html}

Choose this family when matrix multiplication, convolution, attention building blocks, or unusual layouts need more architectural control than Triton exposes. Expect a steeper layout-algebra learning curve and a larger specialization space. Use CUTLASS profiler or an equivalent harness, inspect PTX and SASS, and keep compile-time policies in the benchmark manifest.

CUDA Tile

CUDA Tile changes CUDA's unit of thought from one SIMT thread to one block operating collectively on immutable multidimensional tiles. The compiler chooses the number of threads per block and maps tile operations onto registers, shared memory, tensor cores, and other hardware resources. SIMT and tile kernels can coexist in one application.^{[11]Reference 11Writing Tile Kernelshttps://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html}

Tooling snapshot, verified August 29, 2026: NVIDIA's CUDA 13.3 documentation lists CUDA Tile in Python through cuda.tile and CUDA Tile C++ in the toolkit from 13.3 onward. Tile shapes must be compile-time powers of two in the documented model. Treat earlier toolkits and different language releases as separate environments, not compatible assumptions.^{[11]Reference 11Writing Tile Kernelshttps://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html}

CUDA Tile offers source portability across NVIDIA generations by hiding thread mapping, not cross-vendor portability. Inspect Tile IR or emitted device images where supported, record compiler version, then profile on each architecture you claim to support.

Pallas

Pallas extends JAX with custom kernels for GPU and tensor processing unit (TPU) backends. Kernels use references (Refs) to memory, launch grids, block specifications (BlockSpecs), and backend-specific pipelining or hardware APIs. Current APIs embed a kernel in a larger JAX program through pl.kernel or a backend-specific kernel wrapper.^{[12]Reference 12Pallas: a JAX Kernel Languagehttps://docs.jax.dev/en/latest/pallas/}

Pallas is useful when JAX composition, differentiation, or GPU and TPU integration matters. Shared concepts don't guarantee one source or tuning configuration performs equally across backends. JAX describes Pallas as under active development, and its quickstart still warns that the API is experimental. Mosaic GPU supports Hopper and newer GPUs; the older Triton backend is deprecated in JAX 0.11 and scheduled for removal.^{[12]Reference 12Pallas: a JAX Kernel Languagehttps://docs.jax.dev/en/latest/pallas/} Pin JAX and backend versions, test interpret mode where supported, inspect lowering, and profile each target.

NKI

Neuron Kernel Interface (NKI) targets AWS Trainium, Inferentia2, Trainium2, and Trainium3 NeuronCores. Its high-level nki.language API handles tensor indexing and placement, while lower-level nki.isa exposes hardware operations for computation, data movement, dynamic control flow, and cross-core communication. The documented execution pattern moves inputs from high-bandwidth memory (HBM) into the on-chip state buffer (SBUF), computes on NeuronCore engines, then stores outputs to HBM.^{[13]Reference 13About Neuron Kernel Interfacehttps://awsdocs-neuron.readthedocs-hosted.com/en/latest/nki/get-started/about/index.html}

NKI fits workloads committed to AWS Neuron hardware that need custom operations or tighter control than framework compilation supplies. It isn't a CUDA portability layer. Keep Neuron compiler and instance type in the artifact record, use framework-level correctness comparisons, and collect Neuron profiler evidence.

The comparison is easier to use after each model has a concrete meaning:

A Triton kernel and a CuTe rewrite pass the same correctness fixture. Triton is faster on H100, while CuTe is faster on B200. What must the review artifact contain before either implementation becomes the default?

Answer

Keep both source revisions, compiler and backend versions, specialization keys, compile flags, native targets, correctness tolerances, cold and warm timings, emitted-code evidence, and per-target profiler results. Choose against the deployment GPU mix and latency objective. One architecture's timing or a PTX instruction count can't establish a portable winner.

Diagnose the first failing boundary

Compiler failures become manageable when you ask which representation last satisfied its contract.

Don't skip directly from source to profiler. A profiler identifies an executed bottleneck, but it can't explain a missing image that never loaded. Disassembly proves emitted code, but it can't prove correct answers. Correctness tests prove values for their fixtures, but they can't establish speed. The evidence layers complement one another.

Choose by the decision you must own

Start at the highest abstraction that exposes the decision blocking your kernel:

1. Use a framework or vendor library when an existing operation already meets correctness and performance needs.
2. Use Triton when custom block structure, fusion, or masking is central and compiler-managed thread mapping is acceptable.
3. Use Pallas when the kernel must compose naturally with JAX or span its GPU and TPU ecosystem.
4. Use NKI when NeuronCore memory placement or engines are the target, not an afterthought.
5. Use CUDA Tile when block-level tile semantics fit and NVIDIA's compiler should own intra-block threads.
6. Use CUTLASS/CuTe when tensor layouts, asynchronous copies, and matrix hardware atoms need explicit control.
7. Drop to CUDA C++ when per-thread behavior, synchronization, special instructions, or unsupported irregularity requires it.

That order isn't permanent. A prototype may begin in Triton, expose a compiler limitation, and move one hot kernel to CuTe or CUDA C++. Another kernel may move upward after a compiler learns the missing transformation.

Practice the full artifact path before moving on. Build the vector-add executable with an sm_90 cubin and compute_90 PTX, save its manifest and cuobjdump image inventory, then run forced-PTX and cubin-only checks in fresh processes. Capture the resource report, SASS, correctness output, and separate cold and warm timings in one review artifact. Rebuild the Triton version for at least two shapes and compare the same evidence so a language change keeps the fixture and measurement contract intact.

Mastery check