feat(cuda): optional arch-tiered PTX dispatch with compute_61 fallback - #333
Open
jamesburton wants to merge 1 commit into
Open
feat(cuda): optional arch-tiered PTX dispatch with compute_61 fallback#333jamesburton wants to merge 1 commit into
jamesburton wants to merge 1 commit into
Conversation
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
marked this pull request as ready for review
June 19, 2026 16:50
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Implements the minimal first step from #332: an optional, additive arch-tiered CUDA kernel dispatch mechanism, with
compute_61portable 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>.ptxwhose 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; the0default selects the universal fallback (today's behavior).CudaTransformerModelandHybridTransformerModelpass the already-detectedCudaDevice.ComputeCapabilityMajor/Minorthrough.native/build.sh/build.ps1: opt-in higher-arch emission viaEXTRA_ARCHS(e.g.EXTRA_ARCHS="80 86"). The default invocation produces only thecompute_61PTX — the curated kernel list is empty until a genuinely arch-specific kernel exists.Why
compute_61is a virtual arch, so the shipped PTX cannot encode sm_75+ ISA (tensor-coremma.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).CudaArchVariantSelectionTests(7 cases: base fallback, matching/lower arch, highest-eligible, CC 0, malformed token, no cross-kernel leakage) — pass.--device gpuand 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.