xycld · Nov 2, 2026 · Nov 1, 2026 · Nov 1, 2026 · Nov 2, 2026
@@ -1,5 +1,5 @@
-# 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
+# Reusable build of the four release flavours (Linux/Windows x CUDA 12/13), Windows CUDA 13 with SM75 and Linux
+# CUDA 12 with SM70, 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).
@@ -33,9 +33,9 @@ jobs:
         # 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.
+        # platform/flavour/nvcc); separate SM75/SM70 values add compatibility builds to the four standard jobs.
         sm75: [false]
+        sm70: [false]
         include:
           - os: ubuntu-22.04
             platform: linux
@@ -85,6 +85,17 @@ jobs:
             flavour: cuda13-sm75
             nvcc: "13.3.73"
             sub-packages: '["nvcc", "crt", "nvvm", "cudart", "cublas", "cublas_dev", "curand_dev", "thrust", "documentation"]'
+          - os: ubuntu-22.04
+            platform: linux
+            cuda: "12.8.1"
+            sm70: true
+            flavour: cuda12-sm70
+            nvcc: "12.8.93"
+            apt: >-
+              cuda-nvcc-12-8=12.8.93-1 cuda-nvvm-12-8=12.8.93-1 cuda-crt-12-8=12.8.93-1
+              cuda-cudart-12-8=12.8.90-1 cuda-cudart-dev-12-8=12.8.90-1 cuda-cccl-12-8=12.8.90-1
+              cuda-documentation-12-8=12.8.90-1 libcurand-12-8=10.3.9.90-1 libcurand-dev-12-8=10.3.9.90-1
+              libcublas-12-8=12.8.4.1-1 libcublas-dev-12-8=12.8.4.1-1
     env:
       preset: ${{ matrix.platform }}-${{ matrix.flavour }}
       SCCACHE_GHA_ENABLED: "true"

@@ -121,6 +121,9 @@ jobs:
           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; }
+          sm70="mossinfer-$ARCHIVE_VERSION-linux-cu12-x86_64-sm70.zip"
+          test -f "$sm70" || { echo "missing SM70 package $sm70"; exit 1; }
+          test -f "${sm70%.zip}-system-cuda.zip" || { echo "missing SM70 system-CUDA package"; exit 1; }
       - name: SHA256SUMS
         shell: bash
         working-directory: dist

@@ -31,6 +31,10 @@ 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)
+option(MOSSINFER_SM70 "Build the sm_70 (Volta) kernels and accept compute capability 7.0 (CUDA 12 only)" OFF)
+if(MOSSINFER_SM70 AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
+  message(FATAL_ERROR "MOSSINFER_SM70 needs CUDA 12.x: CUDA 13 removed Volta compilation and library support")
+endif()
 
 find_package(CUDAToolkit 12.4 REQUIRED)
 # sm_120 (Blackwell GeForce) needs CUDA 12.8.
@@ -181,7 +185,10 @@ 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})
-  if(MOSSINFER_SM75)
+  if(MOSSINFER_SM70)
+    target_compile_definitions(mossinfer_tests PRIVATE MOSSINFER_HAVE_SM70=1)
+  endif()
+  if(MOSSINFER_SM75 OR MOSSINFER_SM70)
     # 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()

@@ -81,6 +81,12 @@
       "inherits": [ "release-base", "linux" ],
       "cacheVariables": { "MOSSINFER_EXPECT_CUDA_MAJOR": "12" }
     },
+    {
+      "name": "linux-cuda12-sm70",
+      "displayName": "Linux release, CUDA 12.8+ with Volta (sm_70)",
+      "inherits": "linux-cuda12",
+      "cacheVariables": { "MOSSINFER_SM70": "ON" }
+    },
     {
       "name": "windows-msvc-cuda13",
       "displayName": "Windows release, MSVC + CUDA 13.x (driver >= 580)",
@@ -106,6 +112,7 @@
     { "name": "dev-sccache", "configurePreset": "dev-sccache" },
     { "name": "linux-cuda13", "configurePreset": "linux-cuda13" },
     { "name": "linux-cuda12", "configurePreset": "linux-cuda12" },
+    { "name": "linux-cuda12-sm70", "configurePreset": "linux-cuda12-sm70" },
     { "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" }
@@ -127,6 +134,11 @@
       "description": "bitexact_cu13 tests compare against CUDA 13.4 libdevice and are excluded.",
       "filter": { "exclude": { "label": "bitexact_cu13" } }
     },
+    {
+      "name": "linux-cuda12-sm70",
+      "inherits": "linux-cuda12",
+      "configurePreset": "linux-cuda12-sm70"
+    },
     { "name": "windows-msvc-cuda13", "inherits": "test-base", "configurePreset": "windows-msvc-cuda13" },
     {
       "name": "windows-msvc-cuda13-sm75",
@@ -144,6 +156,7 @@
   "packagePresets": [
     { "name": "linux-cuda13", "configurePreset": "linux-cuda13" },
     { "name": "linux-cuda12", "configurePreset": "linux-cuda12" },
+    { "name": "linux-cuda12-sm70", "configurePreset": "linux-cuda12-sm70" },
     { "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" }

@@ -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.1. 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 Ti and the remaining gaps. The full support matrix is in [docs/integration.md](docs/integration.md#support-status).
+Version 2.1.1. 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 Ti and the remaining gaps. Volta (Tesla V100, sm_70) is supported by the CUDA 12 `linux-cuda12-sm70` preset; see [docs/sm70.md](docs/sm70.md) for its kernels, validation and performance. 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).
 
@@ -28,6 +28,7 @@ Lowest VRAM against normal use at 1024×1024. `min_memory` streams the DiT and c
 ![Step time against peak GPU memory at 1024×1024 in low-VRAM and normal modes for owoinfer and official PyTorch](assets/comparison_1024x1024_low_vram.png)
 
 - [docs/integration.md](docs/integration.md): building the engine into an application. Start here.
+- [docs/sm70.md](docs/sm70.md): Volta (Tesla V100) support: CUDA 12 build, exact INT8 accumulation, validation and performance.
 - [docs/sm75.md](docs/sm75.md): Turing (compute capability 7.5) support: build and package, what runs, measured speed on an RTX 2080 Ti, 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.
@@ -39,6 +40,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`, `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.
+The other presets are `linux-cuda12`, `linux-cuda12-sm70`, `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; the `linux-cuda12-sm70` preset adds Volta with a `-sm70` suffix. The standard presets remain Ampere-or-newer.
 
 Licensed under Apache-2.0 (`LICENSE`, `NOTICE`, `THIRD_PARTY_NOTICES.md`).
@@ -1,7 +1,8 @@
 # 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.
+# MOSSINFER_SM75 appends 75-real here: probe.h's minimum capability follows the same option. MOSSINFER_SM70 appends
+# 70-real, whose SASS also counts for 7.5 (compatible_image in probe.cpp).
 # 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)
@@ -13,6 +14,9 @@ if(_mi_backend_probe IN_LIST MOSSINFER_SOURCES)
   if(MOSSINFER_SM75)
     string(APPEND _mi_arch_string ",75-real")
   endif()
+  if(MOSSINFER_SM70)
+    string(APPEND _mi_arch_string ",70-real")
+  endif()
   set_property(SOURCE ${_mi_backend_probe} APPEND PROPERTY COMPILE_DEFINITIONS
     "MOSSINFER_CUDA_ARCHITECTURES=\"${_mi_arch_string}\"")
   unset(_mi_arch_string)

@@ -292,6 +292,9 @@ set(_mi_sm75_suffix "")
 if(MOSSINFER_SM75)
   set(_mi_sm75_suffix -sm75)
 endif()
+if(MOSSINFER_SM70)
+  string(APPEND _mi_sm75_suffix -sm70)
+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)

@@ -0,0 +1,39 @@
+# Volta tensor-core kernels live outside the sm_80+ core architecture list, like sm75 and nvfp4. CUDA 12 only.
+set(_mi_sm70_sources ${MOSSINFER_SOURCES})
+list(FILTER _mi_sm70_sources INCLUDE REGEX ".*/src/kernels/sm70/[^/]*\\.cu$")
+if(_mi_sm70_sources)
+  get_target_property(_mi_core_sources mossinfer_core SOURCES)
+  list(REMOVE_ITEM _mi_core_sources ${_mi_sm70_sources})
+  set_property(TARGET mossinfer_core PROPERTY SOURCES ${_mi_core_sources})
+  unset(_mi_core_sources)
+  if(MOSSINFER_SM70)
+    add_library(mossinfer_sm70 OBJECT ${_mi_sm70_sources})
+    set_target_properties(mossinfer_sm70 PROPERTIES CUDA_ARCHITECTURES 70-real POSITION_INDEPENDENT_CODE ON
+                          CUDA_SEPARABLE_COMPILATION OFF CUDA_VISIBILITY_PRESET hidden VISIBILITY_INLINES_HIDDEN ON)
+    target_include_directories(mossinfer_sm70 PRIVATE ${CMAKE_SOURCE_DIR}/include ${CMAKE_SOURCE_DIR}/src)
+    target_include_directories(mossinfer_sm70 SYSTEM PRIVATE ${CMAKE_SOURCE_DIR}/third_party/cutlass/include
+                               ${CMAKE_SOURCE_DIR}/third_party/cutlass/tools/util/include)
+    target_compile_options(mossinfer_sm70 PRIVATE --expt-relaxed-constexpr -lineinfo ${MOSSINFER_CUDA_HOST_WARN})
+    target_sources(mossinfer_core PRIVATE $<TARGET_OBJECTS:mossinfer_sm70>)
+    set_property(SOURCE ${CMAKE_SOURCE_DIR}/src/backend/probe.cpp
+                        ${CMAKE_SOURCE_DIR}/src/kernels/sm70/sm70.cpp
+                        ${CMAKE_SOURCE_DIR}/src/kernels/attention/flash_attention.cu
+                        ${CMAKE_SOURCE_DIR}/src/kernels/vae/attention.cu
+                        ${CMAKE_SOURCE_DIR}/src/kernels/vae/conv2d.cu
+                        ${CMAKE_SOURCE_DIR}/src/kernels/basic/gemm.cpp
+                        ${CMAKE_SOURCE_DIR}/src/kernels/gdn/chunk_rule.cu
+                        ${CMAKE_SOURCE_DIR}/src/kernels/quant/int8_gemm.cu
+                        ${CMAKE_SOURCE_DIR}/src/kernels/quant/int8_gemm_residual.cu
+                        ${CMAKE_SOURCE_DIR}/src/kernels/quant/int8_gemm_situ.cu
+                 APPEND PROPERTY COMPILE_DEFINITIONS MOSSINFER_HAVE_SM70=1)
+    if(EXISTS ${CMAKE_SOURCE_DIR}/tools/sm75_check.cu)
+      add_executable(sm70_check ${CMAKE_SOURCE_DIR}/tools/sm75_check.cu)
+      set_target_properties(sm70_check PROPERTIES CUDA_ARCHITECTURES 70-real)
+      target_include_directories(sm70_check PRIVATE ${CMAKE_SOURCE_DIR}/include ${CMAKE_SOURCE_DIR}/src)
+      target_compile_options(sm70_check PRIVATE --expt-relaxed-constexpr -lineinfo ${MOSSINFER_CUDA_HOST_WARN})
+      target_link_libraries(sm70_check PRIVATE mossinfer_core)
+    endif()
+    message(STATUS "SM70: Volta kernel family built; compute capability 7.0 accepted")
+  endif()
+endif()
+unset(_mi_sm70_sources)
@@ -10,6 +10,7 @@
 #   - 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.
+# MOSSINFER_SM70 (cmake/modules/sm70.cmake) turns on the same paths, and also builds convert and vae_sm75 for sm_70.
 # 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$")
@@ -19,22 +20,40 @@ if(_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})")
+  set(_mi_sm75_present TRUE)
+  if(MOSSINFER_SM70)
+    # Conversion and VAE decode are SIMT; share their source with Volta without compiling Turing MMA for sm_70.
+    set(_mi_sm7x_sources ${_mi_sm75_sources})
+    list(FILTER _mi_sm7x_sources INCLUDE REGEX "/(convert|vae_sm75)\\.cu$")
+    if(_mi_sm7x_sources)
+      list(REMOVE_ITEM _mi_sm75_sources ${_mi_sm7x_sources})
+      add_library(mossinfer_sm7x OBJECT ${_mi_sm7x_sources})
+      set_target_properties(mossinfer_sm7x PROPERTIES CUDA_ARCHITECTURES "70-real;75-real"
+                            POSITION_INDEPENDENT_CODE ON CUDA_SEPARABLE_COMPILATION OFF)
+      target_include_directories(mossinfer_sm7x PRIVATE ${CMAKE_SOURCE_DIR}/include ${CMAKE_SOURCE_DIR}/src)
+      target_compile_options(mossinfer_sm7x PRIVATE --expt-relaxed-constexpr -lineinfo ${MOSSINFER_CUDA_HOST_WARN})
+      target_sources(mossinfer_core PRIVATE $<TARGET_OBJECTS:mossinfer_sm7x>)
+    endif()
+    unset(_mi_sm7x_sources)
+  endif()
+  if(_mi_sm75_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()
 endif()
 
-if(MOSSINFER_SM75 AND _mi_sm75_sources)
+if((MOSSINFER_SM75 OR MOSSINFER_SM70) AND _mi_sm75_present)
   # 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).
@@ -71,6 +90,9 @@ if(MOSSINFER_SM75 AND _mi_sm75_sources)
     set_property(TARGET mossinfer_core PROPERTY SOURCES ${_mi_core_sources})
     unset(_mi_core_sources)
     set(_mi_portable_archs 75 ${CMAKE_CUDA_ARCHITECTURES})
+    if(MOSSINFER_SM70)
+      list(PREPEND _mi_portable_archs 70-real)
+    endif()
     list(REMOVE_DUPLICATES _mi_portable_archs)
     add_library(mossinfer_sm75_portable OBJECT ${_mi_portable_present})
     set_target_properties(mossinfer_sm75_portable PROPERTIES
@@ -99,7 +121,7 @@ if(MOSSINFER_SM75 AND _mi_sm75_sources)
 
   # 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)
+  if(MOSSINFER_SM75 AND 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)
@@ -108,3 +130,4 @@ if(MOSSINFER_SM75 AND _mi_sm75_sources)
   endif()
 endif()
 unset(_mi_sm75_sources)
+unset(_mi_sm75_present)