diff --git a/.github/workflows/build-cuda-ubuntu.yml b/.github/workflows/build-cuda-ubuntu.yml index 30029887c3..3dc9255be3 100644 --- a/.github/workflows/build-cuda-ubuntu.yml +++ b/.github/workflows/build-cuda-ubuntu.yml @@ -145,7 +145,7 @@ jobs: musa: runs-on: ubuntu-22.04 - container: mthreads/musa:rc4.3.0-devel-ubuntu22.04-amd64 + container: registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64 steps: - name: Clone @@ -156,7 +156,7 @@ jobs: id: depends run: | apt-get update - apt-get install -y build-essential git cmake libssl-dev jq + apt-get install -y build-essential git cmake libssl-dev jq python3-venv - name: ccache uses: ggml-org/ccache-action@v1.2.24 @@ -178,8 +178,8 @@ jobs: run: | cmake -B build -S . \ -DGGML_MUSA=ON \ - -DMUSA_ARCHITECTURES=21 - time cmake --build build --config Release -j $(nproc) + -DMUSA_ARCHITECTURES=31 + cmake --build build --config Release -j $(nproc) - name: ccache-buckets-save if: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }} diff --git a/ci/README-MUSA.md b/ci/README-MUSA.md index c5e24c5d9e..2101dfa2c0 100644 --- a/ci/README-MUSA.md +++ b/ci/README-MUSA.md @@ -21,7 +21,7 @@ docker run --privileged -it \ -v $HOME/llama.cpp/ci-cache:/ci-cache \ -v $HOME/llama.cpp/ci-results:/ci-results \ -v $PWD:/ws -w /ws \ - mthreads/musa:rc4.3.0-devel-ubuntu22.04-amd64 + registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64 ``` Inside the container, execute the following commands: diff --git a/ci/run.sh b/ci/run.sh index 0595fac5a1..fc661f9767 100755 --- a/ci/run.sh +++ b/ci/run.sh @@ -158,8 +158,8 @@ if [ ! -z ${GG_BUILD_WEBGPU} ]; then fi if [ ! -z ${GG_BUILD_MUSA} ]; then - # Use qy1 by default (MTT S80) - MUSA_ARCH=${MUSA_ARCH:-21} + # Use ph1 by default (MTT S5000) + MUSA_ARCH=${MUSA_ARCH:-31} CMAKE_EXTRA="${CMAKE_EXTRA} -DGGML_MUSA=ON -DMUSA_ARCHITECTURES=${MUSA_ARCH}" fi diff --git a/docs/build.md b/docs/build.md index 70fc17af24..bd666c145e 100644 --- a/docs/build.md +++ b/docs/build.md @@ -323,11 +323,11 @@ cmake --build build --config Release By default, all supported compute capabilities are enabled. To customize this behavior, you can specify the `MUSA_ARCHITECTURES` option in the CMake command: ```bash -cmake -B build -DGGML_MUSA=ON -DMUSA_ARCHITECTURES="21" +cmake -B build -DGGML_MUSA=ON -DMUSA_ARCHITECTURES="31" cmake --build build --config Release ``` -This configuration enables only compute capability `2.1` (MTT S80) during compilation, which can help reduce compilation time. +This configuration enables only compute capability `3.1` (MTT S5000) during compilation, which can help reduce compilation time. #### Compilation options diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 2e78ae4fac..5363dd001c 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -111,9 +111,9 @@ #define GGML_CUDA_CC_IS_QY2(cc) (cc >= GGML_CUDA_CC_QY2 && cc < GGML_CUDA_CC_PH1) #define GGML_CUDA_CC_IS_PH1(cc) (cc >= GGML_CUDA_CC_PH1) -#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070 +#if !defined(GGML_USE_HIP) && (defined(GGML_USE_MUSA) || CUDART_VERSION >= 11070) # define GGML_CUDA_USE_CUB -#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070 +#endif // !defined(GGML_USE_HIP) && (defined(GGML_USE_MUSA) || CUDART_VERSION >= 11070) // PDL host-side support (cudaLaunchKernelEx) requires CUDART >= 11.8. // However, this has been bugged in CTK < 12.3 for MSVC builds, see @@ -237,7 +237,7 @@ static const char * cu_get_error_str(CUresult err) { #define CU_CHECK(err) CUDA_CHECK_GEN(err, CUDA_SUCCESS, cu_get_error_str) #endif -#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) +#if !defined(GGML_USE_HIP) # define CUDA_SET_SHARED_MEMORY_LIMIT(kernel, nbytes) \ do { \ static bool shared_memory_limit_raised[GGML_CUDA_MAX_DEVICES] = { false }; \ @@ -252,7 +252,7 @@ static const char * cu_get_error_str(CUresult err) { do { \ GGML_UNUSED(nbytes); \ } while (0) -#endif // !(defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) +#endif // !defined(GGML_USE_HIP) #if CUDART_VERSION >= 11010 || defined(GGML_USE_MUSA) #define GGML_CUDA_ASSUME(x) __builtin_assume(x) @@ -397,7 +397,7 @@ static constexpr __device__ int ggml_cuda_get_physical_warp_size() { // Maximum number of bytes that can be copied in a single instruction. static constexpr __device__ int ggml_cuda_get_max_cpy_bytes() { -#ifdef GGML_USE_HIP +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) return 16; #else #if __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA @@ -405,7 +405,7 @@ static constexpr __device__ int ggml_cuda_get_max_cpy_bytes() { #else return 8; #endif // __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA -#endif // GGML_USE_HIP +#endif // defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) } @@ -424,10 +424,6 @@ static __device__ void no_device_code( __trap(); GGML_UNUSED(no_device_code); // suppress unused function warning - -#if defined(GGML_USE_MUSA) - __builtin_unreachable(); -#endif // defined(GGML_USE_MUSA) } #ifdef __CUDA_ARCH__ @@ -696,16 +692,11 @@ static __device__ __forceinline__ half2 ggml_cuda_hmax2(const half2 a, const hal template static __device__ __forceinline__ half2 warp_reduce_max(half2 x) { -#if !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL || defined(GGML_USE_HIP) #pragma unroll for (int offset = width/2; offset > 0; offset >>= 1) { x = ggml_cuda_hmax2(x, __shfl_xor_sync(0xffffffff, x, offset, width)); } return x; -#else - GGML_UNUSED(x); - NO_DEVICE_CODE; -#endif // !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL || defined(GGML_USE_HIP) } #if (defined(CUDART_VERSION) && CUDART_VERSION < CUDART_HMASK) || defined(GGML_USE_HIP) || \ diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh index 449a77c5b0..abb99a354a 100644 --- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh +++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh @@ -2091,26 +2091,22 @@ void ggml_cuda_flash_attn_ext_mma_f16_case(ggml_backend_cuda_context & ctx, ggml constexpr bool use_sparse_kernel = false; fattn_kernel = flash_attn_ext_f16; -#if !defined(GGML_USE_MUSA) static bool shared_memory_limit_raised[GGML_CUDA_MAX_DEVICES] = {false}; if (!shared_memory_limit_raised[id]) { CUDA_CHECK(cudaFuncSetAttribute(reinterpret_cast(fattn_kernel), cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes_shared_total)); shared_memory_limit_raised[id] = true; } -#endif // !defined(GGML_USE_MUSA) } } else { constexpr bool use_logit_softcap = true; constexpr bool use_sparse_kernel = false; fattn_kernel = flash_attn_ext_f16; -#if !defined(GGML_USE_MUSA) static bool shared_memory_limit_raised[GGML_CUDA_MAX_DEVICES] = {false}; if (!shared_memory_limit_raised[id]) { CUDA_CHECK(cudaFuncSetAttribute(reinterpret_cast(fattn_kernel), cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes_shared_total)); shared_memory_limit_raised[id] = true; } -#endif // !defined(GGML_USE_MUSA) } launch_fattn diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 36c6a93b18..e76ff3128f 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -310,13 +310,9 @@ static ggml_cuda_device_info ggml_cuda_init() { info.devices[id].smpb = prop.sharedMemPerBlock; info.devices[id].warp_size = prop.warpSize; -#ifndef GGML_USE_MUSA int supports_coop_launch = 0; CUDA_CHECK(cudaDeviceGetAttribute(&supports_coop_launch, cudaDevAttrCooperativeLaunch, physical_id)); info.devices[id].supports_cooperative_launch = !!supports_coop_launch; -#else - info.devices[id].supports_cooperative_launch = false; -#endif // !(GGML_USE_MUSA) #if defined(GGML_USE_HIP) info.devices[id].smpbo = prop.sharedMemPerBlock; @@ -337,8 +333,6 @@ static ggml_cuda_device_info ggml_cuda_init() { device_vmm ? "yes" : "no", prop.warpSize, device_vram_mib); #elif defined(GGML_USE_MUSA) - // FIXME: Ensure compatibility with varying warp sizes across different MUSA archs. - info.devices[id].warp_size = 32; info.devices[id].smpbo = prop.sharedMemPerBlockOptin; info.devices[id].cc = GGML_CUDA_CC_OFFSET_MTHREADS + prop.major * 0x100; info.devices[id].cc += prop.minor * 0x10; @@ -5581,12 +5575,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_OP_RWKV_WKV7: return true; case GGML_OP_GATED_DELTA_NET: - //TODO: enable once MUSA compiler is solved https://github.com/ggml-org/llama.cpp/pull/19504#issuecomment-4018634327 -#ifdef GGML_USE_MUSA - return false; -#else return true; -#endif // GGML_USE_MUSA case GGML_OP_DSV4_HC_COMB: return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->src[2]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32; diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index b13b34ee9c..f68b3df181 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -389,5 +389,10 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t return n_experts > 0; } + // MUSA: the MMQ kernels compute wrong values on PH1 (MTT S5000). + if (cc == GGML_CUDA_CC_PH1) { + return false; + } + return (!GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE; } diff --git a/ggml/src/ggml-cuda/ssm-scan.cu b/ggml/src/ggml-cuda/ssm-scan.cu index 40cb38dee7..e9a1043fb6 100644 --- a/ggml/src/ggml-cuda/ssm-scan.cu +++ b/ggml/src/ggml-cuda/ssm-scan.cu @@ -1,6 +1,6 @@ -#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070 +#if !defined(GGML_USE_HIP) && (defined(GGML_USE_MUSA) || CUDART_VERSION >= 11070) #define USE_CUB -#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070 +#endif // !defined(GGML_USE_HIP) && (defined(GGML_USE_MUSA) || CUDART_VERSION >= 11070) #ifdef USE_CUB #include @@ -342,7 +342,7 @@ static void ssm_scan_f32_cuda(const float * src0, const float * src1, const floa } } -#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) +#if !defined(GGML_USE_HIP) // ============================================================================ // SSD (State Space Duality) kernels for Mamba-2 prefill (n_tok > SSM_SSD_MIN_TOKENS) // @@ -821,7 +821,7 @@ void ggml_cuda_op_ssm_scan(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { GGML_ASSERT(src5->nb[2] <= (size_t)INT_MAX); GGML_ASSERT(src5->nb[3] <= (size_t)INT_MAX); -#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) +#if !defined(GGML_USE_HIP) // Mamba-2 with scalar A per head: use SSD matmul path for long sequences. // Requires NVIDIA Turing+ otherwise fallback to scan. const bool is_mamba2 = (src3->nb[1] == sizeof(float)); diff --git a/ggml/src/ggml-cuda/topk-moe.cu b/ggml/src/ggml-cuda/topk-moe.cu index dadcd601cb..ee903bf21c 100644 --- a/ggml/src/ggml-cuda/topk-moe.cu +++ b/ggml/src/ggml-cuda/topk-moe.cu @@ -98,7 +98,12 @@ __global__ void topk_moe_cuda(const float * logits, const float clamp_val, const float scale_val, const topk_moe_config config) { +#if defined(GGML_USE_MUSA) + // MUSA: every warp of a partially filled block must reach the barrier below. + const int row = MIN(blockIdx.x * blockDim.y + threadIdx.y, n_rows - 1); +#else const int row = blockIdx.x * blockDim.y + threadIdx.y; +#endif // defined(GGML_USE_MUSA) if (row >= n_rows) { return; } diff --git a/ggml/src/ggml-cuda/vendors/musa.h b/ggml/src/ggml-cuda/vendors/musa.h index 6d725c7ec1..4243caab44 100644 --- a/ggml/src/ggml-cuda/vendors/musa.h +++ b/ggml/src/ggml-cuda/vendors/musa.h @@ -44,6 +44,7 @@ #define cudaDeviceGetPCIBusId musaDeviceGetPCIBusId #define cudaDeviceProp musaDeviceProp #define cudaDeviceSynchronize musaDeviceSynchronize +#define cudaDeviceGetAttribute musaDeviceGetAttribute #define cudaError_t musaError_t #define cudaErrorMemoryAllocation musaErrorMemoryAllocation #define cudaErrorPeerAccessAlreadyEnabled musaErrorPeerAccessAlreadyEnabled @@ -114,6 +115,7 @@ #define cuMemRelease muMemRelease #define cuMemSetAccess muMemSetAccess #define cuMemUnmap muMemUnmap +#define cudaDevAttrCooperativeLaunch musaDevAttrCooperativeLaunch #define cudaFuncAttributeMaxDynamicSharedMemorySize musaFuncAttributeMaxDynamicSharedMemorySize #define cudaFuncSetAttribute musaFuncSetAttribute #define cudaMemcpy3DPeerParms musaMemcpy3DPeerParms @@ -145,6 +147,9 @@ #define cudaStreamCaptureModeRelaxed musaStreamCaptureModeRelaxed #define cudaStreamBeginCapture musaStreamBeginCapture #define cudaStreamEndCapture musaStreamEndCapture +#define cudaStreamCaptureStatus musaStreamCaptureStatus +#define cudaStreamCaptureStatusNone musaStreamCaptureStatusNone +#define cudaStreamIsCapturing musaStreamIsCapturing #define cudaOccupancyMaxActiveBlocksPerMultiprocessor musaOccupancyMaxActiveBlocksPerMultiprocessor typedef __mt_bfloat16 nv_bfloat16; diff --git a/ggml/src/ggml-cuda/wkv.cu b/ggml/src/ggml-cuda/wkv.cu index 2361112124..0bf9977605 100644 --- a/ggml/src/ggml-cuda/wkv.cu +++ b/ggml/src/ggml-cuda/wkv.cu @@ -79,9 +79,7 @@ static __global__ void rwkv_wkv7_f32(const int B, const int T, const int C, cons float state[head_size]; __shared__ float _r[head_size], _w[head_size], _k[head_size], _a[head_size], _b[head_size]; -#ifndef GGML_USE_MUSA #pragma unroll -#endif for (int i = 0; i < head_size; i++) { state[i] = s[batch_i * state_size + head_i * head_size * head_size + tid * head_size + i]; }