DeepSeek V4.1 上 4×RTX 3090：视觉链路的两个坑和修复

把带视觉的多模态大模型（DeepSeek V4.1，748B）部署在一台 4 张消费级 RTX 3090 的机器上，做成能聊天、能发图的服务。过程中踩了两个"不报错、只是悄悄错"的坑：一个是视觉编码器里的线程竞态，数值偶尔会飘；另一个是发一张稍大的图直接崩出 cuBLAS status 13，小图却全部正常。两个都修掉了，一个一行代码，一个五行。最终效果：1554×1078 的大图完整跑通，25 秒出 601 个 token（31.2 tok/s），模型对图片内容的描述准确无误。

结论一览

项目	结果
模型	DeepSeek-V4.1-Flash-1017（748B MoE，MIT），Maya-S 量化（专家 2-3 bit，约 340GB，与原版 92.6% 同 token）
硬件	4× RTX 3090 24G（GDDR6X），Linux + CUDA 12.1，单机
软件	开源项目 Project Maya（引擎 maya-ds4，基于 llama.cpp/ggml，参考 exllamav3），不是纯自研
部署方式	setup 脚本自动下载模型并自调，起服务后网页聊天 + OpenAI 兼容 API 两用
坑 1	视觉编码器 qkv_rope 核内线程竞态，数值误差一度 ~10.5%，一行修复后降到 0.15%
坑 2	大图触发 attention 分块缓冲区越界写，异步非法访存污染 CUDA context，报 cuBLAS status 13，五行修复
最终效果	1554×1078 大图（8547 个视觉 token）全链路正常：601 tokens / 25s，31.2 tok/s，图文描述准确，日志零报错

一、部署：模型、引擎和启动方式

先说清楚这不是纯自研：跑模型用的是开源项目 Project Maya（Peasant Smith 主导、社区共同维护），它就是专门干一件事的——让这类几百 B 的 MoE 大模型跑在消费级显卡上。我们的引擎构建基于它的 v1.1.30，另外做了几处本地修改（数值对齐调试用），两个坑的修复也只在本地构建里，还没有提交上游。

模型：DeepSeek 开源的 DeepSeek-V4.1-Flash-1017（748B 参数，MoE，MIT 协议）。用的量化包是 Project Maya 自研的 Maya-S：专家层压到 2-3 bit，连同 Engram 表约 340GB，和原版模型 92.6% 的 token 选择一致。视觉部分用的是 DeepSeek 官方 vision tower（约 1GB 的 checkpoint，在第一张卡上占约 320MB 显存）。

硬件：一台 4× RTX 3090 24G 的服务器（GDDR6X 消费级卡，无 ECC），Linux + CUDA 12.1。96GB 显存装不下整个模型，靠的是 Project Maya 的专家分层思路：MoE 每个 token 只激活一小部分专家，所以把最常用的专家放 GPU、次热的放系统内存、剩下的留在 NVMe SSD 上按聊天需求随取随换，显存和内存越大越快。

引擎的技术底座（README 里都写明了，也列在 third_party 里）：张量层基于 llama.cpp / ggml，EXL3 量化路径参考了 exllamav3，视觉编码器是引擎自己写的 CUDA 核。

部署步骤其实很轻：

1. 克隆 Project Maya，跑 ./setup.sh：自动下载对应模型包，测量你的 GPU / CPU / 内存 / SSD，自动选配置（我们这台选了 256K 上下文）；
2. 启动后本地起一个端口（我们用 8001），同一个地址上同时挂网页聊天界面和标准 OpenAI 兼容 API（/v1/chat/completions）。我们这台的引擎参数：--tp 2（张量并行）+ --layer-split auto（层自动切分到多卡）+ --vision on + --dspark on（用模型自带的 DSpark drafter 做推测解码）。

发一张图的请求大概长这样：

curl -X POST http://127.0.0.1:8001/v1/chat/completions \
  -H "Authorization: Bearer $API_KEY" -H "Content-Type: application/json" \
  -d @req.json   # content 里放 {"type":"image_url","image_url":{"url":"data:image/png;base64,..."}}

模型加载大约 7~20 分钟（看磁盘热度），加载完浏览器打开端口就能聊天。

部署本身很顺，坑都出在视觉链路——也就是图片进视觉编码器、算完交给语言模型的那一段。下面两个坑都在这段里。

二、坑 1：数值"偶尔会飘"的线程竞态

症状：视觉编码器逐块和 CPU 参考实现做数值对比（把 GPU 算出的每层输出 dump 下来，和 CPU 版本算 L2 相对误差），发现 qkv + RoPE 那一步的块误差高达 10.5%（正常应该小于 0.2%）。诡异的是它不是每次都错——有的块好，有的块坏。

原因：典型的 CUDA 核内竞态。qkv_rope_kernel 里两个线程会同时碰同一块内存，谁先谁后不确定，于是结果看运气。这种 bug 用"跑一遍看对不对"永远抓不全，只能靠逐块数值对比暴露出来。

修复：一行代码（调整该处的同步/访问方式，消除共享写入的竞争）。

验证：同一张图、同一块，L2 相对误差从 1.048e-01（10.5%）降到 1.498e-03（0.15%），改善约 70 倍，且后续所有块稳定达标。

这个坑的教训放最后说。先讲更要命的第二个。

三、坑 2：大图一崩，报 cuBLAS status 13

症状：小图（视觉 token 数 N 较小）一切正常。换一张 1554×1078 的大图，视觉编码器跑完后，紧接着的线性层矩阵乘直接报错：

vision: cublasGemmEx(...): cuBLAS status 13

status 13 是 CUBLAS_STATUS_EXECUTION_FAILED。生成中止，0 token。而且报错的位置在视觉编码器之后的第一个 cuBLAS 调用（l=0 的 w1 矩阵乘）——但那里代码没问题，它只是"第一个倒霉蛋"。

排查：为什么小图没事、大图必崩？

关键数字：这张 1554×1078 的图切成 8547 个视觉 token（111×77 的 patch 网格）。而小图最多才 1500 来个。

引擎的 attention 是分块算的：一次只处理 qb 行查询，算完这一"块"再算下一块（专业点叫 piece/tiling）。对应的中间缓冲区 s_（注意力分数）和 p_（softmax 概率）也按"一块"的大小分配——桌面只铺得下一张纸。

问题出在索引上：部署构建里的代码给这两个缓冲区寻址时用的是 q0 * N（“你是第几块 × 总行数”），等于按全矩阵 N×N 的方式去摆一个只有一块大小的缓冲区。

• N ≤ 3535 时：N² 不超过缓冲区大小，而且永远只有一块（q0 恒为 0，偏移量是 0），天衣无缝。之前所有测试图都在这个安全区里，所以一直没暴露；
• N = 8547 时：qb = 1462，一共 6 块。第 2 块（q0=1462）开始，写入起点 1462×8547 = 1250 万个 float，已经超出 1249.7 万 float 的分配；最后一块要写到 7300 万 float（约 292MB）——整个缓冲区被反复越界写穿。

为什么报的是 cuBLAS 13？

GPU 的越界写不一定会当场崩。它的真实路径是：

1. 越界先踩坏相邻的已提交显存——静默破坏，毫无报错；
2. 踩到未提交/受保护的地址时，触发异步非法访存（illegal memory access）；
3. 一旦发生过非法访存，整个 CUDA context 就被"污染"了——这个进程之后做任何 CUDA 操作都会失败；
4. 于是视觉编码之后的第一个 cuBLAS 调用（w1 矩阵乘）成为第一个"替罪羊"，报出 status 13。

也就是说：错误发生在 attention 的 S/P 矩阵上，爆炸却在下一个矩阵乘，而且报错信息里没有任何"越界"字样。这就是它难查的地方。

顺带揪出的"误导报错"

排查中还有一条 LAUNCH FAIL S-gemm: invalid argument 的日志，一度以为是启动参数问题。查下来是另一处 bug 的连带效应：调试用的 dump 把尺寸写成了 N×N（实际缓冲区只有"一块"大），cudaMemcpy 读越界被驱动拦下返回 invalid argument——但 dump 函数静默吞掉了这个错误，错误状态留在 CUDA 里，被下一次操作（S-gemm 的探针）冒领报告。表象在 A 处，根因在 B 处，还隔了一步。

修复：五行

位置	改动	原因
S-gemm 的 C 指针	s_ + q0*N → s_	缓冲区是"一块"大小，从块首开始写，内核内部自己按块号偏移
softmax 调用	s_ + q0*N、p_ + q0*N → 基址 s_、p_	同上，softmax 内核按 blockIdx.x 内部寻址
PV-gemm 的 B 指针	p_ + q0*N → p_	同上
y0 dump 尺寸	N*H3 → rb_*H3	dump 范围不能超过分配
s0/p0 dump 尺寸	N*N → qb*N	同上

核心就一句话：piece 缓冲区按块复用，索引就该从块首（基址）开始，q0*N 的全局偏移是多余的。每块完整算完（S → softmax → PV → 输出映射回全局行号）才换下一块，所以去掉偏移后复用是安全的。

四、验证

修完重编（增量编译，md5 变化确认新二进制生效）、重启引擎，用同尺寸 1554×1078 的图（8547 个视觉 token、6 个块，和原图完全相同的代码路径）重发：

检查项	结果
HTTP 状态	200，自然结束（非报错中止）
生成量	601 tokens / 25 秒，31.2 tok/s，专家缓存命中 79.6%
内容正确性	模型准确描述图中三个色块：左上红矩形、中部绿椭圆、右下蓝矩形，位置颜色全对
日志	LAUNCH FAIL、cuBLAS status、illegal memory 出现次数全为 0
softmax 数学	dump 出的 1462 行概率，每行和都在 0.9997~1.0003，注意力计算数学上完全正确
系统层	内核日志无 Xid、无非法访存记录
后续健康	紧接着发小图请求，同样正常，context 无污染

各级输出的 dump 文件都留档备查了，想复查任何一层的数值都可以。

五、总结

1. 越界写是最阴险的 bug：它不在出错的地方崩，而是先静默踩坏一片内存，再让"下一个操作"替它报错。cuBLAS status 13 里没有任何"越界"字样，顺着报错位置查会一直在死胡同里打转——得反过来想"谁污染了 context"。
2. "小数据上正常"不等于正确。N ≤ 3535 这个安全阈值把 bug 完美掩护了很久，直到一张 1554×1078 的图。边界条件（分块数 > 1）必须被测试覆盖。
3. CUDA 调试的命根子是数值对比：把每层输出 dump 下来和 CPU 参考实现比，两个坑都是这么逮住的。竞态表现为"误差忽高忽低"，越界表现为"误差爆表 + 系统级非法访存"。
4. 静默吞错是大忌：dump 函数把 invalid argument 吞掉不记日志，才让排查多绕了一步。调试代码也必须有日志。

最后说结论：748B 的 DeepSeek V4.1 用 4 张消费级 3090 单机跑起来是完全可行的，靠的是 Project Maya 的专家分层（GPU→内存→SSD）+ 低比特量化；大图视觉链路修完全通，解码 31 tok/s 左右，网页和 API 两种入口都正常。如果你也在消费级显卡上跑这类带视觉的大模型、又碰上 cuBLAS status 13，大概率是视觉链路里有越界写，可以按本文的路子从"谁污染了 context"查起。