oob-detection
Detect out-of-bounds memory accesses in CPU or GPU code using static interval analysis and runtime assertions/printfs. Use when investigating OOB, buffer overrun, invalid memory access, HIP/ROCm illegal address, CUDA illegal memory access, silent tensor corruption, or suspicious buffer_load/store address arithmetic.
来源信息
- 仓库
- ROCm/FlyDSL
- 最近来源活动
- 2026年7月10日 13:33
- 检测到的 SKILL.md 语言
- 英语
- 星标
- 283
- 分支
- 120
安装方式
默认使用会先检查来源的 Prompt；你也可以切换为直接命令，或下载本地副本。
检查来源文件
决定是否安装前，请先阅读 SKILL.md，以及 SkillsMP 当前展示的配套文件。
正在显示 SKILL.md
SKILL.md
来源说明 · 只读预览- name
- oob-detection
- description
- Detect out-of-bounds memory accesses in CPU or GPU code using static interval analysis and runtime assertions/printfs. Use when investigating OOB, buffer overrun, invalid memory access, HIP/ROCm illegal address, CUDA illegal memory access, silent tensor corruption, or suspicious buffer_load/store address arithmetic.
- allowed-tools
- Read Edit Bash Grep Glob Agent
# OOB Detection
Use this skill when a kernel, runtime, or host path may read or write outside
its intended buffer or logical tile. Prefer proving the address range first
instead of relying only on runtime failures.
## 1. Classify the OOB
Decide which boundary may be violated:
| Boundary | Meaning | Typical detector |
|---|---|---|
| Physical allocation OOB | Address leaves the allocated tensor/buffer | HIP illegal address, runtime failure |
| Logical object OOB | Address stays in allocation but crosses a row/head/tile | Static interval analysis, explicit runtime check |
| Lane/thread ownership OOB | Thread reads another lane's slot | Static interval analysis, debug printf/assert |
| LDS/shared-memory OOB | Address exceeds allocated shared-memory region | Static interval analysis, LDS index guard |
Physical tools usually miss logical OOB because the access can still be inside
the same allocation.
## 2. Static Interval Analysis
For each memory access, write the exact element range:
```text
start = base + offset_expression
end = start + vec_width - 1
legal = [object_base, object_base + object_extent - 1]
```
Then substitute known ranges:
- Thread/lane ids: `threadIdx.x`, `lane`, `lane16id`, `warp_id`, `rowid`
- Compile-time loops: `range_constexpr(N)` gives `i in [0, N-1]`
- Vector widths: `buffer_load(..., vec_width=W, dtype=T)` reads `W` elements of `T`
- Strides and shapes: tensor `.shape`, `.stride`, layout shape/stride, tile extents
- Masks/clamps: `select(valid, value, safe_value)` changes the range only if it dominates the load/store
If `max(end) > object_base + object_extent - 1`, the OOB is statically proven.
If `min(start) < object_base`, the lower-bound OOB is statically proven.
### FlyDSL Example
For a Q head with `HEAD_SIZE = 128`:
```text
q_elem = q_base + lane16id * 8
load_start = q_elem + qwi * 4
load_end = load_start + 3
lane16id in [0, 15], qwi in [0, 3]
max(load_end) = q_base + 15*8 + 3*4 + 3 = q_base + 135
legal head end = q_base + 127
```
This proves logical OOB for the head. If each lane owns only 8 elements, then
`qwi in [2, 3]` also proves lane-slot OOB for every lane.
## 3. Add Runtime Logical Checks
When static proof is not enough or the formula depends on runtime values, add a
temporary guard immediately before the load/store. In FlyDSL kernels, prefer a
small `printf` with the failing coordinates:
```python
load_start = q_elem + fx.Int32(qwi * 4)
load_end = load_start + fx.Int32(vec_width - 1)
legal_end = q_base + fx.Int32(HEAD_SIZE - 1)
if load_end > legal_end:
fx.printf(
"OOB q load: lane=%d qwi=%d q_base=%d load=[%d,%d] legal_end=%d\n",
lane16id,
fx.Int32(qwi),
q_base,
load_start,
load_end,
legal_end,
)
```
For stores, include the output tensor/tile coordinates and the flattened offset.
Keep debug prints narrow; too many lanes printing can hide the useful signal.
## 4. Fix Strategy
Prefer fixing the invariant, not masking the fault:
- If the loop trip count is wrong, reduce the loop or vector width.
- If the per-lane ownership changed, update the layout, LDS store, and reader together.
- If a boundary tile is partial, clamp or predicate before the load/store.
- If a descriptor/resource range is too large or offset overflows i32, chunk the buffer resource or widen arithmetic before truncation.
- After the fix, rerun both the focused failing test and one neighboring shape that exercises the boundary.
## 5. Report Format
When reporting an OOB investigation, include:
- Access expression and element units
- Proven or observed failing range
- Legal range and which boundary was violated
- Minimal fix and validation command
在 GitHub 查看