Repository navigation
musa: Fix for FWHT shared memory bug - #30167
Conversation
|
@yeahdongcn my stance is that I am fine with keeping the MUSA stuff in the ggml CUDA backend but only insofar it does not cause a significant maintenance burden. Having to cater to a reduced max. shared memory size is explicitly not something that I want to do. |
It does build MUSA, but only arch 31 to save time, we might consider changing that to 21 to catch these issues. To that effect this change can probably be gated to arch <= 21 only as it's likely to work on other, but @yeahdongcn can follow up on that. |
This is indeed a hardware limitation. As mentioned in #30080, all MUSA architectures except I'm thinking of updating the Docker build to drop Please see: #30203 |
Overview
#29100 broke the MUSA Docker build. The wide FWHT kernel
fwht_cuda_block<N, 256, T>keeps a full row in a static__shared__ float s[N]array, so at N=8192 it needs 33300 bytes of shared memory per block. MUSAmp_21allows at most 28672, so both the F32 and F16 instantiations fail to compile:CUDA allows 48 KB of static shared memory by default, so CUDA builds were not affected, and the PR CI does not build MUSA.
This PR compiles the 8192 case out on MUSA. On MUSA, an 8192-wide FWHT returns false from
ggml_cuda_op_fwht_impland takes the existing generic fallback, the same path widths without a kernel already use (covered by the 16384 case in #29100). Widths 1024 to 4096 keep the wide kernel on MUSA; 4096 needs about 16.9 KB. CUDA and HIP are unchanged.Additional information
Failing job: https://github.com/ggml-org/llama.cpp/actions/runs/37727833228/job/113151497895
I don't have a MUSA toolchain, so this was not compiled for MUSA locally. I checked that preprocessing removes the 8192 case only when
GGML_USE_MUSAis defined. @yeahdongcn, could you confirm the build on your side?Requirements