Skip to content

[FIX] Resolve the base tensor of 2-D pointer tiles built from loaded indices - #516

Open
mark14wu wants to merge 1 commit into
mainfrom
fix/sanitizer-nested-pointer-base
Open

mark14wu wants to merge 1 commit into
mainfrom
fix/sanitizer-nested-pointer-base

Conversation

@mark14wu

Copy link
Copy Markdown
Collaborator

Summary

When the pointer of a load or store depends on loaded values, the sanitizer concretizes it. For a 2-D pointer tile, such as

rows = tl.load(idx_ptr + start_n + offs_n)
v = tl.load(v_ptr + rows[:, None] * D + offs_d[None, :])

_collect_tensor_base returns a nested list of addresses. SymbolicClient._resolve_tensor handled only a flat list, so the sanitizer crashed with

TypeError: int() argument must be a string, a bytes-like object or a real number, not 'list'

This PR flattens nested candidates before they are matched against the registered tensors. The flat-list and scalar cases are unchanged.

Tests

  • New test_sanitizer_resolves_2d_pointer_tile_from_loaded_indices in tests/end_to_end/test_sanitizer.py, with an in-bounds case (no record) and a case with one loaded row index past the end of v (record reported). Both cases fail with the TypeError above without the fix. Because both cases launch the same kernel with the same signature, the test clears _fn_symbolic_cache_set so that the second case is analyzed rather than replayed.
  • pytest tests/: 368 passed, 6 skipped (Python 3.12, Triton 3.8.0, tilelens installed in editable mode with its CLIs on PATH).
  • ruff check and ruff format --check (v0.1.9, as pinned in .pre-commit-config.yaml) pass on both changed files.

🤖 Generated with Claude Code

https://claude.ai/code/session_016HMpCLXdUVNfMbJTj8aHE2

…indices

When a pointer tile depends on loaded values, e.g.
`tl.load(v_ptr + rows[:, None] * D + offs_d[None, :])` with `rows`
loaded from memory, the sanitizer concretizes it and
`_collect_tensor_base` returns a nested list of addresses.
`_resolve_tensor` only handled a flat list and crashed with
`TypeError: int() argument must be ... not 'list'`. Flatten nested
candidates before matching them against registered tensors.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_016HMpCLXdUVNfMbJTj8aHE2
@github-actions

Copy link
Copy Markdown

Performance Benchmark

Benchmark main (min) PR (min) Change Samples
gemm N/A N/A N/A 0 / 0
gemm_oob N/A N/A N/A 0 / 0
indirect_load N/A N/A N/A 0 / 0
nested_loop N/A N/A N/A 0 / 0
block_pointer_loop_advance N/A N/A N/A 0 / 0
liger_jsd N/A N/A N/A 0 / 0
flaggems_layernorm N/A N/A N/A 0 / 0
swiglu 0.162s 0.162s -0.0% 20 / 20
cross_entropy N/A N/A N/A 0 / 0
fused_linear_jsd N/A N/A N/A 0 / 0
Total N/A N/A N/A N/A

Iterations: 1 warmup + 20 measured
Samples are shown as main / PR; long pytest benchmarks may use fewer samples.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant