cuda: cap host MMQ tile size on Volta to match device kernels (#1638)
get_mmq_x_max_host() returned 128 for NVIDIA Volta when built with
GGML_CUDA_FORCE_MMQ, but get_mmq_x_max_device() caps at
MMQ_DP4A_MAX_BATCH_SIZE (= 64) for the same configuration because the
wider mmq_x device kernels are not instantiated on sm_70. The host
dispatcher selects an mmq_x in [8, mmq_x_max] and instantiates the
matching template; any pick > 64 hits NO_DEVICE_CODE at runtime and
the kernel returns without writing output, producing zero-valued
activations and incoherent generation.
Repro on Tesla V100 (sm_70) with a Q6_K Qwen3.5-MoE model:
cmake -B build -DGGML_CUDA=ON -DGGML_CUDA_FORCE_MMQ=ON \
-DGGML_CUDA_NO_PEER_COPY=ON -DCMAKE_CUDA_ARCHITECTURES=70
./llama-cli -m model.gguf -ngl 99 -p "<~200-token prompt>"
# stderr: mmq.cuh:3941: ERROR: CUDA kernel mul_mat_q has no
# device code compatible with CUDA arch 700.
Cap the host bound at MMQ_DP4A_MAX_BATCH_SIZE on Volta regardless of
GGML_CUDA_FORCE_MMQ; this matches the device template availability and
yields correct output. Measured on a single V100-SXM2:
before patch: PP unusable (NO_DEVICE_CODE), TG ~100 t/s
after patch: PP ~630 t/s, TG ~102 t/s
The change is a no-op on architectures that have int8 MMA (Turing+),
on AMD/HIP, and on builds without GGML_CUDA_FORCE_MMQ.
This commit is contained in:
parent
55d3c05bf7
commit
1163af96cf
|
|
@ -121,12 +121,15 @@ struct tile_x_sizes {
|
|||
};
|
||||
|
||||
static constexpr int get_mmq_x_max_host(const int cc) {
|
||||
// On NVIDIA Volta the device-side MMQ kernels are only instantiated up to
|
||||
// MMQ_DP4A_MAX_BATCH_SIZE when GGML_CUDA_FORCE_MMQ is defined (see
|
||||
// get_mmq_x_max_device() below). If the host dispatch picks a larger tile,
|
||||
// the kernel hits NO_DEVICE_CODE at runtime. Without GGML_CUDA_FORCE_MMQ
|
||||
// the host already capped at MMQ_DP4A_MAX_BATCH_SIZE so cuBLAS handles
|
||||
// larger batches. Either way the bound on Volta is the same, so collapse
|
||||
// the two paths and document the constraint.
|
||||
return int8_mma_available(cc) ? 128 :
|
||||
#ifdef GGML_CUDA_FORCE_MMQ
|
||||
cc >= CC_VOLTA && cc < CC_OFFSET_AMD ? 128 : 64;
|
||||
#else
|
||||
cc >= CC_VOLTA && cc < CC_OFFSET_AMD ? MMQ_DP4A_MAX_BATCH_SIZE : 64;
|
||||
#endif // GGML_CUDA_FORCE_MMQ
|
||||
}
|
||||
|
||||
static constexpr __device__ int get_mmq_x_max_device() {
|
||||
|
|
|
|||
Loading…
Reference in New Issue