From a7c90181972f550e8fc595b7c3dc369951d79910 Mon Sep 17 00:00:00 2001 From: NaLan ZeYu Date: Thu, 27 Aug 2026 15:32:32 +0800 Subject: [PATCH 1/2] Add NVIDIA Turing GPU support to FlashMMA --- cpp/neuralnet/cudaandrocmbackend.inc | 9 +++++--- cpp/neuralnet/cudaflashmma.cuh | 34 +++++++++++++++++++++++----- 2 files changed, 34 insertions(+), 9 deletions(-) diff --git a/cpp/neuralnet/cudaandrocmbackend.inc b/cpp/neuralnet/cudaandrocmbackend.inc index 03b349315..cd19a7e7d 100644 --- a/cpp/neuralnet/cudaandrocmbackend.inc +++ b/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) diff --git a/cpp/neuralnet/cudaflashmma.cuh b/cpp/neuralnet/cudaflashmma.cuh index 10a7f165f..52fa18eb6 100644 --- a/cpp/neuralnet/cudaflashmma.cuh +++ b/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(gmemSrc); + *reinterpret_cast(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; From 23302df82049357f55a1efc4ba83d1c7c08abf1a Mon Sep 17 00:00:00 2001 From: NaLan ZeYu Date: Fri, 28 Aug 2026 17:01:33 +0800 Subject: [PATCH 2/2] Fix runcudaoptests.sh fails on Turing --- cpp/runcudaopttests.sh | 29 +++++++++++++++++++++++------ 1 file changed, 23 insertions(+), 6 deletions(-) diff --git a/cpp/runcudaopttests.sh b/cpp/runcudaopttests.sh index 9646f8498..f0aaabae2 100755 --- a/cpp/runcudaopttests.sh +++ b/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 # 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" \