ワンクリックでManusで任意のスキルを実行

始める

ptx-debug

スター3

フォーク0

更新日2026年7月1日 08:49

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

インストール

Codex または Claude でインストールこの Prompt をコピーして Codex、Claude、または他のアシスタントに貼り付けると、Skill ページを確認してインストールできます。

Manusで実行

ソース

chisuhua

chisuhua/PTX-EMU

GitHub リポジトリを開く Creator のリポジトリを見る

ダウンロード

Manusで実行

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

このリポジトリの他の Skills

このリポジトリの他の Skills

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` 的分工