[mark14wu]

Summary

When the sanitizer concretizes a data-dependent value, it uses the running block's values: _materialize_memory_operand replaces an operand that contains a load with constants, and _materialize_loop_value does the same for a loop bound. _on_data_dependent_value then sets need_full_grid, so every block of the launch runs. The checks, however, still add only pid_ok, which lets pid_0..2 range over the whole grid. One block's concretized values were therefore combined with every other block's addresses, which gives false positives such as:

• A bounds mask AND-ed with a condition on loaded values. mask = (offs < n) & (flag == 0) is concretized as a whole. In block 0, offs < n is true in every lane, so the concretized mask is all ones, and the access p + offs was checked for the last block without its bounds mask.
• A loop bound loaded per block. In for j in range(0, tl.load(lens_ptr + pid)), block 0's bound was applied to every block.

This PR pins pid_0..2 to the running block (grid_idx) in _check_range_satisfiable once the block has concretized a value (need_full_grid is set). Every block of such a launch runs and checks its own accesses, so nothing is skipped. Outside loops, the check cache also keys on the block in that case. Otherwise a later block that builds the same concretized expression would be skipped as a repeat of block 0's check.

Checks made before a block concretizes anything are unchanged: they involve no concretized value and still cover every program id.

Tests

• New in tests/end_to_end/test_sanitizer.py:
  • test_data_dependent_mask_is_checked_per_block: the masked-add pattern above, three blocks, with the last block partly masked off. It fails on main and passes with this PR.
  • test_loop_bound_loaded_per_block[fits]: per-block loop bounds [4, 3] over a buffer that holds 7 rows. It fails on main and passes with this PR.
  • test_loop_bound_loaded_per_block[overruns]: bounds [3, 4], where block 1 reads row 7. It is reported on main and with this PR.
  • test_data_dependent_mask_without_bounds_reports_last_block: the masked-add kernel without the bounds mask, so the last block reads past p. Blocks 0-2 build the same concretized mask, so this test fails if the cache does not key on the block.
• pytest tests/: 366 passed, 6 skipped on main; 370 passed, 6 skipped with this PR (Python 3.12, Triton 3.8.0).
• ruff check and ruff format --check (v0.1.9, as pinned in .pre-commit-config.yaml) pass on both changed files. mypy (v1.8.0) reports no new errors in sanitizer.py.
• CI: the build job installs Triton 3.9.0 and then Triton main. With those, about 96 tests of the existing suite fail whatever the change (the run on [FIX] Stop the sanitizer from skipping repeated launches #518 also has 96 failures), mostly with 'SymbolicScalarDType' object has no attribute 'scalar' or 'SymbolicExprDataWrapper' object has no attribute 'astype'. The new end-to-end tests fail there with the same errors as the existing sanitizer tests; the new unit tests pass.

🤖 Generated with Claude Code

https://claude.ai/code/session_016HMpCLXdUVNfMbJTj8aHE2

[github-actions]

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.120s	0.121s	+0.5%	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.