musa: fix __dp4a incorrect result

Author: Huaishun Hu

Makefile

@@ -847,7 +847,7 @@ ifdef GGML_MUSA
 	CXX := $(MUSA_PATH)/bin/clang++
 	MCC := $(CCACHE) $(MUSA_PATH)/bin/mcc
 
-	MUSAFLAGS  = -x musa -mtgpu
+	MUSAFLAGS  = -fsigned-char -x musa -mtgpu
 	MUSAFLAGS += $(foreach arch,$(subst ;, ,$(MUSA_ARCHITECTURES)),--cuda-gpu-arch=mp_$(arch))
 
 ifdef GGML_CUDA_FORCE_MMQ

ggml/src/ggml-cuda/common.cuh

@@ -353,7 +353,7 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i
 
 #else // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
 
-#if __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A
+#if __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A || defined(GGML_USE_MUSA)
     return __dp4a(a, b, c);
 #else // __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A
     const int8_t * a8 = (const int8_t *) &a;

ggml/src/ggml-cuda/mmq.cu

@@ -136,6 +136,10 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) {
         return true;
     }
 
+#if defined(GGML_USE_MUSA)
+    return true;
+#endif
+
     if (cc < GGML_CUDA_CC_DP4A) {
         return false;
     }

ggml/src/ggml-cuda/vendors/musa.h

@@ -137,32 +137,35 @@
 typedef mt_bfloat16 nv_bfloat16;
 
 /** FIXME: MUSA arch should match CUDA 11.4 */
-// #define CC_OFFSET_MT  99999 // should < CC_OFFSET_AMD
-// #define __CUDA_ARCH__ CC_OFFSET_MT
-// #define __CUDA_ARCH__ 800 // AMPERE
+// #define GGML_CUDA_CC_PASCAL     600
+// #define GGML_CUDA_CC_DP4A       610
+// #define GGML_CUDA_CC_VOLTA      700
+// #define GGML_CUDA_CC_TURING     750
+// #define GGML_CUDA_CC_AMPERE     800
 
-#define __MUSA_CC__   800
+#define __MUSA_CC__   610
+// #define __CUDA_ARCH__ __MUSA_CC__
 
 
-/** TODO: following apis not supported yet by musa sdk: ***********
+/** TODO: following apis not supported yet by musa sdk: ***********/
 
-__device__ __half hexp(const __half a) {
-    float f_a = __half2float(a);
-    float f_result = expf(f_a);
-    return __float2half(f_result);
-}
+// __device__ __half hexp(const __half a) {
+//     float f_a = __half2float(a);
+//     float f_result = expf(f_a);
+//     return __float2half(f_result);
+// }
 
-__host__ __device__ __half2 h2exp(const __half2 a) {
-    // Extract lower and upper halves
-    __half lower = __low2half(a);
-    __half upper = __high2half(a);
+// __host__ __device__ __half2 h2exp(const __half2 a) {
+//     // Extract lower and upper halves
+//     __half lower = __low2half(a);
+//     __half upper = __high2half(a);
 
-    // Compute exp for each half
-    __half exp_lower = hexp(lower);
-    __half exp_upper = hexp(upper);
+//     // Compute exp for each half
+//     __half exp_lower = hexp(lower);
+//     __half exp_upper = hexp(upper);
 
-    // Combine back into __half2
-    return __halves2half2(exp_lower, exp_upper);
-}
+//     // Combine back into __half2
+//     return __halves2half2(exp_lower, exp_upper);
+// }
 
-******************************************************************/
+/******************************************************************/

ggml/src/ggml-musa/CMakeLists.txt

@@ -49,7 +49,7 @@ if (MUSAToolkit_FOUND)
 
     set_source_files_properties(${GGML_SOURCES_MUSA} PROPERTIES LANGUAGE CXX)
     foreach(SOURCE ${GGML_SOURCES_MUSA})
-        set(COMPILE_FLAGS "-x musa -mtgpu")
+        set(COMPILE_FLAGS "-fsigned-char -x musa -mtgpu")
         foreach(ARCH ${MUSA_ARCHITECTURES})
             set(COMPILE_FLAGS "${COMPILE_FLAGS} --cuda-gpu-arch=mp_${ARCH}")
         endforeach()