Run any Skill in Manus with one click

Get Started

ptx-debug

Stars3

Forks0

UpdatedJuly 1, 2026 at 08:49

PTX-EMU 专用调试技能 - 自动化调试配置选择和场景化调试方法

Installation

Install with Codex or Claude Copy this prompt, paste it into Codex, Claude, or another assistant, and let it review the skill page and install it for you.

Run Skill in Manus

Source

chisuhua

chisuhua/PTX-EMU

View GitHub Repository View Creator Repositories

Download

Run Skill in Manus

Related occupationsSOC

Based on SOC occupation classification

Software DevelopersComputer and Mathematical Occupations·SOC 15-1252

SKILL.md

readonly

PTX-EMU 调试技能

核心原则

自动化调试配置选择：根据问题类型自动选择最合适的调试配置，避免手动配置。

场景化调试方法：针对不同调试场景（解析错误、内存问题、指令错误、性能问题）提供专门的调试流程。

证据驱动：使用适当的调试配置收集证据，定位问题根源。

调试配置自动选择

配置选择矩阵

问题类型	关键词	自动选择配置	日志级别	输出目标
PTX 解析错误	"解析错误", "语法错误", "ANTLR", "grammar"	`configs/verbose_trace_config.ini`	trace	file
测试失败	"测试失败", "test failed", "ctest"	`configs/dev_debug_config.ini`	debug	both
程序崩溃	"崩溃", "segfault", "SIGSEGV", "core"	`configs/verbose_trace_config.ini`	trace	file
内存问题	"内存", "memory", "越界", "非法访问"	`configs/memory_debug_config.ini`	debug/trace	both
指令错误	"指令", "instruction", "结果不对", "执行错误"	`configs/instruction_debug_config.ini`	info/trace	both
性能问题	"性能", "慢", "优化", "benchmark"	`configs/perf_config.ini`	warning	console
日常调试	"调试", "debug", "看一下"	`configs/dev_debug_config.ini`	debug	both
默认/未知	其他情况	`configs/dev_debug_config.ini`	debug	both

自动选择流程

1. 分析用户问题描述
   ↓
2. 提取关键词匹配问题类型
   ↓
3. 根据配置选择矩阵确定调试配置
   ↓
4. 使用 ./scripts/debug-run.sh 选择配置
   ↓
5. 运行程序收集调试信息
   ↓
6. 分析日志文件，定位问题

场景化调试方法

场景 1: PTX 解析错误

触发条件: 用户报告 "PTX 解析错误", "语法错误", "ANTLR 错误"

自动行动:

选择配置: configs/verbose_trace_config.ini

 ./scripts/debug-run.sh verbose ./build/bin/dummy-args

运行解析器:

# 或使用 test-ptx 直接解析
./build/bin/test-ptx source.ptx 2>&1 | grep -E "parser|lexer"

收集证据:

# 查看解析器日志
grep "PTX version\|PTX target\|Address size" ptx_emu_trace.log

# 查看 ANTLR 解析过程
grep "Visiting PTX\|parser\|lexer" ptx_emu_trace.log

分析重点:
- PTX 版本号是否正确
- 目标架构是否支持
- 解析到哪条指令时失败
- ANTLR 错误消息的具体内容

定位语法文件:

# 查找相关的语法文件
grep -r "错误中提到的指令名" src/grammar/

生成测试用例:
- 提取导致错误的 PTX 代码
- 添加到 tests/ptx/ 目录
- 运行 ./tests/ptx/test_all_ptx.sh 验证

场景 2: 测试失败

触发条件: 用户报告 "测试失败", "test failed", "ctest 不通过"

自动行动:

选择配置: debug 配置（configs/dev_debug_config.ini）
```
./scripts/debug-run.sh debug ./build/bin/程序
```
运行失败测试:
```
cd build && ctest -R 测试名 -V
```

收集证据:

# 查看测试输出
ctest -R 测试名 --output-on-failure

# 查看详细日志
tail -100 ptx_emu_debug.log

分析重点:
- 测试期望值 vs 实际值
- 失败前的最后几条指令
- 寄存器状态是否正确
- 内存访问是否越界

对比参考实现:

# 查找类似的通过测试
grep -r "类似功能" tests/

生成调试脚本:

# 创建最小复现
cat > debug_test.sh << 'EOF'
#!/bin/bash

./scripts/debug-run.sh debug ./build/bin/测试程序 EOF chmod +x debug_test.sh


---

### 场景 3: 程序崩溃

**触发条件**: 用户报告 "崩溃", "segfault", "SIGSEGV", "core dumped"

**自动行动**:

1. **选择配置**: `verbose` 配置（`configs/verbose_trace_config.ini`）
```bash
./scripts/debug-run.sh verbose ./build/bin/程序

运行程序获取堆栈:

# 使用 gdb 获取堆栈
gdb -batch -ex "run" -ex "bt" ./build/bin/程序

# 或直接运行获取 core dump
ulimit -c unlimited
./build/bin/程序

收集证据:

# 查看崩溃前的指令序列
grep "PC\[" ptx_emu_trace.log | tail -50

# 查看内存访问
grep "Memory allocated\|Memory freed" ptx_emu_trace.log

# 查看寄存器状态
grep "Register.*contains" ptx_emu_trace.log

分析重点:
- 崩溃时的 PC 值
- 最后执行的指令
- 访问的内存地址是否合法
- 寄存器值是否异常

定位崩溃点:

# 使用 addr2line 定位
addr2line -e ./build/bin/程序 崩溃地址

生成修复方案:
- 检查空指针
- 检查数组越界
- 检查 use-after-free
- 检查资源未初始化

场景 4: 内存问题

触发条件: 用户报告 "内存错误", "越界", "非法访问"

自动行动:

选择配置: memory 配置（configs/memory_debug_config.ini）
```
./scripts/debug-run.sh memory ./build/bin/程序
```
运行程序:
```
./build/bin/程序
```

收集证据:

# 查看所有内存操作
grep "\[mem\]" ptx_emu_memory_debug.log

# 查看特定地址的访问
grep "0x具体地址" ptx_emu_memory_debug.log

# 查看内存分配/释放配对
grep "Memory allocated\|Memory freed" ptx_emu_memory_debug.log

分析重点:
- 内存分配大小 vs 访问大小
- 释放后是否再次访问
- 是否越界访问
- 地址对齐是否正确

使用工具验证:

# 如果有 valgrind 支持
valgrind --leak-check=full --track-origins=yes ./build/bin/程序

# 或使用 AddressSanitizer（需要重新编译）
# cmake -DCMAKE_BUILD_TYPE=Debug -DENABLE_ASAN=ON ..

场景 5: 指令执行错误

触发条件: 用户报告 "指令错误", "结果不对", "执行异常"

自动行动:

选择配置: instruction 配置（configs/instruction_debug_config.ini）
```
./scripts/debug-run.sh instruction ./build/bin/程序
```
运行程序:
```
./build/bin/程序
```

收集证据:

# 查看指令执行序列
grep "\[instr\]" ptx_emu_instr_debug.log

# 查看特定指令类型
grep "st.global\|ld.global\|add" ptx_emu_instr_debug.log

# 查看寄存器变化
grep "Commit.*%" ptx_emu_instr_debug.log

分析重点:
- 指令的输入操作数
- 指令的输出结果
- 前后寄存器状态对比
- 指令语义是否正确实现

对比参考:

# 查找 PTX ISA 文档
# 或查看其他实现
grep -r "指令名" src/ptxsim/instructions/

验证修复:
- 创建单元测试
- 验证各种边界情况
- 对比硬件行为

场景 6: 性能问题

触发条件: 用户报告 "性能慢", "需要优化", "benchmark"

自动行动:

选择配置: perf 配置（configs/perf_config.ini）
```
./scripts/debug-run.sh perf ./build/bin/程序
```

性能测量:

# 测量执行时间
time ./build/bin/程序

# 多次运行取平均
for i in {1..10}; do time ./build/bin/程序; done

性能分析:

# 使用 perf（如果可用）
perf record ./build/bin/程序
perf report

# 或使用 gprof
cmake -DCMAKE_BUILD_TYPE=Profiling ..
make
./build/bin/程序
gprof ./build/bin/程序 gmon.out > profile.txt

分析重点:
- 热点函数

场景 7: Barrier 同步卡住 (SIMT 特定)

触发条件:

测试超时或长时间无响应
barrier 相关测试卡住 (test_syncthreads, test_warp_divergence 等)
日志显示 threads blocked at barrier

自动行动:

快速诊断 (类型 1 - 不需要仿真):

# 1. 验证 PTX 解析是否正确
/usr/local/cuda/bin/cuobjdump -ptx ./build/bin/test_xxx 2>/dev/null | \
  tail -n+12 > /tmp/test_xxx.ptx
./build/bin/test-ptx /tmp/test_xxx.ptx 2>&1 | grep "bar.warp.sync"

启用详细日志:

./scripts/debug-run.sh verbose ./build/bin/test_xxx

运行收集日志:

timeout 120 ./build/bin/test_xxx 2>&1 | tee /tmp/test_output.log

分析关键模式:

grep "Initialized wbar" /tmp/test_output.log
grep "Barrier complete" /tmp/test_output.log
grep "blocked at bar.warp.sync" /tmp/test_output.log | tail -10

预期行为:

# 正常:
bar.warp.sync: Initialized wbar[0] with mask=0xFFFFFFFF, reconvergence_pc=N
Lane 0 arrived (arrived=1/32)
...
Lane 31 arrived (arrived=32/32)
bar.warp.sync: Barrier complete, releasing 32 threads to PC=N

# 异常:
- 停在 arrived=X/32 (X<32) → 部分线程未到达
- Barrier complete 但仍卡住 → PC 更新或调度问题

常见问题诊断表:

现象	可能原因	排查步骤
一直 blocked	arrive() 未被所有线程调用	检查指令调度逻辑
is_complete 总 false	arrived_mask 未递增	检查 lane_id 是否正确
PC 未更新	set_thread_pc() 未调用	检查 complete 后代码路径
唤醒后错误指令	pc_stack 未同步	检查 update_pc_stack()

场景 7.5: Divergent Execution 路径不同步 (高级)

触发条件:

test_nested_sync 或类似 divergent execution 测试卡住
日志显示不同的 lane 到达不同的 reconvergence_pc
arrived 计数停滞 (如 arrived=16/32)

根因分析:

PTX 语法: @%p1 bra $L__BB2_2 创建两条执行路径
Visitor 翻译: bar.sync → bar.warp.sync 时使用固定 reconvergence_pc = size + 1
问题: 当 divergent 后再次汇合时，不同路径的下一条指令位置不同

调试步骤:

# 1. 查看每个 lane 到达的 PC
grep "arrived at bar.warp.sync" /tmp/test_output.log | 
  awk -F' '"{print
### 场景 8: 数值异常

**触发条件**: 测试报告 Numerical Exception

**自动行动**:

1. **选择配置**: `configs/dev_debug_config.ini`
2. **运行分析**: `grep -E "Numerical|NaN|Inf" ptx_emu_debug.log`
3. **分析重点**: 哪个指令产生异常值，输入操作数是否合法

---

   - 内存访问模式
   - 缓存命中率
   - 分支预测

---

## 自动化调试流程

### 标准调试流程（自动触发）

接收用户问题 ↓
分析问题类型，选择调试配置 ↓
使用 debug-run.sh 选择配置：./scripts/debug-run.sh {配置名} ./build/bin/{程序} ↓
运行程序收集日志 ↓
分析日志文件，提取关键信息 ↓
定位问题根源 ↓
生成修复方案 ↓
验证修复


### 快捷脚本使用

**技能会自动使用或指导用户使用快捷脚本**:

```bash
# 技能自动执行
./scripts/debug-run.sh {配置名} ./build/bin/{程序}

# 示例
./scripts/debug-run.sh memory ./build/bin/dummy-args
./scripts/debug-run.sh trace ./build/bin/RAY 512 512

日志分析技术

常用 grep 命令

# 查看错误
grep "ERROR\|FATAL" ptx_emu_*.log

# 查看特定组件
grep "\[mem\]\|\[instr\]\|\[exec\]" ptx_emu_*.log

# 查看特定指令
grep "st.global\|ld.global" ptx_emu_*.log

# 查看寄存器
grep "Register.*contains\|Commit.*%" ptx_emu_*.log

# 查看内存操作
grep "Memory allocated\|Memory freed" ptx_emu_*.log

# 时间线分析
grep "\[CLK:" ptx_emu_*.log | head -100

日志分析模式

模式 1: 崩溃分析

# 找到崩溃前的最后状态
tac ptx_emu_trace.log | grep -A 50 "最后的消息"

模式 2: 内存泄漏

# 对比分配和释放
grep "Memory allocated" ptx_emu.log > alloc.txt
grep "Memory freed" ptx_emu.log > free.txt
diff alloc.txt free.txt

模式 3: 指令追踪

# 查看指令执行序列
grep "PC\[" ptx_emu_instr_debug.log | awk '{print $5}'

调试输出位置

日志文件

配置文件	日志文件	用途
`configs/dev_debug_config.ini`	`ptx_emu_debug.log`	日常调试
`configs/verbose_trace_config.ini`	`ptx_emu_trace.log`	详细跟踪
`configs/memory_debug_config.ini`	`ptx_emu_memory_debug.log`	内存调试
`configs/instruction_debug_config.ini`	`ptx_emu_instr_debug.log`	指令调试

快速查看

# 实时查看
tail -f ptx_emu_*.log

# 查看最新 N 行
tail -N ptx_emu_debug.log

# 分页查看
less -R ptx_emu_trace.log

# 搜索查看
grep "关键词" ptx_emu_debug.log | less

与其他技能协作

依赖技能

cpp-debug: 提供通用 C++ 调试方法
systematic-debugging: 提供系统化调试流程
cuda-ptx: 提供 CUDA/PTX 专业知识

⚠️ 与 `ptx-grammar-modification` 的分工

重要: 本项目有专门的 PTX 语法修复技能 ptx-grammar-modification（位于 docs/skills/）。

职责划分:

问题类型	使用技能	说明
ANTLR 解析错误	`ptx-grammar-modification`	语法文件修改、重新生成解析器
语法文件修改	`ptx-grammar-modification`	修改 `.g4` 文件
`no viable alternative`	`ptx-grammar-modification`	语法规则修复
`mismatched input`	`ptx-grammar-modification`	Token 定义修复
测试失败	`ptx-debug`	运行时错误、逻辑错误
内存问题	`ptx-debug`	内存访问、越界
指令错误	`ptx-debug`	指令执行结果不对
程序崩溃	`ptx-debug`	segfault、core dump
性能问题	`ptx-debug`	性能优化、benchmark

自动识别逻辑:

当用户提到 "PTX 解析错误" 时:
  1. ptx-debug 技能触发
  2. 检查错误类型
  3. IF 错误是 ANTLR 解析错误 (no viable alternative, mismatched input):
       → 🛑 停止调试流程
       → 建议用户使用 ptx-grammar-modification 技能
       → 提示运行 ./tests/ptx/test_all_ptx.sh
  4. ELSE 错误是运行时错误:
       → 继续调试流程
       → 选择 verbose_trace 配置
       → 收集日志分析

协作场景

内存泄漏: 使用 cpp-debug 的 valgrind/ASan 方法
解析错误:
- 语法问题 → ptx-grammar-modification
- 日志分析 → ptx-debug 辅助
系统调试: 遵循 systematic-debugging 的四阶段流程
PTX ISA 查询: 使用 cuda-ptx 的本地文档

验证标准

调试完成的标志:

示例使用

示例 1: 解析错误

用户: "PTX 解析失败，报告语法错误"

技能自动行动:

识别为"PTX 解析错误"
选择 verbose_trace_config.ini

执行：

./scripts/debug-run.sh verbose ./build/bin/dummy-args
grep "parser\|lexer" ptx_emu_trace.log

分析错误位置
修复语法文件
验证修复

示例 2: 测试失败

用户: "test_memory_manager 测试失败"

技能自动行动:

识别为"测试失败"
选择 debug_config.ini

执行：

./scripts/debug-run.sh debug ./build/bin/测试程序
cd build && ctest -R test_memory_manager -V
tail -100 ptx_emu_debug.log

分析失败原因
生成修复

示例 3: 性能优化

用户: "RAY benchmark 运行太慢"

技能自动行动:

识别为"性能问题"
选择 perf_config.ini

执行：

./scripts/debug-run.sh perf ./build/bin/RAY 512 512

分析热点
提出优化建议

诊断方法分类

核心原则: 80% 的问题可以通过不需要仿真的快速诊断定位，只在必要时运行完整仿真

快速诊断 (<30s) - 不需要仿真

适用场景:

PTX 语法解析错误
Label 注册问题
CFG 分析警告
指令翻译问题 (如 bar.sync → bar.warp.sync)

验证方法:

方法	命令	耗时	适用
test-ptx 解析	`./build/bin/test-ptx tests/ptx/test.ptx`	<1s	语法、翻译验证
手动提取 PTX	`cuobjdump -ptx binary > test.ptx`	<5s	检查生成的 PTX
最小化 PTX	手动创建简化 PTX 文件	<10s	隔离特定指令
CFG 日志	添加日志到 cfg_builder.cpp	<30s	控制流分析

快速验证流程:

# 1. 从二进制提取 PTX
/usr/local/cuda/bin/cuobjdump -ptx ./build/bin/test_xxx 2>/dev/null | tail -n+12 > /tmp/test.ptx

# 2. 用 test-ptx 验证解析
./build/bin/test-ptx /tmp/test.ptx 2>&1 | grep -E "bar|bra|label|PASS|FAIL"

# 3. 查看特定指令翻译
./build/bin/test-ptx /tmp/test.ptx 2>&1 | grep -A2 "bar.warp.sync"

成功案例:

✅ PTX label 语法修复：通过 test-ptx 快速验证 33/33 语法测试
✅ bra 指令 predicate: 通过解析输出确认 target 正确注册
✅ CFG 空 target 警告：通过添加日志快速定位并修复

仿真诊断 (>30s) - 需要完整运行

适用场景:

Barrier 同步卡住
线程调度问题
内存访问错误
寄存器值错误
numerical exception

验证方法:

方法	命令	耗时	适用
完整单元测试	`ctest -R test_xxx -V`	30-120s	功能回归
日志分析	启用 PTX_DEBUG	30-60s	执行跟踪
GDB 调试	`gdb --args ./bin/test_xxx`	5-10min	崩溃定位
ASan/UBSan	Debug 构建 + sanitizer	2-5min	内存错误

调试流程:

# 1. 启用详细日志
export PTX_LOG_LEVEL=debug
./build/bin/test_xxx 2>&1 | grep -E "bar|sync|blocked|complete"

# 2. 定位卡住点
./build/bin/test_xxx 2>&1 | tail -50

# 3. GDB 调试 (如果崩溃)
gdb --args ./build/bin/test_xxx
(gdb) run
(gdb) bt  # 崩溃时获取堆栈

决策树:

遇到问题
    │
    ├─ 解析错误/语法问题？
    │   └─→ 使用 test-ptx (1s) ← 优先
    │
    ├─ Label/CFG 警告？
    │   └─→ 添加日志，重解析 (10s) ← 优先
    │
    ├─ 指令执行卡住？
    │   └─→ 启用调试日志 (30s)
    │       └─ 仍无法定位？
    │           └─→ GDB 调试 (5min)
    │
    └─ 崩溃/SegFault？
        └─→ GDB + ASan (2min)

Barrier 调试最佳实践

屏障调试流程 (类型 2 - 需要仿真)

# 1. 验证 PTX 解析 (类型 1 - 快速)
./build/bin/test-ptx test.ptx | grep "bar.warp.sync"

# 2. 启用 barrier 日志
grep -r "PTX_DEBUG" src/ptxsim/instructions/barrier.cpp
# 确保相关日志启用

# 3. 运行测试并抓取日志
./build/bin/test_xxx 2>&1 | grep -E "arrived|complete|blocked|released"

# 4. 分析关键指标
# - participation_mask 是否正确？
# - arrived_mask 是否递增？
# - is_complete() 何时返回 true？
# - set_thread_pc() 是否被调用？
# - update_pc_stack() 是否同步？

关键日志模式:

bar.warp.sync: Initialized wbar[0] with mask=0xFFFFFFFF, reconvergence_pc=N
Lane X arrived at bar.warp.sync (mask=0xFFFFFFFF, pc=N)
bar.warp.sync: Barrier complete, releasing N threads to PC=N
Lane X blocked at bar.warp.sync (arrived=N/32)

预期行为:

初始化 wbar，participation_mask = 0xFFFFFFFF
每个 lane 调用 arrive()，arrived_mask 递增
当 arrived_mask == participation_mask 时，barrier complete
所有参与线程 PC 更新为 reconvergence_pc
is_blocked 清除，status 设为 Active

常见问题诊断表:

现象	可能原因	排查步骤
一直 blocked	arrive() 未被调用	检查指令调度器是否调度过 Blocked 线程
is_complete 总 false	arrived_mask 未递增	检查 lane_id 是否正确传入 arrive()
PC 未更新	set_thread_pc() 未被调用	检查 is_complete() 返回后的代码路径
唤醒后执行错误指令	pc_stack 未同步	检查 update_pc_stack() 调用点
participation_mask=0	指令翻译错误	检查 visitor 中 barrier operand 提取

关键代码位置:

src/ptx_parser/ptx_visitor_barrier.cpp  - bar.sync → bar.warp.sync 翻译
src/ptxsim/instructions/barrier.cpp     - Wbar 实现与 arrive()/is_complete()
include/ptxsim/wbar.h                    - Wbar 数据结构定义

自愈型 Bug 检测（重要）

问题：update_active_mask() 自愈

execute_warp_instruction() 末尾调用 update_active_mask()，从 warp_state.threads[i].is_active 重新计算 active_mask[]。这意味着：

一条指令内临时破坏 active_mask[]（如 barrier set_active_mask(arrived_mask) 覆写）
下一条指令时 update_active_mask() 自动从线程状态恢复
端到端集成测试只看 is_finished() → 可能 PASS 即使有 bug

检测方法

单元测试必须：在特定调用后直接断言 get_active_mask()，不要只检查 is_finished()
集成测试补充：检查所有 lane 的最终 PC — 如果某些 lane 停在 barrier PC 而非 post-barrier PC，说明有 lane 丢失
生产症状检查：simpleGEMM-* 不再卡住 = fix 生效；但结果是否正确 = divergence reconciliation 是否完整（独立问题）

诊断命令

# 1. 单元测试 — 验证 set_active_mask 行为（RED guard）
ctest -R "post_barrier_two_halves" -V

# 2. 集成测试 — 仅看 warp 是否完成（smoke test，可能被骗过）
ctest -R "integration_post_barrier_two_halves" -V

# 3. 真实症状 — 看 lane 是否全部推进
for t in int float double; do
  timeout 5 ./build/bin/simpleGEMM-$t 2>&1 | tail -1
done

典型症状对照表

检查	含义
`ctest -R "post_barrier"` RED	bug 存在 — set_active_mask 行为不对
`ctest -R "integration_post_barrier"` PASS 但 simpleGEMM 仍输出 0	自愈型 bug — warp 完成但数据错误（divergence reconciliation 独立 bug）
`simpleGEMM` 卡住	bug 存在 — active_mask 被破坏后没有 lane 可调度
`simpleGEMM` 退出但结果非 0	fix 生效，divergence reconciliation 正常

经验教训

日志文件大小: trace 级别可能产生数百 MB 日志
性能影响: 详细日志会降低 10-100 倍性能
及时清理: 定期清理旧日志文件
配置恢复: 调试后记得恢复到 release 配置

更新记录

2026-03-23: 初始版本，包含 6 种调试场景的自动化方法
后续更新：根据新场景添加调试配置和方法

场景 7.5: Divergent Execution 路径不同步 (高级)

触发条件:

test_nested_sync 或类似 divergent execution 测试卡住
日志显示不同的 lane 到达不同的 reconvergence_pc
arrived 计数停滞 (如 arrived=16/32)

根因分析:

PTX 语法: @%p1 bra $L__BB2_2 创建两条执行路径
Visitor 翻译: bar.sync → bar.warp.sync 时使用固定 reconvergence_pc = size + 1
问题: 当 divergent 后再次汇合时，不同路径的下一条指令位置不同

调试步骤:

# 1. 查看每个 lane 到达的 PC
grep "arrived at bar.warp.sync" /tmp/test_output.log | \
  grep -oP 'pc=\K\d+' | sort -n | uniq -c

# 2. 检查是否所有 lane 都到达同一个 barrier
grep "arrived=" /tmp/test_output.log | \
  grep -oP 'arrived=\K\d+/\d+' | sort | uniq -c

# 3. 对比 label 注册位置
grep "Registering label" /tmp/test_output.log

实际案例 (test_nested_sync):

Lane 13 arrived pc=26  # 在分支路径中
Lane 14 arrived pc=12  # 直接跳转路径
arrived=16/32          # 停滞

问题根因: threads 16-31 跳转到 L__BB2_2 (PC=23)，然后执行第二个 barrier (PC=25)，但实际到达的 reconvergence_pc=26 不正确。

修复建议:

检查 visitor 中的 barrier PC 计算逻辑
确保 divergent 路径正确汇合到同一个 label
使用 CFG analysis 验证 reconvergence 点
查看 ptx_visitor_barrier.cpp 中 next_pc 的计算

场景 9: S_BAR 多 warp CTA 级 barrier 死锁 (新增 2026-07-01)

触发条件:

CTA 级 barrier 测试（integration_cta_barrier_*, integration_cute_rmsnorm_*）卡住或 FAIL
日志显示线程 blocked at barrier 但 barrier 永不完成

调试方法:

# 1. 启用线程追踪（替代全局 fprintf，受 debug config 控制）
# configs/debug_config.ini 中 component.thread=debug
./build/bin/test_xxx 2>&1 | grep "\[thread\]" | grep -E "bar.sync|barrier|blocked|released" | head -30

# 2. 检查 CTABarrier 到达情况
./build/bin/test_xxx 2>&1 | grep -E "arrive|complete|uninitialized"

# 3. 检查调度器是否能看到释放的线程
./build/bin/test_xxx 2>&1 | grep -E "lanes_by_pc|active_mask" | head -20

常见问题诊断表:

现象	可能原因	排查步骤
`arrive called on uninitialized barrier`	`CTABarrier::reset()` 清除了 `is_initialized_`	检查 reset() 是否保持初始化状态
barrier complete 但 warp 无法前进	`is_active` 未被 barrier release 恢复	检查 `release_cta_barrier` 中的 `ts.is_active` 设置
warp 跳过 barrier handler	测试 driver 提前返回 `post_barrier_pc`	添加 `is_cta_barrier_complete()` 门控
divergence + bar.sync 卡住	`reconvergence_pc` 指向错误位置	检查 `make_bra_pred` 的 reconv 参数是否等于真实汇聚点

关键代码位置:

src/ptxsim/barrier/barrier_module.cpp  - release_cta_barrier (is_active 恢复)
src/ptxsim/barrier/cta_barrier.cpp     - reset() (保留初始化状态)
include/ptxsim/testing/scheduler_utils.h - step_warp (PC 边界检查)

调试命令速查:

# 打印 per-lane 指令执行序列
./build/bin/test_xxx 2>&1 | grep "\[thread\]" | grep -oP 'lane=\d+ pc=\d+' | sort | uniq -c

# 打印所有 barrier 操作时间线
./build/bin/test_xxx 2>&1 | grep -E "\[CLK:\d+.*arrive|complete|release|reset" | awk '{print $1,$2,$NF}'

name	ptx-debug
description	PTX-EMU 专用调试技能 - 自动化调试配置选择和场景化调试方法
when_to_use	当用户在 PTX-EMU 项目中遇到以下问题时自动触发： - "PTX 解析错误", "语法错误", "ANTLR 错误" - "测试失败", "单元测试不通过" - "程序崩溃", "segfault", "core dumped" - "内存错误", "非法访问", "越界" - "指令执行错误", "结果不对" - "性能慢", "优化" - "调试这个", "分析一下问题" 触发关键词： - "ptx", "cubin", "GPU", "kernel" - "test_", "ctest", "单元测试" - "debug", "调试", "analyze", "分析" - "为什么失败", "怎么回事", "哪里错了"
skills_required	["cpp-debug","systematic-debugging","regression-bisect"]

PTX-EMU 调试技能

核心原则

自动化调试配置选择：根据问题类型自动选择最合适的调试配置，避免手动配置。

场景化调试方法：针对不同调试场景（解析错误、内存问题、指令错误、性能问题）提供专门的调试流程。

证据驱动：使用适当的调试配置收集证据，定位问题根源。

调试配置自动选择

配置选择矩阵

问题类型	关键词	自动选择配置	日志级别	输出目标
PTX 解析错误	"解析错误", "语法错误", "ANTLR", "grammar"	`configs/verbose_trace_config.ini`	trace	file
测试失败	"测试失败", "test failed", "ctest"	`configs/dev_debug_config.ini`	debug	both
程序崩溃	"崩溃", "segfault", "SIGSEGV", "core"	`configs/verbose_trace_config.ini`	trace	file
内存问题	"内存", "memory", "越界", "非法访问"	`configs/memory_debug_config.ini`	debug/trace	both
指令错误	"指令", "instruction", "结果不对", "执行错误"	`configs/instruction_debug_config.ini`	info/trace	both
性能问题	"性能", "慢", "优化", "benchmark"	`configs/perf_config.ini`	warning	console
日常调试	"调试", "debug", "看一下"	`configs/dev_debug_config.ini`	debug	both
默认/未知	其他情况	`configs/dev_debug_config.ini`	debug	both

自动选择流程

1. 分析用户问题描述
   ↓
2. 提取关键词匹配问题类型
   ↓
3. 根据配置选择矩阵确定调试配置
   ↓
4. 使用 ./scripts/debug-run.sh 选择配置
   ↓
5. 运行程序收集调试信息
   ↓
6. 分析日志文件，定位问题

场景化调试方法

场景 1: PTX 解析错误

触发条件: 用户报告 "PTX 解析错误", "语法错误", "ANTLR 错误"

自动行动:

选择配置: configs/verbose_trace_config.ini

 ./scripts/debug-run.sh verbose ./build/bin/dummy-args

运行解析器:

# 或使用 test-ptx 直接解析
./build/bin/test-ptx source.ptx 2>&1 | grep -E "parser|lexer"

收集证据:

# 查看解析器日志
grep "PTX version\|PTX target\|Address size" ptx_emu_trace.log

# 查看 ANTLR 解析过程
grep "Visiting PTX\|parser\|lexer" ptx_emu_trace.log

分析重点:
- PTX 版本号是否正确
- 目标架构是否支持
- 解析到哪条指令时失败
- ANTLR 错误消息的具体内容

定位语法文件:

# 查找相关的语法文件
grep -r "错误中提到的指令名" src/grammar/

生成测试用例:
- 提取导致错误的 PTX 代码
- 添加到 tests/ptx/ 目录
- 运行 ./tests/ptx/test_all_ptx.sh 验证

场景 2: 测试失败

触发条件: 用户报告 "测试失败", "test failed", "ctest 不通过"

自动行动:

选择配置: debug 配置（configs/dev_debug_config.ini）
```
./scripts/debug-run.sh debug ./build/bin/程序
```
运行失败测试:
```
cd build && ctest -R 测试名 -V
```

收集证据:

# 查看测试输出
ctest -R 测试名 --output-on-failure

# 查看详细日志
tail -100 ptx_emu_debug.log

分析重点:
- 测试期望值 vs 实际值
- 失败前的最后几条指令
- 寄存器状态是否正确
- 内存访问是否越界

对比参考实现:

# 查找类似的通过测试
grep -r "类似功能" tests/

生成调试脚本:

# 创建最小复现
cat > debug_test.sh << 'EOF'
#!/bin/bash

./scripts/debug-run.sh debug ./build/bin/测试程序 EOF chmod +x debug_test.sh


---

### 场景 3: 程序崩溃

**触发条件**: 用户报告 "崩溃", "segfault", "SIGSEGV", "core dumped"

**自动行动**:

1. **选择配置**: `verbose` 配置（`configs/verbose_trace_config.ini`）
```bash
./scripts/debug-run.sh verbose ./build/bin/程序

运行程序获取堆栈:

# 使用 gdb 获取堆栈
gdb -batch -ex "run" -ex "bt" ./build/bin/程序

# 或直接运行获取 core dump
ulimit -c unlimited
./build/bin/程序

收集证据:

# 查看崩溃前的指令序列
grep "PC\[" ptx_emu_trace.log | tail -50

# 查看内存访问
grep "Memory allocated\|Memory freed" ptx_emu_trace.log

# 查看寄存器状态
grep "Register.*contains" ptx_emu_trace.log

分析重点:
- 崩溃时的 PC 值
- 最后执行的指令
- 访问的内存地址是否合法
- 寄存器值是否异常

定位崩溃点:

# 使用 addr2line 定位
addr2line -e ./build/bin/程序 崩溃地址

生成修复方案:
- 检查空指针
- 检查数组越界
- 检查 use-after-free
- 检查资源未初始化

场景 4: 内存问题

触发条件: 用户报告 "内存错误", "越界", "非法访问"

自动行动:

选择配置: memory 配置（configs/memory_debug_config.ini）
```
./scripts/debug-run.sh memory ./build/bin/程序
```
运行程序:
```
./build/bin/程序
```

收集证据:

# 查看所有内存操作
grep "\[mem\]" ptx_emu_memory_debug.log

# 查看特定地址的访问
grep "0x具体地址" ptx_emu_memory_debug.log

# 查看内存分配/释放配对
grep "Memory allocated\|Memory freed" ptx_emu_memory_debug.log

分析重点:
- 内存分配大小 vs 访问大小
- 释放后是否再次访问
- 是否越界访问
- 地址对齐是否正确

使用工具验证:

# 如果有 valgrind 支持
valgrind --leak-check=full --track-origins=yes ./build/bin/程序

# 或使用 AddressSanitizer（需要重新编译）
# cmake -DCMAKE_BUILD_TYPE=Debug -DENABLE_ASAN=ON ..

场景 5: 指令执行错误

触发条件: 用户报告 "指令错误", "结果不对", "执行异常"

自动行动:

选择配置: instruction 配置（configs/instruction_debug_config.ini）
```
./scripts/debug-run.sh instruction ./build/bin/程序
```
运行程序:
```
./build/bin/程序
```

收集证据:

# 查看指令执行序列
grep "\[instr\]" ptx_emu_instr_debug.log

# 查看特定指令类型
grep "st.global\|ld.global\|add" ptx_emu_instr_debug.log

# 查看寄存器变化
grep "Commit.*%" ptx_emu_instr_debug.log

分析重点:
- 指令的输入操作数
- 指令的输出结果
- 前后寄存器状态对比
- 指令语义是否正确实现

对比参考:

# 查找 PTX ISA 文档
# 或查看其他实现
grep -r "指令名" src/ptxsim/instructions/

验证修复:
- 创建单元测试
- 验证各种边界情况
- 对比硬件行为

场景 6: 性能问题

触发条件: 用户报告 "性能慢", "需要优化", "benchmark"

自动行动:

选择配置: perf 配置（configs/perf_config.ini）
```
./scripts/debug-run.sh perf ./build/bin/程序
```

性能测量:

# 测量执行时间
time ./build/bin/程序

# 多次运行取平均
for i in {1..10}; do time ./build/bin/程序; done

性能分析:

# 使用 perf（如果可用）
perf record ./build/bin/程序
perf report

# 或使用 gprof
cmake -DCMAKE_BUILD_TYPE=Profiling ..
make
./build/bin/程序
gprof ./build/bin/程序 gmon.out > profile.txt

分析重点:
- 热点函数

场景 7: Barrier 同步卡住 (SIMT 特定)

触发条件:

测试超时或长时间无响应
barrier 相关测试卡住 (test_syncthreads, test_warp_divergence 等)
日志显示 threads blocked at barrier

自动行动:

快速诊断 (类型 1 - 不需要仿真):

# 1. 验证 PTX 解析是否正确
/usr/local/cuda/bin/cuobjdump -ptx ./build/bin/test_xxx 2>/dev/null | \
  tail -n+12 > /tmp/test_xxx.ptx
./build/bin/test-ptx /tmp/test_xxx.ptx 2>&1 | grep "bar.warp.sync"

启用详细日志:

./scripts/debug-run.sh verbose ./build/bin/test_xxx

运行收集日志:

timeout 120 ./build/bin/test_xxx 2>&1 | tee /tmp/test_output.log

分析关键模式:

grep "Initialized wbar" /tmp/test_output.log
grep "Barrier complete" /tmp/test_output.log
grep "blocked at bar.warp.sync" /tmp/test_output.log | tail -10

预期行为:

# 正常:
bar.warp.sync: Initialized wbar[0] with mask=0xFFFFFFFF, reconvergence_pc=N
Lane 0 arrived (arrived=1/32)
...
Lane 31 arrived (arrived=32/32)
bar.warp.sync: Barrier complete, releasing 32 threads to PC=N

# 异常:
- 停在 arrived=X/32 (X<32) → 部分线程未到达
- Barrier complete 但仍卡住 → PC 更新或调度问题

常见问题诊断表:

现象	可能原因	排查步骤
一直 blocked	arrive() 未被所有线程调用	检查指令调度逻辑
is_complete 总 false	arrived_mask 未递增	检查 lane_id 是否正确
PC 未更新	set_thread_pc() 未调用	检查 complete 后代码路径
唤醒后错误指令	pc_stack 未同步	检查 update_pc_stack()

场景 7.5: Divergent Execution 路径不同步 (高级)

触发条件:

test_nested_sync 或类似 divergent execution 测试卡住
日志显示不同的 lane 到达不同的 reconvergence_pc
arrived 计数停滞 (如 arrived=16/32)

根因分析:

PTX 语法: @%p1 bra $L__BB2_2 创建两条执行路径
Visitor 翻译: bar.sync → bar.warp.sync 时使用固定 reconvergence_pc = size + 1
问题: 当 divergent 后再次汇合时，不同路径的下一条指令位置不同

调试步骤:

# 1. 查看每个 lane 到达的 PC
grep "arrived at bar.warp.sync" /tmp/test_output.log | 
  awk -F' '"{print
### 场景 8: 数值异常

**触发条件**: 测试报告 Numerical Exception

**自动行动**:

1. **选择配置**: `configs/dev_debug_config.ini`
2. **运行分析**: `grep -E "Numerical|NaN|Inf" ptx_emu_debug.log`
3. **分析重点**: 哪个指令产生异常值，输入操作数是否合法

---

   - 内存访问模式
   - 缓存命中率
   - 分支预测

---

## 自动化调试流程

### 标准调试流程（自动触发）

接收用户问题 ↓
分析问题类型，选择调试配置 ↓
使用 debug-run.sh 选择配置：./scripts/debug-run.sh {配置名} ./build/bin/{程序} ↓
运行程序收集日志 ↓
分析日志文件，提取关键信息 ↓
定位问题根源 ↓
生成修复方案 ↓
验证修复


### 快捷脚本使用

**技能会自动使用或指导用户使用快捷脚本**:

```bash
# 技能自动执行
./scripts/debug-run.sh {配置名} ./build/bin/{程序}

# 示例
./scripts/debug-run.sh memory ./build/bin/dummy-args
./scripts/debug-run.sh trace ./build/bin/RAY 512 512

日志分析技术

常用 grep 命令

# 查看错误
grep "ERROR\|FATAL" ptx_emu_*.log

# 查看特定组件
grep "\[mem\]\|\[instr\]\|\[exec\]" ptx_emu_*.log

# 查看特定指令
grep "st.global\|ld.global" ptx_emu_*.log

# 查看寄存器
grep "Register.*contains\|Commit.*%" ptx_emu_*.log

# 查看内存操作
grep "Memory allocated\|Memory freed" ptx_emu_*.log

# 时间线分析
grep "\[CLK:" ptx_emu_*.log | head -100

日志分析模式

模式 1: 崩溃分析

# 找到崩溃前的最后状态
tac ptx_emu_trace.log | grep -A 50 "最后的消息"

模式 2: 内存泄漏

# 对比分配和释放
grep "Memory allocated" ptx_emu.log > alloc.txt
grep "Memory freed" ptx_emu.log > free.txt
diff alloc.txt free.txt

模式 3: 指令追踪

# 查看指令执行序列
grep "PC\[" ptx_emu_instr_debug.log | awk '{print $5}'

调试输出位置

日志文件

配置文件	日志文件	用途
`configs/dev_debug_config.ini`	`ptx_emu_debug.log`	日常调试
`configs/verbose_trace_config.ini`	`ptx_emu_trace.log`	详细跟踪
`configs/memory_debug_config.ini`	`ptx_emu_memory_debug.log`	内存调试
`configs/instruction_debug_config.ini`	`ptx_emu_instr_debug.log`	指令调试

快速查看

# 实时查看
tail -f ptx_emu_*.log

# 查看最新 N 行
tail -N ptx_emu_debug.log

# 分页查看
less -R ptx_emu_trace.log

# 搜索查看
grep "关键词" ptx_emu_debug.log | less

与其他技能协作

依赖技能

cpp-debug: 提供通用 C++ 调试方法
systematic-debugging: 提供系统化调试流程
cuda-ptx: 提供 CUDA/PTX 专业知识

⚠️ 与 `ptx-grammar-modification` 的分工

重要: 本项目有专门的 PTX 语法修复技能 ptx-grammar-modification（位于 docs/skills/）。

职责划分:

问题类型	使用技能	说明
ANTLR 解析错误	`ptx-grammar-modification`	语法文件修改、重新生成解析器
语法文件修改	`ptx-grammar-modification`	修改 `.g4` 文件
`no viable alternative`	`ptx-grammar-modification`	语法规则修复
`mismatched input`	`ptx-grammar-modification`	Token 定义修复
测试失败	`ptx-debug`	运行时错误、逻辑错误
内存问题	`ptx-debug`	内存访问、越界
指令错误	`ptx-debug`	指令执行结果不对
程序崩溃	`ptx-debug`	segfault、core dump
性能问题	`ptx-debug`	性能优化、benchmark

自动识别逻辑:

当用户提到 "PTX 解析错误" 时:
  1. ptx-debug 技能触发
  2. 检查错误类型
  3. IF 错误是 ANTLR 解析错误 (no viable alternative, mismatched input):
       → 🛑 停止调试流程
       → 建议用户使用 ptx-grammar-modification 技能
       → 提示运行 ./tests/ptx/test_all_ptx.sh
  4. ELSE 错误是运行时错误:
       → 继续调试流程
       → 选择 verbose_trace 配置
       → 收集日志分析

协作场景

内存泄漏: 使用 cpp-debug 的 valgrind/ASan 方法
解析错误:
- 语法问题 → ptx-grammar-modification
- 日志分析 → ptx-debug 辅助
系统调试: 遵循 systematic-debugging 的四阶段流程
PTX ISA 查询: 使用 cuda-ptx 的本地文档

验证标准

调试完成的标志:

示例使用

示例 1: 解析错误

用户: "PTX 解析失败，报告语法错误"

技能自动行动:

识别为"PTX 解析错误"
选择 verbose_trace_config.ini

执行：

./scripts/debug-run.sh verbose ./build/bin/dummy-args
grep "parser\|lexer" ptx_emu_trace.log

分析错误位置
修复语法文件
验证修复

示例 2: 测试失败

用户: "test_memory_manager 测试失败"

技能自动行动:

识别为"测试失败"
选择 debug_config.ini

执行：

./scripts/debug-run.sh debug ./build/bin/测试程序
cd build && ctest -R test_memory_manager -V
tail -100 ptx_emu_debug.log

分析失败原因
生成修复

示例 3: 性能优化

用户: "RAY benchmark 运行太慢"

技能自动行动:

识别为"性能问题"
选择 perf_config.ini

执行：

./scripts/debug-run.sh perf ./build/bin/RAY 512 512

分析热点
提出优化建议

诊断方法分类

核心原则: 80% 的问题可以通过不需要仿真的快速诊断定位，只在必要时运行完整仿真

快速诊断 (<30s) - 不需要仿真

适用场景:

PTX 语法解析错误
Label 注册问题
CFG 分析警告
指令翻译问题 (如 bar.sync → bar.warp.sync)

验证方法:

方法	命令	耗时	适用
test-ptx 解析	`./build/bin/test-ptx tests/ptx/test.ptx`	<1s	语法、翻译验证
手动提取 PTX	`cuobjdump -ptx binary > test.ptx`	<5s	检查生成的 PTX
最小化 PTX	手动创建简化 PTX 文件	<10s	隔离特定指令
CFG 日志	添加日志到 cfg_builder.cpp	<30s	控制流分析

快速验证流程:

# 1. 从二进制提取 PTX
/usr/local/cuda/bin/cuobjdump -ptx ./build/bin/test_xxx 2>/dev/null | tail -n+12 > /tmp/test.ptx

# 2. 用 test-ptx 验证解析
./build/bin/test-ptx /tmp/test.ptx 2>&1 | grep -E "bar|bra|label|PASS|FAIL"

# 3. 查看特定指令翻译
./build/bin/test-ptx /tmp/test.ptx 2>&1 | grep -A2 "bar.warp.sync"

成功案例:

✅ PTX label 语法修复：通过 test-ptx 快速验证 33/33 语法测试
✅ bra 指令 predicate: 通过解析输出确认 target 正确注册
✅ CFG 空 target 警告：通过添加日志快速定位并修复

仿真诊断 (>30s) - 需要完整运行

适用场景:

Barrier 同步卡住
线程调度问题
内存访问错误
寄存器值错误
numerical exception

验证方法:

方法	命令	耗时	适用
完整单元测试	`ctest -R test_xxx -V`	30-120s	功能回归
日志分析	启用 PTX_DEBUG	30-60s	执行跟踪
GDB 调试	`gdb --args ./bin/test_xxx`	5-10min	崩溃定位
ASan/UBSan	Debug 构建 + sanitizer	2-5min	内存错误

调试流程:

# 1. 启用详细日志
export PTX_LOG_LEVEL=debug
./build/bin/test_xxx 2>&1 | grep -E "bar|sync|blocked|complete"

# 2. 定位卡住点
./build/bin/test_xxx 2>&1 | tail -50

# 3. GDB 调试 (如果崩溃)
gdb --args ./build/bin/test_xxx
(gdb) run
(gdb) bt  # 崩溃时获取堆栈

决策树:

遇到问题
    │
    ├─ 解析错误/语法问题？
    │   └─→ 使用 test-ptx (1s) ← 优先
    │
    ├─ Label/CFG 警告？
    │   └─→ 添加日志，重解析 (10s) ← 优先
    │
    ├─ 指令执行卡住？
    │   └─→ 启用调试日志 (30s)
    │       └─ 仍无法定位？
    │           └─→ GDB 调试 (5min)
    │
    └─ 崩溃/SegFault？
        └─→ GDB + ASan (2min)

Barrier 调试最佳实践

屏障调试流程 (类型 2 - 需要仿真)

# 1. 验证 PTX 解析 (类型 1 - 快速)
./build/bin/test-ptx test.ptx | grep "bar.warp.sync"

# 2. 启用 barrier 日志
grep -r "PTX_DEBUG" src/ptxsim/instructions/barrier.cpp
# 确保相关日志启用

# 3. 运行测试并抓取日志
./build/bin/test_xxx 2>&1 | grep -E "arrived|complete|blocked|released"

# 4. 分析关键指标
# - participation_mask 是否正确？
# - arrived_mask 是否递增？
# - is_complete() 何时返回 true？
# - set_thread_pc() 是否被调用？
# - update_pc_stack() 是否同步？

关键日志模式:

bar.warp.sync: Initialized wbar[0] with mask=0xFFFFFFFF, reconvergence_pc=N
Lane X arrived at bar.warp.sync (mask=0xFFFFFFFF, pc=N)
bar.warp.sync: Barrier complete, releasing N threads to PC=N
Lane X blocked at bar.warp.sync (arrived=N/32)

预期行为:

初始化 wbar，participation_mask = 0xFFFFFFFF
每个 lane 调用 arrive()，arrived_mask 递增
当 arrived_mask == participation_mask 时，barrier complete
所有参与线程 PC 更新为 reconvergence_pc
is_blocked 清除，status 设为 Active

常见问题诊断表:

现象	可能原因	排查步骤
一直 blocked	arrive() 未被调用	检查指令调度器是否调度过 Blocked 线程
is_complete 总 false	arrived_mask 未递增	检查 lane_id 是否正确传入 arrive()
PC 未更新	set_thread_pc() 未被调用	检查 is_complete() 返回后的代码路径
唤醒后执行错误指令	pc_stack 未同步	检查 update_pc_stack() 调用点
participation_mask=0	指令翻译错误	检查 visitor 中 barrier operand 提取

关键代码位置:

src/ptx_parser/ptx_visitor_barrier.cpp  - bar.sync → bar.warp.sync 翻译
src/ptxsim/instructions/barrier.cpp     - Wbar 实现与 arrive()/is_complete()
include/ptxsim/wbar.h                    - Wbar 数据结构定义

自愈型 Bug 检测（重要）

问题：update_active_mask() 自愈

execute_warp_instruction() 末尾调用 update_active_mask()，从 warp_state.threads[i].is_active 重新计算 active_mask[]。这意味着：

一条指令内临时破坏 active_mask[]（如 barrier set_active_mask(arrived_mask) 覆写）
下一条指令时 update_active_mask() 自动从线程状态恢复
端到端集成测试只看 is_finished() → 可能 PASS 即使有 bug

检测方法

单元测试必须：在特定调用后直接断言 get_active_mask()，不要只检查 is_finished()
集成测试补充：检查所有 lane 的最终 PC — 如果某些 lane 停在 barrier PC 而非 post-barrier PC，说明有 lane 丢失
生产症状检查：simpleGEMM-* 不再卡住 = fix 生效；但结果是否正确 = divergence reconciliation 是否完整（独立问题）

诊断命令

# 1. 单元测试 — 验证 set_active_mask 行为（RED guard）
ctest -R "post_barrier_two_halves" -V

# 2. 集成测试 — 仅看 warp 是否完成（smoke test，可能被骗过）
ctest -R "integration_post_barrier_two_halves" -V

# 3. 真实症状 — 看 lane 是否全部推进
for t in int float double; do
  timeout 5 ./build/bin/simpleGEMM-$t 2>&1 | tail -1
done

典型症状对照表

检查	含义
`ctest -R "post_barrier"` RED	bug 存在 — set_active_mask 行为不对
`ctest -R "integration_post_barrier"` PASS 但 simpleGEMM 仍输出 0	自愈型 bug — warp 完成但数据错误（divergence reconciliation 独立 bug）
`simpleGEMM` 卡住	bug 存在 — active_mask 被破坏后没有 lane 可调度
`simpleGEMM` 退出但结果非 0	fix 生效，divergence reconciliation 正常

经验教训

日志文件大小: trace 级别可能产生数百 MB 日志
性能影响: 详细日志会降低 10-100 倍性能
及时清理: 定期清理旧日志文件
配置恢复: 调试后记得恢复到 release 配置

更新记录

2026-03-23: 初始版本，包含 6 种调试场景的自动化方法
后续更新：根据新场景添加调试配置和方法

场景 7.5: Divergent Execution 路径不同步 (高级)

触发条件:

test_nested_sync 或类似 divergent execution 测试卡住
日志显示不同的 lane 到达不同的 reconvergence_pc
arrived 计数停滞 (如 arrived=16/32)

根因分析:

PTX 语法: @%p1 bra $L__BB2_2 创建两条执行路径
Visitor 翻译: bar.sync → bar.warp.sync 时使用固定 reconvergence_pc = size + 1
问题: 当 divergent 后再次汇合时，不同路径的下一条指令位置不同

调试步骤:

# 1. 查看每个 lane 到达的 PC
grep "arrived at bar.warp.sync" /tmp/test_output.log | \
  grep -oP 'pc=\K\d+' | sort -n | uniq -c

# 2. 检查是否所有 lane 都到达同一个 barrier
grep "arrived=" /tmp/test_output.log | \
  grep -oP 'arrived=\K\d+/\d+' | sort | uniq -c

# 3. 对比 label 注册位置
grep "Registering label" /tmp/test_output.log

实际案例 (test_nested_sync):

Lane 13 arrived pc=26  # 在分支路径中
Lane 14 arrived pc=12  # 直接跳转路径
arrived=16/32          # 停滞

问题根因: threads 16-31 跳转到 L__BB2_2 (PC=23)，然后执行第二个 barrier (PC=25)，但实际到达的 reconvergence_pc=26 不正确。

修复建议:

检查 visitor 中的 barrier PC 计算逻辑
确保 divergent 路径正确汇合到同一个 label
使用 CFG analysis 验证 reconvergence 点
查看 ptx_visitor_barrier.cpp 中 next_pc 的计算

场景 9: S_BAR 多 warp CTA 级 barrier 死锁 (新增 2026-07-01)

触发条件:

CTA 级 barrier 测试（integration_cta_barrier_*, integration_cute_rmsnorm_*）卡住或 FAIL
日志显示线程 blocked at barrier 但 barrier 永不完成

调试方法:

# 1. 启用线程追踪（替代全局 fprintf，受 debug config 控制）
# configs/debug_config.ini 中 component.thread=debug
./build/bin/test_xxx 2>&1 | grep "\[thread\]" | grep -E "bar.sync|barrier|blocked|released" | head -30

# 2. 检查 CTABarrier 到达情况
./build/bin/test_xxx 2>&1 | grep -E "arrive|complete|uninitialized"

# 3. 检查调度器是否能看到释放的线程
./build/bin/test_xxx 2>&1 | grep -E "lanes_by_pc|active_mask" | head -20

常见问题诊断表:

现象	可能原因	排查步骤
`arrive called on uninitialized barrier`	`CTABarrier::reset()` 清除了 `is_initialized_`	检查 reset() 是否保持初始化状态
barrier complete 但 warp 无法前进	`is_active` 未被 barrier release 恢复	检查 `release_cta_barrier` 中的 `ts.is_active` 设置
warp 跳过 barrier handler	测试 driver 提前返回 `post_barrier_pc`	添加 `is_cta_barrier_complete()` 门控
divergence + bar.sync 卡住	`reconvergence_pc` 指向错误位置	检查 `make_bra_pred` 的 reconv 参数是否等于真实汇聚点

关键代码位置:

src/ptxsim/barrier/barrier_module.cpp  - release_cta_barrier (is_active 恢复)
src/ptxsim/barrier/cta_barrier.cpp     - reset() (保留初始化状态)
include/ptxsim/testing/scheduler_utils.h - step_warp (PC 边界检查)

调试命令速查:

# 打印 per-lane 指令执行序列
./build/bin/test_xxx 2>&1 | grep "\[thread\]" | grep -oP 'lane=\d+ pc=\d+' | sort | uniq -c

# 打印所有 barrier 操作时间线
./build/bin/test_xxx 2>&1 | grep -E "\[CLK:\d+.*arrive|complete|release|reset" | awk '{print $1,$2,$NF}'

ptx-debug

More from this repository

More from this repository

PTX-EMU 调试技能

核心原则

调试配置自动选择

配置选择矩阵

自动选择流程

场景化调试方法

场景 1: PTX 解析错误

场景 2: 测试失败

场景 4: 内存问题

场景 5: 指令执行错误

场景 6: 性能问题

场景 7: Barrier 同步卡住 (SIMT 特定)

场景 7.5: Divergent Execution 路径不同步 (高级)

日志分析技术

常用 grep 命令

日志分析模式

调试输出位置

日志文件

快速查看

与其他技能协作

依赖技能

⚠️ 与 ptx-grammar-modification 的分工

协作场景

验证标准

示例使用

示例 1: 解析错误

示例 2: 测试失败

示例 3: 性能优化

诊断方法分类

快速诊断 (<30s) - 不需要仿真

仿真诊断 (>30s) - 需要完整运行

Barrier 调试最佳实践

屏障调试流程 (类型 2 - 需要仿真)

自愈型 Bug 检测（重要）

问题：update_active_mask() 自愈

检测方法

诊断命令

典型症状对照表

经验教训

更新记录

场景 7.5: Divergent Execution 路径不同步 (高级)

场景 9: S_BAR 多 warp CTA 级 barrier 死锁 (新增 2026-07-01)

PTX-EMU 调试技能

核心原则

调试配置自动选择

配置选择矩阵

自动选择流程

场景化调试方法

场景 1: PTX 解析错误

场景 2: 测试失败

场景 4: 内存问题

场景 5: 指令执行错误

场景 6: 性能问题

场景 7: Barrier 同步卡住 (SIMT 特定)

场景 7.5: Divergent Execution 路径不同步 (高级)

日志分析技术

常用 grep 命令

日志分析模式

调试输出位置

日志文件

快速查看

与其他技能协作

依赖技能

⚠️ 与 ptx-grammar-modification 的分工

协作场景

验证标准

示例使用

示例 1: 解析错误

示例 2: 测试失败

示例 3: 性能优化

诊断方法分类

快速诊断 (<30s) - 不需要仿真

仿真诊断 (>30s) - 需要完整运行

Barrier 调试最佳实践

屏障调试流程 (类型 2 - 需要仿真)

自愈型 Bug 检测（重要）

问题：update_active_mask() 自愈

⚠️ 与 `ptx-grammar-modification` 的分工

⚠️ 与 `ptx-grammar-modification` 的分工