thor-hacks

Measured facts about the NVIDIA Jetson AGX Thor (T5000, sm_110, Blackwell, 128 GB LPDDR5X @ 273 GB/s) that are not documented anywhere obvious, plus the scripts that produce them.

Everything here was run on a real Thor. Where a number comes from someone else it is cited. Where a probe was malformed it is reported as untested, not as a capability finding.

---

Landed here from an error message?

Click for the exact strings — and where the answer is

ptxas fatal   : Value 'sm_101' is not defined for option 'gpu-name'
nvcc fatal    : Unsupported gpu architecture 'sm_101'

Thor was sm_101 before CUDA 13.0 and the name is now hard-rejected. Use sm_110. → docs/00

ptxas fatal   : PTX .version 8.8 does not support .target sm_110

PTX 9.0 is the floor for sm_110. Anything below reports every feature as unsupported, which looks like "Thor has no features" and is not. → docs/00

error: class "cudaDeviceProp" has no member "clockRate"
error: class "cudaDeviceProp" has no member "memoryClockRate"

Both were removed in CUDA 13. Use cudaDeviceGetAttribute with cudaDevAttrClockRate / cudaDevAttrMemoryClockRate. Working CUDA 13 example: probes/device-props.cu → docs/00

ptxas error   : Instruction 'mma with block scale' not supported on .target 'sm_110'

True of the warp-level form only. Thor has the tcgen05 form, which needs -arch=sm_110a. This is the error behind the widespread "sm_110 lacks FP4" belief. → docs/01

Instruction 'wgmma.mma_async ...' not supported on .target 'sm_110a'

wgmma is Hopper-only (sm_90a). Blackwell replaced warpgroup MMA with tcgen05. → docs/01

CUTLASS kernel for SM110 does not run
Using TRITON backend for Unquantized MoE
KeyError: layers.0.mlp.experts.w2_input_scale

CUTLASS needs -DCUTLASS_NVCC_ARCHS=110a and ≥ 4.2.0; the NVFP4 MoE KeyError is a vLLM version floor. → docs/02, docs/05

Silent failure with no error at all: -arch=sm_87 (Orin) still compiles cleanly on the Thor toolkit and produces a binary that will not run. → docs/00

---

The headline: "sm_110 lacks FP4" is wrong

A widespread belief — repeated in NVIDIA's own forums, and in our notes until we tested it — is that Thor cannot reach the block-scaled FP4 tensor-core path. That is true only of the warp-level mma form. Thor has the tcgen05 form, which is the one the hardware actually uses, and it needs -arch=sm_110a.

Also undocumented: warp-level 2:4 structured sparsity (mma.sp) works on plain sm_110, including s4 at m16n8k128.

ptxas 13.2, CUDA 13.2. OK = assembles, no = explicit "not supported on .target", ERR = our probe was malformed.

probe	sm_110	sm_110a	sm_90a	sm_100a	sm_120a
mma s8 m16n8k32 (dense control)	OK	OK	OK	OK	OK
mma.sp s8 m16n8k64	OK	OK	OK	OK	OK
mma.sp::ordered_metadata s8 m16n8k64	OK	OK	OK	OK	OK
mma.sp::ordered_metadata f16 m16n8k32	OK	OK	OK	OK	OK
mma.sp::ordered_metadata s4 m16n8k128	OK	OK	OK	OK	OK
tcgen05.alloc	no	OK	no	OK	no
tcgen05.alloc cta_group::2	no	OK	no	OK	no
tcgen05.ld	no	OK	no	OK	no
tcgen05.mma cta_group::1 kind::f16	no	OK	no	OK	no
tcgen05.mma.sp cta_group::1 kind::f16	no	OK	no	OK	no
tcgen05.mma cta_group::2 kind::mxf4nvf4.block_scale.scale_vec::4X	no	OK	no	OK	no
tcgen05.mma.sp cta_group::2 kind::mxf4nvf4.block_scale.scale_vec::4X	no	OK	no	OK	no
mma ... kind::mxf8f6f4.block_scale (warp-level)	ERR	ERR	ERR	ERR	ERR

Three things to take from it:

1. 2:4 sparsity is available on plain sm_110. No special arch flag.
2. tcgen05 — including the sparse and NVFP4 block-scaled forms — is sm_110a-only. tcgen05.mma.sp.cta_group::2.kind::mxf4nvf4.block_scale.scale_vec::4X is exactly the shape behind the marketed FP4-sparse peak, and it assembles.
3. sm_120a rejects the entire tcgen05 family. Thor (sm_110) and DGX Spark's GB10 (sm_121) are in different tensor families — see docs/02. Do not assume Spark results transfer.

The last row is ours failing, not the hardware's: that probe was rejected on all five targets including the known-good controls, so it is our syntax. Reported so nobody mistakes it for a capability gap.

Reproduce in about ten seconds

bash probes/sm110-sparse-tcgen05.sh

ptxas only — no GPU, no driver, no root. Our reference output is in results/sm110-probe-2026-08-17.txt.

---

Contents

doc	what is in it
00 — CUDA 11/12 → 13	sm_101 is hard-rejected, sm_87 silently isn't, PTX .version 9.0 floor, the two cudaDeviceProp fields CUDA 13 deleted
01 — Tensor cores	the full matrix, what sm_110a costs, why IMMA is the wrong tool for batch-1 decode
02 — What sm_110a unlocks	tcgen05, TMEM, 2-SM MMA, and the cluster shape that decides 95% vs 50% of peak
03 — CPU ISA	SVE is 128-bit (= NEON), no L3, and why codebook quants are gather-bound at 5.5 Gw/s
04 — Bandwidth and the roofline	262.8 GB/s measured, and why the marketed FP4 TFLOPS are irrelevant to single-stream decode
05 — Platform gotchas	multi-arch NGC containers, NVFP4 version floors, tag resolution, robotics coexistence

Probes

probe	needs	what it answers
sm110-sparse-tcgen05.sh	ptxas only	the capability matrix above
cuda-version-migration.sh	ptxas + nvcc	what breaks moving from CUDA 11/12 to 13
device-props.cu	nvcc + GPU	every property the spec pages disagree about

Coming from Orin? Start with docs/00 — Thor was sm_101 before CUDA 13.0 and that name is now hard-rejected, while sm_87 still compiles cleanly and produces a binary that will not run.

Method

Two rules everything here follows, both learned the hard way:

• Probe every target, not just yours. A probe that fails on the known-good controls is your syntax error, not a capability gap. Ours did, twice, and both are labelled ERR rather than quietly dropped.
• A kernel probe measures the kernel, not the engine. Cache-resident microbenchmarks flatter themselves. Where we quote a rate, we say which.

License

MIT. Findings are facts; do what you like with them.