Skip to content

feat(cuda): optional arch-tiered PTX dispatch with compute_61 fallback - #333

Open
jamesburton wants to merge 1 commit into
kkokosa:mainfrom
jamesburton:feat/cuda-arch-tiered-dispatch
Open

feat(cuda): optional arch-tiered PTX dispatch with compute_61 fallback#333
jamesburton wants to merge 1 commit into
kkokosa:mainfrom
jamesburton:feat/cuda-arch-tiered-dispatch

Conversation

@jamesburton

Copy link
Copy Markdown

Implements the minimal first step from #332: an optional, additive arch-tiered CUDA kernel dispatch mechanism, with compute_61 portable PTX retained as the universal default and fallback.

Zero behavioral change by default. No higher-arch PTX is shipped, so the loader resolves to exactly the same files as today on every GPU. This PR is the extension point only — no new/optimized kernels — so the mechanism can be reviewed in isolation. Arch-specific kernels (e.g. tensor-core prefill GEMM) would be follow-ups once the approach is agreed.

What it does

  • CudaModule.ResolveArchVariantPath / LoadForArch: for a kernel, prefer the highest <kernel>.sm_<arch>.ptx whose arch ≤ the device compute capability; otherwise use the universal <kernel>.ptx. Pure, allocation-light file-system logic.
  • CudaKernels(string ptxDir, int ccMajor = 0, int ccMinor = 0): backward-compatible overload; the 0 default selects the universal fallback (today's behavior). CudaTransformerModel and HybridTransformerModel pass the already-detected CudaDevice.ComputeCapabilityMajor/Minor through.
  • native/build.sh / build.ps1: opt-in higher-arch emission via EXTRA_ARCHS (e.g. EXTRA_ARCHS="80 86"). The default invocation produces only the compute_61 PTX — the curated kernel list is empty until a genuinely arch-specific kernel exists.

Why

compute_61 is a virtual arch, so the shipped PTX cannot encode sm_75+ ISA (tensor-core mma.sync, cp.async, bf16) regardless of the physical GPU. This adds a per-kernel opt-in path to ship higher-arch variants while preserving the single-portable-PTX deployment model and full Pascal+ back-compat. Rationale, the fatbin alternative, and the recommendation are in #332.

Validation

  • dotnet build -c Release — clean (0 warnings).
  • New CudaArchVariantSelectionTests (7 cases: base fallback, matching/lower arch, highest-eligible, CC 0, malformed token, no cross-kernel leakage) — pass.
  • End-to-end GPU smoke (RTX 3060, sm_86): a Q4_K_M model loads via --device gpu and generates correctly, confirming the arch-aware loader is behavior-preserving on the universal fallback path.

Marked draft — happy to adjust scope/naming or fold in a first arch-specific kernel if you'd prefer the mechanism land alongside a concrete user. Refs #332.

Add an optional, additive arch-tiered kernel dispatch so sm_75+ GPUs can
later load kernels built for newer ISAs, while compute_61 portable PTX
stays the universal default AND fallback. Zero behavioral change by
default: no higher-arch PTX is shipped, so the loader resolves to exactly
today's files.

- CudaModule.ResolveArchVariantPath / LoadForArch: select the highest
  "<kernel>.sm_<arch>.ptx" whose arch <= device compute capability, else
  fall back to the universal "<kernel>.ptx".
- CudaKernels(ptxDir, ccMajor=0, ccMinor=0): backward-compatible; 0 keeps
  today's behavior. CudaTransformerModel / HybridTransformerModel thread the
  already-detected device CC through.
- build.sh / build.ps1: opt-in higher-arch emission via EXTRA_ARCHS; the
  default invocation produces only the compute_61 PTX (curated list empty).
- Add CudaArchVariantSelectionTests (pure file-system logic, no GPU).

Refs #332

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@jamesburton
jamesburton marked this pull request as ready for review June 19, 2026 16:50
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant