Eval bug: Gemma 4 with-sm Tensor crashes post load on a request with NVlink Turing cards

快速结论：该错误发生在使用 -sm tensor 在 NVLink 连接的 Turing 显卡（如 Quadro RTX5000/RTX4000）上加载 Gemma 4 模型后，进行推理请求时。优先排查 NCCL 的 AllReduce 操作中的非法内存访问问题，可尝试使用 -sm layer 或 -sm row 替代。

问题场景

用户在使用 llama-cli 或 llama-server 通过 -sm tensor 参数加载 Gemma 4 31B 模型（Unsloth Q5_K_XL 量化）时，模型加载成功，但发送少量请求（通常第一个请求成功，第二个请求崩溃）后出现 CUDA 错误并硬崩溃。该问题在单次 llama-cli 交互中可能只在第二个 prompt 后触发，在 llama-server 外部 UI 下几乎每次请求都会崩溃。使用 -sm layer 或 -sm row 时未出现该问题。Qwen3.5 MoE 模型在相同参数下运行正常。

报错原文

E ggml_cuda_compute_forward: MUL_MAT failed
E CUDA error: an illegal memory access was encountered
E   current device: 0, in function ggml_cuda_compute_forward at /home/user/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:3163

E CUDA error: an illegal memory access was encountered
E   current device: 0, in function ggml_cuda_flash_attn_ext_mma_f16_case at /home/user/llama.cpp/ggml/src/ggml-cuda/template-instances/../fattn-mma-f16.cuh:1945
E   cudaFuncSetAttribute(reinterpret_cast(fattn_kernel), cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes_shared_total)

========= Invalid __global__ write of size 16 bytes
=========     at ncclDevFunc_AllReduce_Sum_bf16_RING_SIMPLE()+0x14440
=========     by thread (221,0,0) in block (0,0,0)
=========     Address 0x1bc00031d0 is out of bounds

原因分析

可能原因：当启用 -sm tensor 且使用 NVLink 连接的 Turing 显卡时，NCCL 库在 AllReduce 操作中产生了非法内存写入（Invalid __global__ write）。具体表现为 NCCL 的 ncclDevFunc_AllReduce_Sum_bf16_RING_SIMPLE 内核尝试写入超出分配内存范围的地址。这可能与 Gemma 4 模型的特定张量布局或 Turing 架构（Compute Capability 7.5）在 tensor 并行分片下的 NCCL 兼容性问题有关。

额外发现：Flash Attention 是 -sm tensor 的必需条件，如果 -fa 0 且设置了 GGML_CUDA_P2P=1，会导致段错误。该问题在 Issue 中被提及为潜在的独立 bug。

环境排查

Python / 工具版本：llama.cpp 版本 9600（commit 263cc04a5）
CUDA 版本：Release 12.4, V12.4.131
GPU：NVIDIA Quadro RTX 5000 (16GB, Turing) × 2 + Quadro RTX 4000 (8GB, Turing) × 1（RTX4000 通过 CUDA_VISIBLE_DEVICES 过滤）
NVLink：启用 NVLink Bridge
模型：Gemma 4 31B（Unsloth Q5_K_XL，测试中也验证了 Q4_K_XL 量化），Qwen3.5 MoE 模型无问题（作为参照）
启动参数：使用 -sm tensor、-fa 1、-rea on、--no-mmap、-ngl 999 等
NCCL 日志：使用 compute-sanitizer --tool memcheck 和 GGML_CUDA_DISABLE_GRAPHS=1 可捕获 NCCL 内存写入越界错误

解决步骤

临时规避方案（可优先尝试）：将 -sm tensor 替换为 -sm layer 或 -sm row。Issue 报告确认这两种模式不会触发此错误。
检查 Flash Attention 是否启用：确保 -fa 1 处于激活状态。如果 -sm tensor 搭配 -fa 0，即便设置 GGML_CUDA_P2P=1 也会导致段错误。
关闭 peer-to-peer（P2P）功能：取消设置 GGML_CUDA_P2P=1。测试表明即使不设置该环境变量，崩溃仍会出现，但去掉可能缓解部分内存错误。
尝试其他 CUDA 调试环境变量：设置 GGML_CUDA_DISABLE_GRAPHS=1 并使用 compute-sanitizer --tool memcheck 诊断更具体的 NCCL 内存越界地址。
验证其他类似模型：测试 Qwen3.5 MoE 等模型在相同参数下的表现，以确认问题是否与 Gemma 4 的特定架构有关。

验证方法

使用 -sm tensor 加载 Gemma 4 模型后，连续发送至少 2 个以上的推理请求（例如通过 llama-cli 交互式输入），如果不再出现 CUDA 错误和核心转储（core dump），且生成正常，则说明问题已解决。也可以切换回 -sm layer 或 -sm row 确认崩溃不再重现。