musa: fix failures in test-backend-ops for mul_mat_id op (ggml-org#15236)

yeahdongcn · web-flow · commit 25ff6f7659f6 · 2025-08-12T10:02:51.000+08:00
* musa: fix failures in test-backend-ops for mul_mat_id op

Signed-off-by: Xiaodong Ye &lt;xiaodong.ye@mthreads.com&gt;

* Address review comments

Signed-off-by: Xiaodong Ye &lt;xiaodong.ye@mthreads.com&gt;

---------

Signed-off-by: Xiaodong Ye &lt;xiaodong.ye@mthreads.com&gt;
diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh
@@ -312,11 +312,11 @@ static bool turing_mma_available(const int cc) {
 }
 
 static bool ampere_mma_available(const int cc) {
-    return cc < GGML_CUDA_CC_OFFSET_AMD && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_AMPERE;
+    return GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_AMPERE;
 }
 
 static bool cp_async_available(const int cc) {
-    return cc < GGML_CUDA_CC_OFFSET_AMD && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_AMPERE;
+    return GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_AMPERE;
 }
 
 static constexpr __device__ int ggml_cuda_get_physical_warp_size() {

Original file line number	Diff line number	Diff line change
`@@ -312,11 +312,11 @@ static bool turing_mma_available(const int cc) {`
`312`	`312`	`}`
`313`	`313`
`314`	`314`	`static bool ampere_mma_available(const int cc) {`
`315`		`- return cc < GGML_CUDA_CC_OFFSET_AMD && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_AMPERE;`
	`315`	`+ return GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_AMPERE;`
`316`	`316`	`}`
`317`	`317`
`318`	`318`	`static bool cp_async_available(const int cc) {`
`319`		`- return cc < GGML_CUDA_CC_OFFSET_AMD && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_AMPERE;`
	`319`	`+ return GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_AMPERE;`
`320`	`320`	`}`
`321`	`321`
`322`	`322`	`static constexpr __device__ int ggml_cuda_get_physical_warp_size() {`