attempt to debug

LostRuins · LostRuins · commit 3829ae728e83 · 2025-05-31T01:02:30.000+08:00
diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh
@@ -1243,10 +1243,11 @@ static __global__ void flash_attn_ext_f16(
     }
 #if __CUDA_ARCH__ == GGML_CUDA_CC_TURING
     if (ncols1*ncols2 > 32) {
+        printf("\nCUDA_ARCH:%d, DKQ:%d, DV:%d, ncols1:%d, ncols2:%d, nwarps:%d, ntiles:%d, ne00:%d, ne01:%d\n",__CUDA_ARCH__,DKQ, DV, ncols1, ncols2, nwarps, ntiles,ne00,ne01);
         NO_DEVICE_CODE;
         return;
     }
-#endif __CUDA_ARCH__ == GGML_CUDA_CC_TURING
+#endif // __CUDA_ARCH__ == GGML_CUDA_CC_TURING
 
     static_assert(!mla || DKQ >= DV, "MLA needs DKQ >= DV");
 

Original file line number	Diff line number	Diff line change
`@@ -1243,10 +1243,11 @@ static __global__ void flash_attn_ext_f16(`
`1243`	`1243`	`}`
`1244`	`1244`	`#if __CUDA_ARCH__ == GGML_CUDA_CC_TURING`
`1245`	`1245`	`if (ncols1*ncols2 > 32) {`
	`1246`	`+ printf("\nCUDA_ARCH:%d, DKQ:%d, DV:%d, ncols1:%d, ncols2:%d, nwarps:%d, ntiles:%d, ne00:%d, ne01:%d\n",__CUDA_ARCH__,DKQ, DV, ncols1, ncols2, nwarps, ntiles,ne00,ne01);`
`1246`	`1247`	`NO_DEVICE_CODE;`
`1247`	`1248`	`return;`
`1248`	`1249`	`}`
`1249`		`-#endif __CUDA_ARCH__ == GGML_CUDA_CC_TURING`
	`1250`	`+#endif // __CUDA_ARCH__ == GGML_CUDA_CC_TURING`
`1250`	`1251`
`1251`	`1252`	`static_assert(!mla \|\| DKQ >= DV, "MLA needs DKQ >= DV");`
`1252`	`1253`