sm_75 (Turing) support: fp16 GEMMs, INT8 GEMMs, exact attention, VAE

.github/workflows/build.yml

@@ -1,7 +1,8 @@
-# Reusable build of the four release flavours (Linux/Windows x CUDA 12/13), called by ci.yml and release.yml.
-# Per flavour: configure + build with the CMakePresets.json release preset, CPU-only tests (hosted runners have no
-# GPU), install, AC1 (no torch/python) + exports + AC17 (size) + cuBLASLt-through-the-shim checks, the find_package
-# consumer, CPack archives, and the system-CUDA variant (the install without the `cuda` component: no cuBLASLt).
+# Reusable build of the four release flavours (Linux/Windows x CUDA 12/13) plus Windows CUDA 13 with SM75, called by
+# ci.yml and release.yml. Per flavour: configure + build with the CMakePresets.json release preset, CPU-only tests
+# (hosted runners have no GPU), install, AC1 (no torch/python) + exports + AC17 (size) + cuBLASLt-through-the-shim
+# checks, the find_package consumer, CPack archives, and the system-CUDA variant (the install without the `cuda`
+# component: no cuBLASLt).
 name: build
 
 on:
@@ -31,6 +32,10 @@ jobs:
         # Bump the 13.x entry to 13.4 once Jimver/cuda-toolkit lists it: the gate toolchain is 13.4.
         # bitexact_cu13 tests are GPU tests and do not run on hosted runners.
         cuda: ["12.8.1", "13.3.1"]
+        # An include entry whose os/cuda match a combination is merged into it (that is how the entries below set
+        # platform/flavour/nvcc); a new combination needs a value no combination has. The Turing flavour is that:
+        # sm75=true does not match any of the four standard jobs, so it adds a fifth one.
+        sm75: [false]
         include:
           - os: ubuntu-22.04
             platform: linux
@@ -73,6 +78,13 @@ jobs:
             cuda: "13.3.1"
             nvcc: "13.3.73"
             sub-packages: '["nvcc", "crt", "nvvm", "cudart", "cublas", "cublas_dev", "curand_dev", "thrust", "documentation"]'
+          - os: windows-2022
+            platform: windows-msvc
+            cuda: "13.3.1"
+            sm75: true
+            flavour: cuda13-sm75
+            nvcc: "13.3.73"
+            sub-packages: '["nvcc", "crt", "nvvm", "cudart", "cublas", "cublas_dev", "curand_dev", "thrust", "documentation"]'
     env:
       preset: ${{ matrix.platform }}-${{ matrix.flavour }}
       SCCACHE_GHA_ENABLED: "true"

.github/workflows/release.yml

@@ -1,6 +1,7 @@
 # Tag vX.Y.Z (a release) or vX.Y.Z-rcN, N >= 1 (a release candidate, published as a pre-release that does not become
-# "latest"): build and check all four flavours (build.yml), then publish their archives (the full package and the
-# -system-cuda variant of each) and a SHA256SUMS file as a GitHub release. X.Y.Z must equal the MI_VERSION_* of
+# "latest"): build and check all four standard flavours and the Windows CUDA 13 SM75 flavour (build.yml), then publish
+# their archives (the full package and the -system-cuda variant of each) and a SHA256SUMS file as a GitHub release.
+# X.Y.Z must equal the MI_VERSION_* of
 # include/mossinfer/mossinfer.h and an rc must come before vX.Y.Z (.github/scripts/release_version.py), the tagged
 # commit must be on main with a successful ci run (.github/scripts/release_gate.sh), and the release must not exist
 # yet (a re-run of the same workflow run may finish a partial upload).
@@ -117,6 +118,9 @@ jobs:
           done
           echo "$n full archives, each with its -system-cuda variant, all version $ARCHIVE_VERSION"
           test "$n" -gt 0
+          sm75="mossinfer-$ARCHIVE_VERSION-windows-cu13-x86_64-sm75.zip"
+          test -f "$sm75" || { echo "missing SM75 package $sm75"; exit 1; }
+          test -f "${sm75%.zip}-system-cuda.zip" || { echo "missing SM75 system-CUDA package"; exit 1; }
       - name: SHA256SUMS
         shell: bash
         working-directory: dist

CMakeLists.txt

@@ -28,6 +28,9 @@ set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
 option(MOSSINFER_WERROR "Treat warnings as errors" ON)
 option(MOSSINFER_ASAN "Build host code with AddressSanitizer" OFF)
 option(MOSSINFER_TESTS "Build unit tests" ON)
+# Turing (compute capability 7.5): builds src/kernels/sm75 for sm_75 and accepts those devices (see
+# cmake/modules/sm75.cmake). Off: the engine is Ampere-or-newer only, exactly as before.
+option(MOSSINFER_SM75 "Build the sm_75 (Turing) kernels and accept compute capability 7.5" OFF)
 
 find_package(CUDAToolkit 12.4 REQUIRED)
 # sm_120 (Blackwell GeForce) needs CUDA 12.8.
@@ -178,8 +181,11 @@ if(MOSSINFER_TESTS)
     list(FILTER MOSSINFER_TEST_SOURCES EXCLUDE REGEX ".*/tests/unit/test_(c_api|c_api_check|c_api_inspect|component_api|component_gpu|progress_model)\\.cpp$")
   endif()
   add_executable(mossinfer_tests ${MOSSINFER_TEST_SOURCES})
-  # Their reference arithmetic is compared bit for bit with the NVFP4 recipes (MOSSINFER_STRICT_FP_SOURCES) and with
-  # the samplers' kernels (the component API's host loops, design §5.3).
+  if(MOSSINFER_SM75)
+    # probe.h's minimum compute capability (test_backend_probe) and the sm_75 kernel-family smoke test.
+    target_compile_definitions(mossinfer_tests PRIVATE MOSSINFER_HAVE_SM75=1)
+  endif()
+  # Its reference arithmetic is compared bit for bit with the NVFP4 recipes (MOSSINFER_STRICT_FP_SOURCES).
   set(_mi_strict_tests ${MOSSINFER_TEST_SOURCES})
   list(FILTER _mi_strict_tests INCLUDE REGEX ".*/tests/unit/test_(nvfp4_convert|component_gpu)\\.cpp$")
   if(_mi_strict_tests)

CMakePresets.json

@@ -87,6 +87,13 @@
       "inherits": [ "release-base", "windows" ],
       "cacheVariables": { "MOSSINFER_EXPECT_CUDA_MAJOR": "13" }
     },
+    {
+      "name": "windows-msvc-cuda13-sm75",
+      "displayName": "Windows release, CUDA 13.x with Turing (sm_75)",
+      "inherits": "windows-msvc-cuda13",
+      "binaryDir": "${sourceDir}/build-${presetName}",
+      "cacheVariables": { "MOSSINFER_SM75": "ON" }
+    },
     {
       "name": "windows-msvc-cuda12",
       "displayName": "Windows release, MSVC + CUDA 12.8+ (driver >= 527)",
@@ -100,6 +107,7 @@
     { "name": "linux-cuda13", "configurePreset": "linux-cuda13" },
     { "name": "linux-cuda12", "configurePreset": "linux-cuda12" },
     { "name": "windows-msvc-cuda13", "configurePreset": "windows-msvc-cuda13" },
+    { "name": "windows-msvc-cuda13-sm75", "configurePreset": "windows-msvc-cuda13-sm75" },
     { "name": "windows-msvc-cuda12", "configurePreset": "windows-msvc-cuda12" }
   ],
   "testPresets": [
@@ -120,6 +128,11 @@
       "filter": { "exclude": { "label": "bitexact_cu13" } }
     },
     { "name": "windows-msvc-cuda13", "inherits": "test-base", "configurePreset": "windows-msvc-cuda13" },
+    {
+      "name": "windows-msvc-cuda13-sm75",
+      "inherits": "test-base",
+      "configurePreset": "windows-msvc-cuda13-sm75"
+    },
     {
       "name": "windows-msvc-cuda12",
       "inherits": "test-base",
@@ -132,6 +145,7 @@
     { "name": "linux-cuda13", "configurePreset": "linux-cuda13" },
     { "name": "linux-cuda12", "configurePreset": "linux-cuda12" },
     { "name": "windows-msvc-cuda13", "configurePreset": "windows-msvc-cuda13" },
+    { "name": "windows-msvc-cuda13-sm75", "configurePreset": "windows-msvc-cuda13-sm75" },
     { "name": "windows-msvc-cuda12", "configurePreset": "windows-msvc-cuda12" }
   ]
 }

README.md

@@ -4,7 +4,7 @@ A native C++/CUDA inference engine for the MOSS-Image text-to-image model (5B Di
 
 Models: MOSS-Image Base and SFT only (`OpenMOSS-Team/MOSS-Image-Base`, `OpenMOSS-Team/MOSS-Image-SFT`); other models are refused ([supported models](docs/integration.md#supported-models)).
 
-Version 2.1.0. Supported: Linux x86-64 and Windows x64 (MSVC), each with a CUDA 12 and a CUDA 13 build, as a full package and a `-system-cuda` package. The full support matrix is in [docs/integration.md](docs/integration.md#support-status).
+Version 2.1.0. Supported: Linux x86-64 and Windows x64 (MSVC), each with a CUDA 12 and a CUDA 13 build, as a full package and a `-system-cuda` package. Turing (compute capability 7.5, e.g. RTX 2080) is supported by builds configured with `-DMOSSINFER_SM75=ON` (the `windows-msvc-cuda13-sm75` preset packages it): fp16 tensor-core GEMMs replace the bf16 ones, INT8 and INT4 ConvRot files run the INT8 GEMMs, the exact attention and the VAE run sm_75 kernels, all three quality tiers work, and an unset attention resolves to `exact`. See [docs/sm75.md](docs/sm75.md) for what runs, the measured speed on an RTX 2080 and the remaining gaps. The full support matrix is in [docs/integration.md](docs/integration.md#support-status).
 
 The caller picks the quality tier: `reference` (bf16, exact), `high` (the loaded DiT's GEMMs with 8-bit attention) or `turbo` (a 4-bit DiT). An INT8 DiT with `high` is the everyday choice: 2.5 times the bf16 speed on an RTX 5090, in 40 % less VRAM ([performance](docs/integration.md#performance)). Upgrading from 1.x: [migration](docs/integration.md#migrating-from-1x).
 
@@ -14,6 +14,7 @@ moss-infer generate --model MOSS-Image/sft --dit moss_image_sft_int8_convrot.saf
 ```
 
 - [docs/integration.md](docs/integration.md): building the engine into an application. Start here.
+- [docs/sm75.md](docs/sm75.md): Turing (compute capability 7.5) support: build and package, what runs, measured speed on an RTX 2080, limitations.
 - [docs/c-api.md](docs/c-api.md): the C library, `include/mossinfer/mossinfer.h`.
 - [docs/server-protocol.md](docs/server-protocol.md): `moss-server`, the JSON-lines sidecar process.
 
@@ -24,6 +25,6 @@ cmake --preset linux-cuda13 && cmake --build --preset linux-cuda13
 CUDA_VISIBLE_DEVICES="" ctest --preset linux-cuda13 -L '^cpu$'
 ```
 
-The other presets are `linux-cuda12`, `windows-msvc-cuda13` and `windows-msvc-cuda12` ([docs/integration.md](docs/integration.md#getting-the-binaries)).
+The other presets are `linux-cuda12`, `windows-msvc-cuda13`, `windows-msvc-cuda13-sm75` and `windows-msvc-cuda12` ([docs/integration.md](docs/integration.md#getting-the-binaries)). The `windows-msvc-cuda13-sm75` preset builds a full CUDA 13 package with Turing (sm_75) support and names it with the `-sm75` suffix; without `MOSSINFER_SM75=ON`, the engine is Ampere-or-newer as before.
 
 Licensed under Apache-2.0 (`LICENSE`, `NOTICE`, `THIRD_PARTY_NOTICES.md`).

cmake/modules/backend.cmake

@@ -1,10 +1,22 @@
 # src/backend: the probe reports which compute capabilities this build has kernel images for, so it is told the
 # CUDA architecture list of mossinfer_core (commas: a semicolon would split the definition). This file only.
+# The sm_75 kernels are their own OBJECT library (cmake/modules/sm75.cmake) outside the core's list, so a build with
+# MOSSINFER_SM75 appends 75-real here: probe.h's minimum capability follows the same option.
 # No-op when MOSSINFER_ONLY filtered the probe out.
 set(_mi_backend_probe ${CMAKE_SOURCE_DIR}/src/backend/probe.cpp)
 if(_mi_backend_probe IN_LIST MOSSINFER_SOURCES)
+  get_target_property(_mi_core_archs mossinfer_core CUDA_ARCHITECTURES)
+  if(NOT _mi_core_archs)
+    set(_mi_core_archs "")
+  endif()
+  string(JOIN "," _mi_arch_string ${_mi_core_archs})
+  if(MOSSINFER_SM75)
+    string(APPEND _mi_arch_string ",75-real")
+  endif()
   set_property(SOURCE ${_mi_backend_probe} APPEND PROPERTY COMPILE_DEFINITIONS
-    "MOSSINFER_CUDA_ARCHITECTURES=\"$<JOIN:$<TARGET_PROPERTY:mossinfer_core,CUDA_ARCHITECTURES>,$<COMMA>>\"")
+    "MOSSINFER_CUDA_ARCHITECTURES=\"${_mi_arch_string}\"")
+  unset(_mi_arch_string)
+  unset(_mi_core_archs)
 endif()
 unset(_mi_backend_probe)
 # src/backend/cuda_driver.cpp loads libcuda and NVML at run time (dlopen; libdl before glibc 2.34).

cmake/modules/packaging.cmake

@@ -288,7 +288,12 @@ set(CPACK_PACKAGE_NAME mossinfer)
 set(CPACK_PACKAGE_VENDOR MossInfer)
 set(CPACK_PACKAGE_VERSION ${PROJECT_VERSION})
 set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "Native MOSS-Image inference engine (C API, moss-infer, moss-server)")
-set(CPACK_PACKAGE_FILE_NAME mossinfer-${PROJECT_VERSION}-${_mi_os}-cu${CUDAToolkit_VERSION_MAJOR}-${_mi_arch})
+set(_mi_sm75_suffix "")
+if(MOSSINFER_SM75)
+  set(_mi_sm75_suffix -sm75)
+endif()
+set(CPACK_PACKAGE_FILE_NAME mossinfer-${PROJECT_VERSION}-${_mi_os}-cu${CUDAToolkit_VERSION_MAJOR}-${_mi_arch}${_mi_sm75_suffix})
+unset(_mi_sm75_suffix)
 set(CPACK_ARCHIVE_THREADS 0)
 set(CPACK_PACKAGE_CHECKSUM SHA256)
 include(CPack)

cmake/modules/sm75.cmake

@@ -0,0 +1,110 @@
+# Turing (sm_75) kernels: the engine's tensor-core work needs bf16 MMA, cp.async, TF32 and up to 99 KiB of shared
+# memory, none of which exist on compute capability 7.x. src/kernels/sm75/ holds the fp16 fallbacks, built as SASS for
+# sm_75 alone in their own OBJECT library (like cmake/modules/nvfp4.cmake does for sm_120a), so the rest of
+# mossinfer_core keeps its sm_80/89/120 architecture list and none of the sm_80-only translation units see --arch=sm_75.
+#
+# The kernel family is built in every build: the host code of the sm_80 translation units references it unconditionally
+# (and only executes it on a 7.x device). MOSSINFER_SM75 (declared in CMakeLists.txt) controls what makes 7.x usable:
+#   - the probe accepts compute capability 7.5 (MOSSINFER_HAVE_SM75 on probe.cpp, and 75-real in its architecture list,
+#     cmake/modules/backend.cmake),
+#   - the portable SIMT translation units are also compiled for sm_75,
+#   - flash_attention.cu / gemm.cpp / chunk_rule.cu dispatch into the sm_75 kernels,
+#   - the sm75_check bring-up tool is built.
+# Tolerates the sources being filtered out by MOSSINFER_ONLY.
+set(_mi_sm75_sources ${MOSSINFER_SOURCES})
+list(FILTER _mi_sm75_sources INCLUDE REGEX ".*/src/kernels/sm75/[^/]*\\.cu$")
+if(_mi_sm75_sources)
+  get_target_property(_mi_core_sources mossinfer_core SOURCES)
+  list(REMOVE_ITEM _mi_core_sources ${_mi_sm75_sources})
+  set_property(TARGET mossinfer_core PROPERTY SOURCES ${_mi_core_sources})
+  unset(_mi_core_sources)
+
+  add_library(mossinfer_sm75 OBJECT ${_mi_sm75_sources})
+  set_target_properties(mossinfer_sm75 PROPERTIES
+    CUDA_ARCHITECTURES 75-real
+    POSITION_INDEPENDENT_CODE ON
+    CUDA_SEPARABLE_COMPILATION OFF)
+  target_include_directories(mossinfer_sm75 PRIVATE ${CMAKE_SOURCE_DIR}/include ${CMAKE_SOURCE_DIR}/src)
+  target_include_directories(mossinfer_sm75 SYSTEM PRIVATE
+    ${CMAKE_SOURCE_DIR}/third_party/flash-attention/csrc/flash_attn/src
+    ${CMAKE_SOURCE_DIR}/third_party/cutlass/include
+    ${CMAKE_SOURCE_DIR}/third_party/cutlass/tools/util/include)
+  target_compile_options(mossinfer_sm75 PRIVATE --expt-relaxed-constexpr -lineinfo ${MOSSINFER_CUDA_HOST_WARN})
+  target_sources(mossinfer_core PRIVATE $<TARGET_OBJECTS:mossinfer_sm75>)
+  message(STATUS "SM75: Turing kernel family built (${_mi_sm75_sources})")
+endif()
+
+if(MOSSINFER_SM75 AND _mi_sm75_sources)
+  # Portable kernels: pure SIMT translation units (no cp.async, no tensor-core asm, no CUTLASS), which run unchanged on
+  # Turing but were only compiled for the core's sm_80+ list. They are built for sm_75 too, in their own OBJECT library
+  # (the core keeps its list, so the mixed-precision kernels around them never see --arch=sm_75).
+  set(_mi_portable_sources
+      src/kernels/basic/elementwise.cu
+      src/kernels/basic/gram.cu
+      src/kernels/basic/norm.cu
+      src/kernels/basic/rope.cu
+      src/kernels/fused/qk_norm_rope.cu
+      src/kernels/gdn/chunk_rule.cu
+      src/kernels/gdn/conv_gates.cu
+      src/kernels/lora/lora_kernels.cu
+      src/kernels/quant/rotate_quantize.cu
+      src/kernels/quant/w4a8_decode.cu
+      src/kernels/quant4/convrot_quantize.cu
+      src/kernels/vae/elementwise.cu
+      src/kernels/vae/encoder_elementwise.cu
+      src/kernels/vae/group_norm.cu
+      src/models/flux2_vae/torch_group_norm.cu
+      src/models/latent_preview/latent_preview.cu
+      src/models/moss_dit/dit_kernels.cu
+      src/pipeline/accel/accel_kernels.cu
+      src/pipeline/samplers/sampler_kernels.cu
+      src/pipeline/samplers/torch_cuda_randn.cu)
+  set(_mi_portable_present "")
+  foreach(rel IN LISTS _mi_portable_sources)
+    if(${CMAKE_SOURCE_DIR}/${rel} IN_LIST MOSSINFER_SOURCES)
+      list(APPEND _mi_portable_present ${CMAKE_SOURCE_DIR}/${rel})
+    endif()
+  endforeach()
+  if(_mi_portable_present)
+    get_target_property(_mi_core_sources mossinfer_core SOURCES)
+    list(REMOVE_ITEM _mi_core_sources ${_mi_portable_present})
+    set_property(TARGET mossinfer_core PROPERTY SOURCES ${_mi_core_sources})
+    unset(_mi_core_sources)
+    set(_mi_portable_archs 75 ${CMAKE_CUDA_ARCHITECTURES})
+    list(REMOVE_DUPLICATES _mi_portable_archs)
+    add_library(mossinfer_sm75_portable OBJECT ${_mi_portable_present})
+    set_target_properties(mossinfer_sm75_portable PROPERTIES
+      CUDA_ARCHITECTURES "${_mi_portable_archs}"
+      POSITION_INDEPENDENT_CODE ON
+      CUDA_SEPARABLE_COMPILATION OFF)
+    target_include_directories(mossinfer_sm75_portable PRIVATE ${CMAKE_SOURCE_DIR}/include ${CMAKE_SOURCE_DIR}/src)
+    target_compile_options(mossinfer_sm75_portable PRIVATE --expt-relaxed-constexpr -lineinfo
+                           ${MOSSINFER_CUDA_HOST_WARN})
+    target_sources(mossinfer_core PRIVATE $<TARGET_OBJECTS:mossinfer_sm75_portable>)
+    message(STATUS "SM75: portable kernels for ${_mi_portable_archs}")
+    unset(_mi_portable_archs)
+  endif()
+  unset(_mi_portable_present)
+  unset(_mi_portable_sources)
+
+  # src/backend/probe.h: the minimum compute capability and the architecture list this build accepts. Defined on the
+  # sources that read it rather than on the whole target, so a build without MOSSINFER_SM75 keeps its old objects and
+  # adding the option only recompiles the dispatchers.
+  set_property(SOURCE ${CMAKE_SOURCE_DIR}/src/backend/probe.cpp
+                      ${CMAKE_SOURCE_DIR}/src/kernels/attention/flash_attention.cu
+                      ${CMAKE_SOURCE_DIR}/src/kernels/basic/gemm.cpp
+                      ${CMAKE_SOURCE_DIR}/src/kernels/gdn/chunk_rule.cu
+               APPEND PROPERTY COMPILE_DEFINITIONS MOSSINFER_HAVE_SM75=1)
+  message(STATUS "SM75: compute capability 7.5 accepted")
+
+  # Bring-up harness (not installed): the engine's GPU tests compile their reference kernels for the build's CUDA
+  # architectures, so a Turing device cannot run them; this tool builds its own host reference.
+  if(EXISTS ${CMAKE_SOURCE_DIR}/tools/sm75_check.cu)
+    add_executable(sm75_check ${CMAKE_SOURCE_DIR}/tools/sm75_check.cu)
+    set_target_properties(sm75_check PROPERTIES CUDA_ARCHITECTURES 75)
+    target_include_directories(sm75_check PRIVATE ${CMAKE_SOURCE_DIR}/include ${CMAKE_SOURCE_DIR}/src)
+    target_link_libraries(sm75_check PRIVATE mossinfer_core)
+    target_compile_options(sm75_check PRIVATE --expt-relaxed-constexpr -lineinfo ${MOSSINFER_CUDA_HOST_WARN})
+  endif()
+endif()
+unset(_mi_sm75_sources)