Add NVIDIA Turing support to FlashMMA kernel

cpp/neuralnet/cudaandrocmbackend.inc

@@ -4806,7 +4806,7 @@ struct ComputeContext {
   bool cudaDisableGraphSDPA;
   // Whether 1x1 NHWC convs use the cuBLAS GEMM path. Auto = matmul iff FP16.
   enabled_t use1x1MatmulMode;
-  // Tensor-core mma flash attention kernel. Auto = on for compute capability >= 8.0.
+  // Tensor-core mma flash attention kernel. Auto = on for compute capability >= 7.5.
   enabled_t useMmaAttentionMode;
   // Share large read-only weight buffers between the handles of NN server threads on the
   // same GPU, roughly halving weight memory for two-thread configs. Off by default because
@@ -4971,12 +4971,15 @@ struct ComputeHandle {
     // Decided before Model construction: TransformerAttentionBlock commits to the combined
     // QKV weight layout at construction when the mma attention path will handle its shapes.
     {
+      bool hasMma =
+        majorComputeCapability >= 8 ||
+        (majorComputeCapability == 7 && minorComputeCapability >= 5);
       bool wantMmaAttention =
         context->useMmaAttentionMode == enabled_t::True ||
-        (context->useMmaAttentionMode == enabled_t::Auto && majorComputeCapability >= 8);
+        (context->useMmaAttentionMode == enabled_t::Auto && hasMma);
       cudaHandles->mmaAttentionEnabled = wantMmaAttention && customCudaFlashAttentionMmaSupported();
       if(wantMmaAttention && !cudaHandles->mmaAttentionEnabled) {
-        // The probe failing on a compute-capability-8.0+ GPU means the kernel did not survive
+        // The probe failing on a compute-capability-7.5+ GPU means the kernel did not survive
         // compilation or JIT for this device (e.g. this binary lacks both SASS and JIT-able PTX
         // for it). Losing the mma path is a large silent slowdown on some GPUs, so be loud.
         if(context->useMmaAttentionMode == enabled_t::True)

cpp/neuralnet/cudaflashmma.cuh

@@ -19,7 +19,8 @@
 // whole-net effect was neutral to slightly negative everywhere except H100 at +1-4%, not
 // worth the extra code paths and config surface.
 //
-// Requires sm_80+ at runtime (mma.sync.aligned.m16n8k16 f32.f16.f16.f32). For older archs the
+// Requires sm_80+ at runtime (mma.sync.aligned.m16n8k16 f32.f16.f16.f32). sm_75 falls back
+// to a pair of mma.sync.aligned.m16n8k8 and synchronous copies. For older archs the
 // kernels compile to empty stubs, so the caller must check
 // flashAttentionMmaSupportedOnCurrentDevice() before dispatching here.
 
@@ -45,6 +46,17 @@ __device__ __forceinline__ void fmmaMma16816(float c[4], const uint32_t a[4], co
     "{%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%0,%1,%2,%3};\n"
     : "+f"(c[0]), "+f"(c[1]), "+f"(c[2]), "+f"(c[3])
     : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), "r"(b[0]), "r"(b[1]));
+#elif __CUDA_ARCH__ >= 750
+  asm volatile(
+    "mma.sync.aligned.m16n8k8.row.col.f32.f16.f16.f32 "
+    "{%0,%1,%2,%3}, {%4,%5}, {%6}, {%0,%1,%2,%3};\n"
+    : "+f"(c[0]), "+f"(c[1]), "+f"(c[2]), "+f"(c[3])
+    : "r"(a[0]), "r"(a[1]), "r"(b[0]));
+  asm volatile(
+    "mma.sync.aligned.m16n8k8.row.col.f32.f16.f16.f32 "
+    "{%0,%1,%2,%3}, {%4,%5}, {%6}, {%0,%1,%2,%3};\n"
+    : "+f"(c[0]), "+f"(c[1]), "+f"(c[2]), "+f"(c[3])
+    : "r"(a[2]), "r"(a[3]), "r"(b[1]));
 #endif
 }
 
@@ -65,7 +77,7 @@ __device__ __forceinline__ float fmmaExp2(float x) {
 __device__ __forceinline__ void fmmaLdmatrixX4(
   uint32_t& r0, uint32_t& r1, uint32_t& r2, uint32_t& r3, const half* rowPtr
 ) {
-#if __CUDA_ARCH__ >= 800
+#if __CUDA_ARCH__ >= 750
   uint32_t addr = (uint32_t)__cvta_generic_to_shared(rowPtr);
   asm volatile("ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%0,%1,%2,%3}, [%4];\n"
     : "=r"(r0), "=r"(r1), "=r"(r2), "=r"(r3) : "r"(addr));
@@ -74,7 +86,7 @@ __device__ __forceinline__ void fmmaLdmatrixX4(
 __device__ __forceinline__ void fmmaLdmatrixX4Trans(
   uint32_t& r0, uint32_t& r1, uint32_t& r2, uint32_t& r3, const half* rowPtr
 ) {
-#if __CUDA_ARCH__ >= 800
+#if __CUDA_ARCH__ >= 750
   uint32_t addr = (uint32_t)__cvta_generic_to_shared(rowPtr);
   asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%0,%1,%2,%3}, [%4];\n"
     : "=r"(r0), "=r"(r1), "=r"(r2), "=r"(r3) : "r"(addr));
@@ -83,21 +95,31 @@ __device__ __forceinline__ void fmmaLdmatrixX4Trans(
 
 // 16-byte async global->shared copy (sm_80+). With valid == false the destination is
 // zero-filled without touching global memory (src-size 0).
+// sm_75 have no cp.async, fallback to plain sync copy, the commit/wait become no-op
 __device__ __forceinline__ void fmmaCpAsync16(void* smemDst, const void* gmemSrc, bool valid) {
 #if __CUDA_ARCH__ >= 800
   uint32_t dst = (uint32_t)__cvta_generic_to_shared(smemDst);
   int srcSize = valid ? 16 : 0;
   asm volatile("cp.async.cg.shared.global [%0], [%1], 16, %2;\n" :: "r"(dst), "l"(gmemSrc), "r"(srcSize));
+#elif __CUDA_ARCH__ >= 750
+  uint4 val = make_uint4(0, 0, 0, 0);
+  if(valid)
+    val = *reinterpret_cast<const uint4*>(gmemSrc);
+  *reinterpret_cast<uint4*>(smemDst) = val;
 #endif
 }
 __device__ __forceinline__ void fmmaCpAsyncCommit() {
 #if __CUDA_ARCH__ >= 800
   asm volatile("cp.async.commit_group;\n");
+#elif __CUDA_ARCH__ >= 750
+  // no-op
 #endif
 }
 __device__ __forceinline__ void fmmaCpAsyncWaitAll() {
 #if __CUDA_ARCH__ >= 800
   asm volatile("cp.async.wait_group 0;\n");
+#elif __CUDA_ARCH__ >= 750
+  // no-op
 #endif
 }
 
@@ -115,7 +137,7 @@ flashAttentionMmaKernel(
   const half* __restrict__ mask, half* __restrict__ out,
   int seqLen, int numHeads, int numKVHeads, float scale, int qStride, int kvStride
 ) {
-#if __CUDA_ARCH__ >= 800
+#if __CUDA_ARCH__ >= 750
   constexpr int BQ = FMMA_BLOCK_Q;
   constexpr int BKV = FMMA_BLOCK_KV;
   constexpr int NTHREADS = FMMA_NWARPS * 32;
@@ -416,7 +438,7 @@ flashAttentionMmaKernel(
       }
     }
   }
-#endif  // __CUDA_ARCH__ >= 800
+#endif  // __CUDA_ARCH__ >= 750
 }
 
 #undef FMMA_NEG_BIG
@@ -425,7 +447,7 @@ flashAttentionMmaKernel(
 
 namespace flashmma {
 __global__ void flashAttentionMmaSupportProbeKernel(int* out) {
-#if __CUDA_ARCH__ >= 800
+#if __CUDA_ARCH__ >= 750
   *out = 1;
 #else
   (void)out;

cpp/runcudaopttests.sh

@@ -16,7 +16,8 @@
 # Environment/arguments:
 #   KATAGO_BIN       binary to test (default ./katago), must be a CUDA-backend build
 #   EXPECT_FUSED_FFN set to 0 for a build without the CUTLASS fused FFN kernel
-#                    (-DNO_CUTLASS_FUSED_FFN=1, or a toolchain that cannot build it):
+#                    (-DNO_CUTLASS_FUSED_FFN=1, or a toolchain that cannot build it),
+#                    or a pre-sm_80 GPU where the kernel never engages in Auto mode:
 #                    the fused FFN markers become forbidden instead of expected
 #   EXPECT_SDPA      set to 0 for a build without the cudnn graph SDPA path
 #                    (-DNO_CUDNN_SDPA=1, or cuDNN older than 8.9.3):
@@ -175,6 +176,8 @@ run_case() {
 
 # run_fail_case <name> <model> <boardsize> <overrides> <expected-error-substring>
 # The command must FAIL at startup and the output must contain the expected error message.
+# Within errpat, '|' separates alternatives of which any one suffices (same convention as
+# run_case's expectation markers).
 run_fail_case() {
   local name="$1" model="$2" boardsize="$3" overrides="$4" errpat="$5"
   matches_filter "$name" || return 0
@@ -185,10 +188,22 @@ run_fail_case() {
       -override-config "$overrides" > "$outfile" 2>&1; then
     echo "  expected an error, but the command succeeded"
     ok=0
-  elif ! grep -qF "$errpat" "$outfile"; then
-    echo "  failed, but without the expected message: '$errpat'"
-    tail -5 "$outfile" | sed 's/^/    /'
-    ok=0
+  else
+    local alt found=0
+    local oldifs="$IFS"
+    IFS='|'
+    for alt in $errpat; do
+      [ -z "$alt" ] && continue
+      if grep -qF "$alt" "$outfile"; then
+        found=1
+      fi
+    done
+    IFS="$oldifs"
+    if [ "$found" -eq 0 ]; then
+      echo "  failed, but without the expected message: '$errpat'"
+      tail -5 "$outfile" | sed 's/^/    /'
+      ok=0
+    fi
   fi
   record_result "$name" "$ok" "(expected-failure case)"
 }
@@ -383,9 +398,11 @@ run_fail_case n_transformer_nchw "$TMODEL" rectangle \
   "transformer models require NHWC, but cudaUseNHWC=false was set"
 
 if [ "$EXPECT_FUSED_FFN" = "0" ]; then
+  # Two flavors of unusable: no CUTLASS at build time, or CUTLASS built but this GPU/mode
+  # cannot run the kernel (pre-sm_80, or FP32 as forced here).
   run_fail_case n_fusedffn_fp32 "$TMODEL" rectangle \
     "requireMaxBoardSize=False,useFP16=false,cudaUseFusedFFN=true" \
-    "cudaUseFusedFFN=true but this build was compiled without CUTLASS"
+    "cudaUseFusedFFN=true but this build was compiled without CUTLASS|cudaUseFusedFFN=true but the fused FFN kernel is not usable here"
 else
   run_fail_case n_fusedffn_fp32 "$TMODEL" rectangle \
     "requireMaxBoardSize=False,useFP16=false,cudaUseFusedFFN=true" \