Repository navigation
sm_75 (Turing) support: fp16 GEMMs, INT8 GEMMs, exact attention, VAE - #1
Merged
Merged
Conversation
Build an sm_75 kernel family next to the sm_80+ code and dispatch to it by compute capability, behind a MOSSINFER_SM75 option (off: Ampere-or-newer, as before). - cmake: the sm_75 kernel family is an OBJECT library with 75-real SASS (src/kernels/sm75); the portable SIMT translation units are also compiled for 75; the probe accepts compute capability 7.5 and reports 75-real in its architecture list when MOSSINFER_SM75 is on. - kernels/basic: a cuBLASLt fp16 GEMM path for the bf16 linears (weights converted once at load or per streamed block, x converted in row chunks, output converted in place with the bias added in fp32). Text encoder, DiT globals and DiT blocks hook the weight conversion. - kernels/sm75: bf16<->fp16 conversions; exact attention on m16n8k8 (DiT joint attention head_dim 128 varlen GQA, Qwen causal head_dim 256); CUTLASS Sm75 m8n8k16 INT8 GEMMs that reuse the sm_80 epilogues; fp32 VAE conv2d, mid attention and conv weight layout. - gdn: a recurrent sm_75 kernel (the chunked prep kernel needs 85 KiB of shared memory, over Turing's 64 KiB). - tier/planner: an unset attention resolves to exact when the device has no 8-bit attention kernel (said in the reason); sm_75 DiT blocks always run op by op (the fused bf16 CUTLASS GEMMs need sm_80+ tensor cores). - vendored FlashAttention: if constexpr on the dropout/return-softmax guards, which otherwise instantiate dropout code that does not compile for the Turing MMA layout. - docs and tests: support matrix, sm_75 smoke/conversion tests, planner fallback test, and tools/sm75_check (bring-up harness with host references). Verified on an RTX 2080 8 GiB with the SFT bundle: the reference, high and turbo tiers generate through moss-infer and moss-server (protocol 2); the INT8 GEMM is bit-exact against a host int32 reference; MOSSINFER_SM75=OFF still builds and keeps the same CPU test baseline.
- docs/sm75.md: what MOSSINFER_SM75 changes (fp16 tensor-core GEMMs in place of bf16, INT8/INT4 ConvRot through the INT8 GEMMs, exact attention, fp32 VAE, the GDN recurrent kernel), how to build and package it, measured RTX 2080 performance, numerics, and what is not supported yet (8-bit attention, native INT4/NVFP4, img2img and the VAE encoder, runtime LoRAs). - README and docs: link the page; state the sm_75 tier, attention and block-path behavior; add the measured RTX 2080 numbers; document the -sm75 package and the launcher's device check for Turing machines. - CMakePresets: windows-msvc-cuda13-sm75 (configure/build/test/package presets). - packaging: append -sm75 to the CPack archive name when MOSSINFER_SM75 is on. - CI: build the Windows CUDA 13 Turing flavour as a fifth matrix job (a distinct sm75 axis value, so the four standard jobs keep their flavour and preset) and require its full and system-CUDA archives in release.yml. Verified: the preset configures with probe architectures 80-real,89-real,120,75-real and CPack name mossinfer-2.0.0-windows-cu13-x86_64-sm75; the matrix expands to five jobs; the workflow YAML parses.
The VAE encoder's convolutions and the LoRA GEMMs are still sm_80 kernels, so an sm_75 build used to fail in the middle of a request with a missing kernel image. Refuse both in Pipeline::run, before anything is planned, read or uploaded: - pipeline::NotSupportedError, mapped to MI_ERR_NOT_SUPPORTED by the C API, so moss-server reports `not_supported` in the done event. - img2img / inpainting: a request with an init image (or a mask) says the VAE encoder has no sm_75 kernel and that txt2img and the hires fix still work. - runtime LoRAs and control adapters (a LoRA): refused with the same code. - docs: sm75.md, integration.md and c-api.md say the three features are refused up front instead of failing mid-request. Verified on an RTX 2080: `moss-infer generate --init-image` and `--lora` exit 1 with the message and write nothing, moss-server answers `not_supported` in done, and txt2img still runs at the reference and high tiers. The CPU suite keeps its baseline (the same nine pre-existing Windows path failures).
The sm_80 tables are filled by a kernel built for sm_80+ only, so the Turing epilogue hit cudaErrorNoKernelImageForDevice when the high tier ran a SiTU-GLU site. Fill a second, per-device copy with a kernel built for 75; the factors are the same functions, so a lookup is bit-identical to the sm_80 table.
…ecks once per device - The hires fix and reference images also encode with the VAE encoder (sm_80+), so on Turing they are refused with not_supported up front, like img2img (the hires message used to say it worked; it failed mid-request). - conv2d_f32_sm75: a 1x1 filter staged 9 taps and read 8*C floats past the end of w. - convert_3d_inplace: rows above 65535 returned without converting (attention then read bf16 as fp16); grid.x holds 2^31 - 1 rows, and more throws. - probe: an sm_75 build's minimum is 7.5, not 7.0 (Volta passed the capability check). - chunk_rule.cu and flash_attention.cu use the cached kernels::sm75::turing(), and the GDN branch is compiled in MOSSINFER_SM75 builds only: no per-call device query on sm_80+. - The unset-attention exact fallback is keyed on the Turing device, so an sm_80+ build without the 8-bit attention kernels still refuses, as before. - test_sm75: the conversion tests skip off Turing (the kernels are 75-real), and the strided test checks the converted state, so a skipped conversion fails it.
…ets, refusals for sessions too - linear_fp32_row reads fp16 weights as well: block_modulation_/head_modulation_ are converted to fp16 at load on sm_75, so the time-embedding path threw on every step. bf16 -> fp16 is exact for in-range weights, so the result is the sm_80 one. - linear_f16: an f32 output's rows past the first chunk were written at r0 * ldo * 2 bytes; f32 rows are 4 bytes per element. - The Turing refusals are one function (detail::refuse_unsupported_on_sm75), called by Pipeline::run and by a Session for its options and every run's spec: a session's LoRAs or control adapter reached the bf16 merge and LoRA GEMMs. - flash_attention.cu: the comment names the kernel that runs on sm_75 (not the vendored FA2 traits).
This was referenced Oct 31, 2026
Merged
Merged
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.
Adds Turing (compute capability 7.5, e.g. RTX 2080 8 GiB) support behind a new
MOSSINFER_SM75option. With theoption off (the default) the engine is unchanged: Ampere/Ada/Blackwell, bf16 tensor cores.
Turing has no bf16 tensor cores, no
cp.async, no TF32, INT8 only atm8n8k16, and 64 KiB of shared memory perblock, so this is a second kernel family dispatched by compute capability, not a flag on the sm_80 one.
What is in it
cmake/modules/sm75.cmake): thesrc/kernels/sm75family as a75-realOBJECT library (always built, soOFF builds still link), the portable SIMT translation units also compiled for 75, a
windows-msvc-cuda13-sm75presetand a
-sm75CPack suffix. The probe accepts cc 7.5 and reports75-realin its architecture list.kernels/basic/gemm.cpp): the bf16 linears run on fp16 tensor cores through cuBLASLt. Weights areconverted once at load (or per streamed block), x in row chunks, accumulation stays fp32, the output is converted
back with the bias added in fp32. Hooks in the DiT (globals, resident blocks, streamed blocks) and in the text
encoder layers.
src/kernels/sm75/attention_sm75.cu): a dedicatedm16n8k8fp16 kernel with an online softmaxfor both engine shapes (DiT joint attention: head_dim 128, non-causal, varlen GQA; Qwen: head_dim 256, causal). The
vendored FlashAttention-2 sm_75 instantiation is numerically wrong (FA2 only maintains Ampere and later) and is not
used; the
if constexprfix it needed still lands here for the build.src/kernels/sm75/int8_gemm_sm75.cu): the same CUTLASS 2.x EVT kernel witharch::Sm75and them8n8k16instruction, 64x128x64 threadblock tiles (24 KiB of shared memory), reusing the sm_80 epilogues(dequantize+bias, gated residual, SiTU-GLU). Results are bit-identical to the sm_80 kernels. SiTU-GLU needs its
factor tables, which the sm_80 TU fills with an sm_80 kernel, so the Turing path fills its own copy.
Hopper and Blackwell use), so
turboworks with no new kernel.src/kernels/sm75/vae_sm75.cu): fp32 NHWC/KRSC convolutions, the mid attention and the conv weight layoutkernel;
conv2d,attention_fp32andconv_weight_krscdispatch by compute capability.kernels/gdn/chunk_rule.cu): the chunked prep kernel needs 85 KiB of shared memory; a per-token recurrentkernel fits 33 KiB and gives the same rule.
exact(the tier reason saysso); the DiT blocks always run op by op on sm_75, with a note (the fused bf16 CUTLASS GEMMs need sm_80+).
with
not_supportedbefore anything is planned or read, instead of failing mid-request on a missing kernel image.docs/sm75.md(build, what runs, performance, numerics, gaps), README/integration/c-api updates,tests/unit/test_sm75.cpp, a planner test for the exact fallback, andtools/sm75_check.cu, a bring-up harness thatruns the Turing kernels against host references.
Verification (RTX 2080 8 GiB, CUDA 13.4, MOSS-Image SFT bundle)
reference(bf16 bundle, fp16 GEMMs)high(INT8 ConvRot DiT, 5.16 GB)turbo(INT4 ConvRot DiT)sm75_check: dtype conversions bit-exact; exact attention cosine >= 0.9999985 (24 heads, 16/4 GQA, varlen, unevenlengths); fp16 GEMM cosine >= 0.999998; INT8 GEMM 0 mismatches against a host int32 reference; VAE conv/attention
cosine 1.0 (max |diff| <= 2e-5).
moss-infer generateandmoss-server(protocol 2) both produce the image, byte-identical perconfiguration; 512x512 in 20 steps takes 22 s, 1024x1024 at INT8 runs 3.8 s/step with a 6.9 GiB peak.
MOSSINFER_SM75=OFFcompiles and links unchanged and keeps the same CPU test baseline.Not included
q8_fp16,q8_fp8):highandturborunexacton sm_75.