# 从"能跑短的"到"能跑真长的"：xLLM v0.7 长上下文攻坚实录

> 10 天、13 个 commit，填掉三个真正卡死长上下文的坑。
> 短上下文 12 个测试格赢了 cuBLAS 基线 10 个；32K 上下文首 token 延迟从 6.9 秒压到 0.9 秒；8B 模型 × 32K prompt 从直接 OOM 变成能跑完。
> 另外还有一个"优化假设被 profiler 当场打脸"的故事，我觉得比性能数字更值得写。

---

## 一、背景：v0.6.6 之后还剩下什么

xLLM 是一个 C 核心 + Python SDK 的 LLM 推理引擎，跑 GGUF 量化模型。

我们的主力测试机是一张 **V100-32GB**。选它不是情怀，恰恰是因为它难啃：V100 是 SM 7.0，没有 tensor core，做量化推理只能靠 `__dp4a` 这条整数点积指令手搓 kernel（一条指令做 4 个 int8 乘加）。TensorRT-LLM 的 weight-only INT 解码从 Turing（SM 7.5）起步，等于把 V100 划到线外了。

xLLM 的做法是：**不把 Q4_K / Q6_K 权重反量化成 fp32 再丢给 cuBLAS，而是直接在 kernel 里用整数点积算完再出 fp32**。好处很直接——8B 模型能省掉约 27 GB 的 fp32 反量化缓存，在 32 GB 卡上这是生死线；代价是所有 kernel 都得自己写、自己调。

到 v0.6.6 收尾时，短上下文（cap ≤ 8K）的成绩单是这样的（单位 ms/token，TPOT）：

| 模型 | bs=1 | bs=2 | bs=4 | bs=8 |
|---|---:|---:|---:|---:|
| Qwen2.5-1.5B | 9.16 | 6.50 | 4.32 | 3.64 |
| Qwen3-8B | 26.76 | 20.39 | 12.48 | 9.64 |
| Llama-3-8B | 24.68 | 18.86 | 11.41 | 9.15 |

看上去"过得去"：bs=1 跟 llama.cpp 打平，bs=8 甚至略优于 cuBLAS 基线（因为 v0.6.6 的 raw cache 大幅省了内存）。

但把代码通读一遍，有三个地方越看越不对劲。

### 坑一：batch ≥ 2 反而不走整数路径

代码里躺着这么一个判断：

```
g_max_batch_hint > 1  →  dim >= 3072  →  dp4a on = 0
```

翻译过来：**batch ≥ 2 而且是 8B 模型（dim 3072 以上）→ 关掉 DP4A，退回 fp32 反量化 + cuBLAS。**

当初这么写不是没理由——batch 大了怕显存不够。问题是**怕错了对象**：raw（未反量化的量化权重）只要 5 GB，fp32 反量化缓存要 27 GB。真正会撑爆显存的是 fp32。结果为了迁就一个本就不该存在的 27 GB，把只需要 5 GB 的路径给关掉了。

这是一个典型的"双缓存共存、但门禁检查错对象"的问题。

### 坑二：首 token 延迟跟着 max_seq_len 线性涨

Qwen3-8B，一个 9 token 的 "Hello" prompt：

- cap = 512：TTFT = 2273 ms
- cap = 32768：TTFT = 6942 ms

生成 9 个 token 的 prompt，为什么要看 cap 的脸色？因为一堆 scratch buffer（`d_qkv_q/k/v`、`d_ff_gate/up/down` 等）全部按 `max_seq_len × dim` 分配，engine 初始化时统一 `cudaMemset(0)` 清一遍。cap=32K 时这些 memset 就是 34 GB，按 900 GB/s 算要 38 ms 一轮——而真正要干的活只有 9 个 token 的前向。

**分配的尺度错了：按"最大可能"分配，为"实际很小"的请求买单。**

### 坑三：8B × 长 prompt 直接 OOM

`decode_mode` 这个字段 pre-B1 就有了，但 forward 侧的 dispatch 从来不看它，一律走 per-slot 路径（fp32 反量化）。8B 模型每层 5 MB × 36 层，反量化成 fp32 就是 27 GB，V100 的 32 GB 直接被撑爆。

三个坑其实是两类问题：**坑一和坑三是同一个病——有整数路径却被强制走 fp32；坑二是设计问题——KV 和 scratch 分配得太早、太大。**

v0.7 就是这三件事的收尾。

---

## 二、Phase A：把 batch ≥ 2 的整数路径打通（3 天，4 个 commit）

### A1：先从小模型入手

dim < 3072 的模型（1.5B），batch ≥ 2 也走 DP4A。核心改动是把 `nxt_q4k_dp4a_matmul_q8_v2_f32_batched` 的 chunk-of-4 wrapper 接进 `forward_layer_batch` 的 7 个 dispatch 点（Q/K/V/attn_o/gate/up/down）。

byte-eq（逐 token 输出与基线逐字节比对）12/12 通过。1.5B 全 batch 档位 −14% ~ −31%。

### A2：真正开锁的一步——engine 初始化时预加载 raw

Qwen3-8B 是个混合精度模型（attn_v 和 ffn_down 是 Q6_K，其余是 Q4_K）。A2.1 先把 Q6_K 的 DP4A prefill dispatch 补上。

但真正解开锁的是 A2.2：**把 raw 权重的加载从 lazy 改成 engine 初始化时预加载**。

之前的逻辑是"用到才 `cudaMalloc + cudaMemcpy`"。8B 模型 36 层 × 7 个 tensor 累积下来，跑到 blk.22 前后就 OOM 了——因为 fp32 反量化已经先把显存吃掉一半。改成 engine create 时一次性把 raw 全部装好：

```
[engine] DP4A raw preload: 36 layers × 7 tensors
[engine] DP4A raw preload OK
```

配套的是 fail-fast 语义：**raw 装不下就直接 loud fail，不再走 retry 兜底**。之前那种"边跑边申请、申请不到就退让重试"的策略，只会把 OOM 推迟到运行中期，变成一个更难查的性能悬崖。

8B 的 bs=2/4 解锁，TPOT −19% / −9%。

### A3：删掉两个全局变量，顺手挖出一个潜伏 bug

A2.2 上线后，两个 file-static 全局变量（`g_max_batch_hint`、`g_engine_dim_hint`）就没人用了——每次 forward 直接看 `n_active` / `total_tokens` 就行。删的时候发现一个潜伏的 bug：

`g_max_batch_hint` 存的是**所有存活 engine 里的全局最大值**。也就是说，一个 bs=8 的 engine 起来之后，后起的 bs=1 engine 会被误判成"bs ≥ 2"，从而走根本不需要的 fp32 fallback 路径。改成 per-call 的 `n_active` 判断后修好。

Phase A 结束时，**12 个测试格里已经有 11 个优于 cuBLAS 基线**，只剩 8B bs=8 有 +26% 的回退。当时怀疑是 chunk-of-4 wrapper 在 M=8 时做了两次 kernel launch——这个假设**后来被 nsight-compute 数据推翻了**，见下文 Phase B2。

---

## 三、Phase B：Q8_0、M=8 tile，和一个 API 抽象漏洞（4 天，5 个 commit）

### B1.1：raw cache 也得是 per-engine 的

又一个跨 engine 的 bug：`g_gpu_q4k_weights[]` 是一个 file-static 的 2048 项池子，同进程内两个 engine 共用，会互相把对方的 raw 权重挤出去。改成 `XllmEngine.q4k_raw_entries` 挂在 engine 上。

byte-eq 12/12 保持——单 engine bench 走的还是同一份数据，只是存储位置变了。

### B1.2：这是 v0.6 到 v0.7 最重要的一处抽象补丁

xLLM 有两条 decode 路径：

- per-slot API（`decode_step_slot`）→ `forward_token`（走 fp32 反量化缓存）
- batched API（`decode_step_batch`）→ `forward_layer_batch`（走 DP4A raw 缓存）

问题在于：**引擎不知道调用方要走哪条**，只能靠 `max_batch` 和一堆环境变量猜。

A3 时我们试过把 DP4A 相关 env 默认打开，结果 per-slot 的 bench 用 `max_batch=1`，却被 batched env 触发了 raw 预加载，raw + fp32 双缓存直接 OOM 级联——**8B bs=1 从 25 ms/tok 崩到 4260 ms/tok**。一百多倍的性能悬崖，而且是静默发生的。

B1.2 的修法很朴素：把"调用方意图"写进配置。

```c
typedef enum {
    XLLM_DECODE_MODE_AUTO     = 0,
    XLLM_DECODE_MODE_BATCHED  = 1,
    XLLM_DECODE_MODE_PER_SLOT = 2,
} XllmDecodeMode;
```

调用方必须显式声明自己要走哪条路，行为矩阵变成一份 loud contract：

| Mode | per-slot API | batched API |
|---|---|---|
| BATCHED | abort（会 OOM） | ✅ eager preload |
| PER_SLOT | ✅ fp32 反量化 | abort（cold raw 太慢） |
| AUTO | ✅（不 preload） | ✅（lazy preload） |

**不匹配就 loud abort，绝不静默降级。**

这不是为了当前这几个 bench 好看，是为了保护未来的自己：以后再有人接错调用，会在第一秒就炸出来，而不是在某个负载下悄悄退化成百分之一的性能。

### B1.3：干掉那个被标注为"已知 bench 假象"的 212 ms

bench 里一直有个怪现象：`XLLM_BENCH_BATCH_DECODE=1` 时，bs=1 的 8B 模型 TPOT 是 **212 ms/tok**——比 llama.cpp 慢 8 倍。这个数字长期被标注成"pre-existing bench artifact"，没人动它。

B1.3 定位到根因：bench 根本没用 `XllmDecodeMode`，一直在传 AUTO。AUTO 模式下 `max_batch=1`，A3 加的 `max_batch >= 2` 保护门让 preload 不触发，于是 prefill 走 fp32 fallback，decode 想 lazy load raw 时显存已经被 fp32 占满 → cudaMalloc OOM → retry 级联 → 212 ms/tok。

修法只有一行：bench 在 `XLLM_BENCH_BATCH_DECODE=1` 时传 `BATCHED`。

- **8B bs=1 batched TPOT：212 → 25.4 ms/tok（−88%）**
- **Llama-3-8B：149 → 23.4 ms/tok（−84%）**

byte-eq 12/12 保持——Q4_K 的 DP4A 在设计上就与 SGEMV 数值一致（Q8 activation 量化会引入约 0.5% 的 RMS 噪声，但不改变 arg-max）。

### B2：一次被 profiler 当场打脸的优化

Phase A 留下的 8B bs=8 +26%。我们的假设是：chunk-of-4 wrapper 在 M=8 时要 launch 两次 kernel，改成原生 M=8 tile 应该能省一半 launch。

代码写完了：`q4k_dp4a_matmul_q8_v2_kernel<M>` 模板扩到 M=1..8，wrapper 改成 chunk-of-8。

结果：**8B bs=8 从 12.90 → 12.28 ms/tok，只快了 1.8%。**

假设错了。拿 nsight-compute 抓 M=8 kernel（数据见 `docs/perf/bs8-profiling.md`）：

| 指标 | 实测 |
|---|---|
| L1/TEX Cache 吞吐 | **85–90%** |
| L1 命中率 | 98% |
| DRAM 吞吐 | 6–7% |
| Compute SM 吞吐 | 65–71% |
| Achieved Occupancy | 42–47%（理论上限 50%，寄存器受限） |

**瓶颈是 L1 带宽，不是 kernel launch 开销。** 数据基本都在 cache 里，但 L1 → register 这一段的带宽已经打满 85%。这张表顺带把三条常见优化路线全否了：

- 上 HMMA / tensor core 没用：SM 计算只用了 65%，再加算力只会让 SM 更闲；
- 更宽的 tile / 上 cp.async 没用：DRAM 才 6%，压根不是访存延迟受限；
- 提高 occupancy 没用：已经寄存器受限，多塞 warp 也挤不出更多 L1 带宽。

**结论：接受 8B bs=8 的 +26%。** 这是 V100 的 L1 带宽天花板，不是 kernel 算法能解决的问题。真有这个需求就该上 A100 / H100。

**教训：改 kernel 之前先 profile。** B2 花半天写代码换 −1.8%，如果先花半天 profiling，直接就能得出"该 accept"的判断——省下的是半天无谓的编码，更重要的是避免了对根因的长期误判。

### B3 + B4：Q8_0 kernel，以及"别把接线往后拖"

Q8_0 是最简单的 DP4A 可用量化格式：对称 int8，每 32 个元素一个 fp16 scale，没有 min 项。B3 加了 kernel（`nxt_q8_0_dp4a_matmul_q8_v2_f32{,_batched}`），11 个单元测试全过。

B4 才是重点——**把 kernel 接进真实的 dispatch 链路**。K3 review 时明确讲过一句话：**别 defer wiring，只有单测、没接进真实链路的代码放久了必然 bit-rot。** B4 就是听劝的产物：requantize 了一个 Qwen2.5-1.5B 的 Q8_0 模型，接入 dispatch，byte-eq 保持。

Q8_0 kernel 当前实测（1.5B, bs=1..8）：**16.97 / 12.06 / 9.21 / 8.04 ms/tok**。比 Q4_K 慢 1.7–3×，原因是 8.5 bpw vs 5 bpw，而且 kernel 只用了 8/32 个 warp lane（Q8_0 block 是 34 字节，不是 4 字节对齐，用 int32 load 会 misaligned）。这个 kernel 还没做性能调优，但 correctness 有"单元测试 + 真实 dispatch"双重保护。

---

## 四、Phase C：让 TTFT 与 cap 解耦，把长 prompt 跑起来（3 天，3 个 commit）

### C.1：34.5 GB 的 memset

Qwen3-8B + 9 token prompt @ cap=32K：**TTFT 6942 ms**。用 nsys 一抓，`[CUDA memset] Total 34.5 GB`，分成两块：

1. `engine_clear_slot_kv_cuda`——每次 slot acquire 都无条件清空整个 `max_blocks_per_seq × block_stride`，8B @ cap=32K 时单次就是 1.2 GB；
2. `d_xn_q8` + 8 个 fp32 scratch buffer 的 engine init memset——一次性 4 GB。

三处修法：

**修法一：lazy KV clear。** 加 `XllmEngine.slot_kv_dirty_blocks[slot]`，release 时记录高水位 block count，下次 acquire 只清那个范围。全新 engine 的第一次 acquire 什么都不用清——KV 本来就是 0。

**修法二：`xllm_engine_prefill_batch` 默认走 batched forward。** 之前 batched prefill 要 `XLLM_USE_BATCHED_PREFILL=1` 才开，默认降级到 per-slot 的 `prefill_slot`。而 per-slot 路径没有 `dp4a_covers_all` 这个 fp32-skip 检查（这个检查只有 batched forward 有），所以即使 DP4A env 打开也会去 load fp32 反量化缓存。改成 batched 路径默认开之后，dispatch 全部走 raw。

**修法三：bench 也默认走 batched prefill**，与 `use_batched_decode` 保持一致——用户既然设了 `XLLM_BENCH_BATCH_DECODE=1`，意图就是让 batched forward 一路走到底。

结果：

| cap | 修复前 | v0.7-c1 | 变化 |
|---|---:|---:|:-:|
| 512 | 2273 ms | 931 ms | −59% |
| 2048 | — | 907 ms | 基本持平 |
| 8192 | — | 953 ms | 基本持平 |
| 32768 | **6942 ms** | **928 ms** | **−87%** |

**TTFT 在 cap 512–32K 范围内几乎拉平在 900–950 ms。** 这 900 ms 是真实工作量：36 层 × 7 个 matmul × 9 token，加上 tokenizer 和 kernel launch 开销。cap 的缩放效应彻底消失。

byte-eq 这次 11/12 精确匹配，有 1 个变了（Qwen2.5-1.5B bs=8）——因为 batched prefill 下 8 个 slot 现在共享一次 Q8 activation 量化（之前是 8 次独立量化），量化噪声的实现不同，8 个 slot 里有 1 个的 first token argmax 翻转了。个体 greedy invariance 保持，aggregate hash 变了，已 rebaseline 到 v0.7-c1。

### C.2：想测长 prompt，先修 tokenizer

TTFT 不再依赖 cap 了，但我们从来没测过**真正的长 prompt**——一直都是 9 个 token 的 "Hello"。C.2 要建 4K / 16K / 32K token 的 prompt bench。

第一步：`XLLM_BENCH_PROMPT_FILE=<path>`，用 `scripts/make_bench_prompts.sh` 生成 16 KB / 64 KB / 128 KB 的三段文本。

第二步：跑 4K prompt，**5 分钟没有任何输出**。

tokenizer 挂了。看代码：

```c
for (int iter = 0; iter < 500; iter++) {              // 500 次迭代
  for (int p = 0; p < n_byte - 1; p++) {              // N=4096 个 pair
    for (int32_t m = 0; m < tok->num_merges; m++) {   // 151K 个 merge 规则
      if (strcmp(tok->merges[m].left, left) != 0) continue;
      ...
    }
  }
}
```

4096 × 151387 × 500 = **3100 亿次 strcmp**。永远跑不完。

修法是两张开放寻址哈希表：

- `vocab_ht`：string → vocab_id（对字符串做 FNV hash）
- `merge_ht`：`(uint64_t)(left_id << 32 | right_id)` → `(priority, merged_id)`（整数 hash）

建表的成本发生在 tokenizer 加载时（一次性），encode 时变成 O(1) 查询。**32K 文本从"永远不完成"变成约 2 秒。** byte-eq 保持——同样的数学，只是换了数据结构。

顺带又挖出两个 bug：

- 跑 16K 时 segfault。`alloca(total_tokens × max_blocks_per_seq × 4)` = 14242 × 2048 × 4 = **116 MB 栈分配**，而默认栈只有 8 MB。改成 malloc/free。
- `byte_tokens[4096]` 是个固定栈数组，超过 4096 字节的输入直接被截断。改成 `malloc(text_len + 16)`。

长 prompt 首轮数据：

| 模型 | 4K TTFT | 16K TTFT | 32K TTFT | 4K/16K/32K TPOT |
|---|---:|---:|---:|:-:|
| Qwen2.5-1.5B | 11.9 s | 68.8 s | 200.5 s | 59 / 213 / **418** ms |
| Qwen3-8B | 48.9 s | OOM | OOM | 88 ms |
| Llama-3-8B | 47.4 s | OOM | OOM | 79 ms |

1.5B 全跑通，8B 在 16K/32K 全 OOM——raw + fp32 + KV 三个加起来装不进 32 GB。

### C.3：打开 int4 KV，把 8B × 32K 救回来

xLLM 有一套完整的 int4 KV 流水线（`XLLM_KV_INT4_*` 系列 env，分 Part A/B/C/D），是 7 月初受 Bakouch 工作启发做的，一直没在长 prompt 上验证过。

C.3 干的事就是把这几个开关打开：

```
XLLM_KV_LAYER_AWARE=1
XLLM_KV_INT4=1
XLLM_KV_INT4_PARTC=1
XLLM_KV_INT4_PARTD=1
```

8B × 32K 的启动日志：

```
[kv_int4_partd] enabled: dropping Phase-B fp32 slab
[kv_int4_partd]   slab bytes: 2320.0 MB (vs 2880.0 MB, -560 MB)
[kv_int4_shadow]   projected @ 32K ctx, bs=1: Part B saves ~1540 MB
```

**总共省下约 2 GB，正好是"装得下"和"OOM"之间的差额。**

完整矩阵：

| 模型 | 4K TTFT | 16K TTFT | 32K TTFT | 4K/16K/32K TPOT |
|---|---:|---:|---:|:-:|
| Qwen2.5-1.5B | 11.9 s | 68.7 s | 200.3 s | 59 / 213 / 416 ms |
| Qwen3-8B | 48.9 s | 317.4 s | **960.8 s** | 88 / 329 / **632** ms |
| Llama-3-8B | 47.3 s | 296.9 s | 880.6 s | 79 / 292 / **559** ms |

**功能上算是打通了**——9 个格子都有数据了。但**性能没打到 roadmap 目标**：Qwen3-8B 32K TPOT 的目标是 ≤100 ms/tok，实测 632 ms/tok，差 6 倍。

根因记在这里，不粉饰：paged attention 要在 28K token 上做 int4 的 K/V 点积，6 GB 工作集装不进 L2，每个 token 都得从 HBM 读一遍。这是 kernel 调优问题，不属于 v0.7 的目标范围，defer 到 v0.8。

---

## 五、完整数据

### 短上下文 vs cuBLAS 基线（TPOT, ms/token）

| 测试格 | v0.6.6 | v0.7-c1 | 变化 |
|---|---:|---:|:-:|
| 1.5B bs=1 | 9.16 | 8.04 | −12% |
| 1.5B bs=2 | 6.50 | 4.74 | −27% |
| 1.5B bs=4 | 4.32 | 2.96 | −31% |
| 1.5B bs=8 | 3.64 | 2.90 | −20% |
| Qwen3-8B bs=1 | 26.76 | 25.29 | −5% |
| Qwen3-8B bs=2 | 20.39 | 16.26 | −20% |
| Qwen3-8B bs=4 | 12.48 | 11.25 | −10% |
| Qwen3-8B bs=8 | 9.64 | 12.35 | **+28%**（已接受） |
| Llama-3-8B bs=1 | 24.68 | 23.37 | −5% |
| Llama-3-8B bs=2 | 18.86 | 14.86 | −21% |
| Llama-3-8B bs=4 | 11.41 | 10.08 | −12% |
| Llama-3-8B bs=8 | 9.15 | 11.46 | **+25%**（已接受） |

**12 格里赢了 10 个。** 剩下两个 8B bs=8 的回退是 L1 带宽天花板，属于硬件特性。

### 长上下文（v0.7 新增覆盖）

- **Qwen3-8B × 32K**：v0.7 之前 OOM → 现在 TTFT 960.8 s / TPOT 632 ms
- **Qwen2.5-1.5B × 32K**：v0.7 之前未测 → 现在 TTFT 200.3 s / TPOT 416 ms

### Commit 列表

- **Phase A**：`01f0597` A1 · `9ba9e14` A2.1 · `b06fee7` A2.2 · `8a3a29d` A3
- **Phase B**：`da45ac2` B1.1 · `a93e4fc` B1.2 · `65b2c5b` B1.3 · `237dc98` B2 · `dd603e7` B3 · `191ce95` B4 · `a5c11a4` bs=8 profiling
- **Phase C**：`0a1e512` C.1 · `da60b59` C.2 · `b316cfe` C.3

测试环境：V100-32GB（SM 7.0），Qwen2.5-1.5B / Qwen3-8B / Meta-Llama-3-8B，GGUF 量化。byte-eq（逐 token 输出与基线逐字节一致）是每个阶段的硬性发布门禁。

---

## 六、v0.7 里最值钱的两个教训

### 第一：改 kernel 之前先 profile

Phase A 结束时我们有一个看起来很合理的假设："chunk-of-4 的 kernel launch 是 8B bs=8 慢的原因"。B2 花了半天写 M=8 tile 去验证，只拿到 −1.8% —— 假设是错的。

真正的原因在 ncu 数据里躺着：L1 带宽打满 85%，是硬件天花板。**同样的半天，如果先花在 profiling 上，直接就能得出"该 accept"的正确判断。**

这条的反面同样重要：如果你的优化假设没有被数据支撑，那它只是一个待验证的猜测，而不是行动依据。

### 第二：把契约写进类型系统，别靠环境变量猜

A3 时我们试图把 DP4A 相关 env 默认打开，结果 per-slot bench 用 `max_batch=1` 却触发了 batched 路径的 raw 预加载，raw + fp32 双缓存 OOM 级联，8B bs=1 从 25 ms/tok 崩到 4260 ms/tok——**一个 100 倍的性能悬崖，而且是静默发生的**。

B1.2 加了 `XllmDecodeMode`，让调用方显式声明意图，误配就 loud abort。之后 B1.3 和 C.1 都在这个契约上扩展，再没翻过车。

环境变量适合做开关，不适合做契约。契约应该在编译期或初始化期就能校验——**能在第一秒炸出来的错误，不要留到生产负载里去查。**

---

## 七、下一步（v0.8）

- Qwen3-8B 32K TPOT 打到 100 ms/tok（int4 KV kernel 的 ncu 调优）
- Continuous batching + 长上下文（bs=8 × 32K）
- 用 v0.7 的数据点刷新与 vLLM 的三家族对比

---

*山野村夫，本文基于 xLLM v0.7 开发记录整理，2026.9.8*


10 天、13 个 commit，填掉三个真正卡死长上下文的坑。 短上下文 12 个测试格赢了 cuBLAS 基线 10 个；32K 上下文首 token 延迟从 6.9 秒压到 0.9 秒；8B 模型 × 32K prompt 从直接 OOM 变成能跑完。 另外还有一个"优化假设被 profiler 当场打脸"的故事，我觉得比性能数字更值得写。

一个优化假设，被 profiler 当场打脸——然后我们决定不优化了

农民伯伯

编程语言

开发工具

人工智能

xLLM v0.7实现长上下文性能突破：10天13个commit解决三大核心问题。32K上下文首token延迟从6.9秒降至0.9秒，8B模型×32K prompt从OOM到成功运行。短上下文12项测试中10项超越cuBLAS基线，TPOT最高提升31%。关键优化包括修复batch≥2整数路径、解耦TTFT与max_seq_len、开启int4 KV缓存。

数据结构

单元测试

代码解释

4核4G3M云服务器 新用户低至38元/年！

文章

问答

视频

教程

学习中心

腾讯云实验室

直播

竞赛

腾讯云代码分析专区

腾讯iOA零信任安全管理系统专区

腾讯云架构师技术同盟交流圈

腾讯云数据库专区

腾讯云智能顾问专区

腾讯云原生专区

腾讯混元专区

腾讯云TCE专区

腾讯云Lighthouse专区

腾讯云HAI专区

腾讯云Edgeone专区

腾讯云存储专区

腾讯云智能专区

腾讯轻联专区 

腾讯云开发专区

TAPD专区

腾讯轻量云游戏服专区

EdgeOne AI 安全实战专区

腾讯云最具价值专家

腾讯云架构师技术同盟

腾讯云创作之星

腾讯云开发者先锋

腾讯云代码助手

云原生构建

TAPD 敏捷项目管理

Cloud Studio

SDK中心

API中心

命令行工具

涵盖代码开发、场景应用、自动测试全流程，助你从零构建专属AI助手

一站式MCP教程库，解锁AI应用新玩法

聚焦“写作效率、视觉美观与运行性能”三方面进行全面升级，为您提供更高效、稳定的创作环境

社区富文本&Markdown编辑器全新改版上线，欢迎大家体验!

诚挚邀请您参与本次调研，分享您的真实使用感受与建议。您的反馈至关重要，感谢您的支持与参与！

社区新版编辑器体验调研

xLLM v0.7实现长上下文性能突破：10天13个commit解决三大核心问题。32K上下文首token延迟从6.9秒降至0.9秒，8B模型×32K prompt从OOM到成功运行。短上下文12项测试中10项超越cuBLAS基线，TPOT最高提升31%。关键优化包括修复batch≥2整数路径、解耦TTFT与max_s...

一个优化假设，被 profiler 当场打脸——然后我们决定不优化了-腾讯云开发者社区-腾讯云

一个优化假设，被 profiler 当场打脸——然后我们决定不优化了

一个优化假设，被 profiler 当场打脸——然后我们决定不优化了

社区

活动

圈层

关于

腾讯云开发者

热门产品

热门推荐

更多推荐