[mark14wu]

Summary

The sanitizer checks an access against the tensor its pointer comes from, so that an access that leaves tensor A and lands inside tensor B is reported. When the offsets are loaded from memory, it does not:

idx = tl.load(idx_ptr + offs)      # offsets from a_ptr, loaded from memory
x = tl.load(a_ptr + idx)           # every address falls inside b: not reported

An address computed from loaded values is concretized before it is checked. _materialize_memory_operand first replaces the loaded subtree with its values, which keeps a_ptr + <values>, and then, when the pointer holds a vector constant, replaces the whole pointer with its concrete addresses, which drops a_ptr. _resolve_tensor then took the tensor that the first concrete address falls in (_collect_tensor_base returns the address list). With offsets into b:

lanes of a_ptr + idx	main	this PR
all in a	no report	no report
all in b	no report	reported on a_ptr
lane 0 in a, the rest in b	reported on a_ptr	reported on a_ptr
lane 0 in b, the rest in a	reported on b_ptr	reported on a_ptr
two in b, two past b	reported on b_ptr (the lanes in b pass)	reported on a_ptr

Loads and stores behave alike.

This PR keeps the pointer argument:

• SymbolicClient._pointer_base follows only the pointer-valued operands of an address (pointer arithmetic, splats and broadcasts, selects), skipping type-spec children, down to a scalar pointer constant. A pointer that comes out of a load, or is cast from an integer, has no such base and returns None.
• Before the full replacement, _materialize_memory_operand takes the address's base and stores it as origin_base on every pointer node on the way (_mark_pointer_base), and replace_subtree carries origin_base to the constant that replaces a node. replace_subtree() works in place, so the subtrees it replaces may be shared with later accesses: with a pointer tile built before a loop and offsets loaded in the loop, the first iteration's concretization would otherwise leave the later iterations without a base (the second commit; the loop test below fails without it).
• _collect_tensor_base returns origin_base when a concrete pointer has one, so the concrete-address route, the deferred loop checks, and reporting all use a_ptr's tensor. signature() includes origin_base, so that two concrete pointers with the same addresses but different bases are not deduplicated into one check.

A pointer loaded from a pointer table is still resolved from its value, as before. A table entry that points to a tensor that is not a kernel argument is still reported, since the sanitizer knows no bounds for it; this PR does not change that.

Tests

• New in tests/end_to_end/test_sanitizer.py:
  • test_sanitizer_checks_loaded_offsets_against_their_base_tensor: the five cases above, for a load and a store (10 cases). The kernels are defined per test, so that no launch repeats an earlier one ([BUG] Sanitizer skips re-launches whose addresses, masks, or branches depend on loaded values #517 is not merged on main). The 6 cases that should be reported on a_ptr but are not fail on main.
  • test_sanitizer_checks_loaded_offsets_in_a_loop_against_their_base_tensor: a pointer tile built before a loop, offsets loaded in it; the second iteration lands entirely in b. Not reported on main, nor with only the first commit.
  • test_sanitizer_resolves_a_loaded_pointer_from_its_value: a pointer loaded from a table into a tensor that is a kernel argument, in bounds; not reported, with and without this PR.
• pytest tests/: 366 passed, 6 skipped on main; 378 passed, 6 skipped with this PR (Python 3.12, Triton 3.8.0).
• ruff check and ruff format --check (v0.1.9) pass on the changed files.

🤖 Generated with Claude Code

[github-actions]

Performance Benchmark

Benchmark	main (min)	PR (min)	Change	Samples
gemm	0.087s	0.087s	+0.0%	20 / 20
gemm_oob	0.097s	0.098s	+0.6%	20 / 20
indirect_load	0.018s	0.018s	+1.2%	20 / 20
nested_loop	0.181s	0.182s	+0.7%	20 / 20
block_pointer_loop_advance	0.098s	0.099s	+0.5%	20 / 20
liger_jsd	0.121s	0.122s	+0.5%	20 / 20
flaggems_layernorm	0.316s	0.315s	-0.3%	20 / 20
swiglu	0.146s	0.146s	+0.5%	20 / 20
cross_entropy	0.815s	0.824s	+1.1%	20 / 20
fused_linear_jsd	0.185s	0.186s	+0.4%	20 / 20
Total	2.065s	2.077s	+0.6%	N/A

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