[mark14wu]

Summary

Under every client, and/or on tensors follow Python semantics instead of Triton's:

in_range = (offs >= lo) and (offs < hi)   # TileLens: offs < hi
x = tl.load(x_ptr + offs - lo, mask=in_range)

Triton's interpreter, which TileLens builds on, runs kernel code as Python and treats any tensor with more than one element as true, so a and b returns b and a or b returns a. Compiled Triton short-circuits only on compile-time values and combines runtime operands with logical_and/logical_or. In the kernel above, the sanitizer reports the load of lanes offs < lo as out of bounds, though the compiled kernel masks them off, and the tracer computes different values. Many kernels in the wild write masks this way; on TritonBench, with the other open sanitizer fixes applied, it causes all the remaining false positives of the sanitizer.

triton-lang/triton#12150 fixes this in the interpreter; this PR applies the same rewrite under TileLens for Tritons without it (3.8.0 included):

• _visit_triton_function_def runs _BoolOpRewriter after Triton's own AST rewrite. It turns each a and b into the native chain op(True) < op(True, a) < op(True, b) over _BoolOpChain objects, so operands are still evaluated once, left to right and in their own scope, and the chain stops only at a compile-time value that decides the result; runtime tensors are combined with logical_and/logical_or.
• It runs after Triton's rewrite because the interpreter's visit_Assign does not visit assigned values (so a visit_BoolOp on Triton's transformer never sees x = a and b), and so that a Triton visit_Compare cannot rewrite the chains again. If Triton's ASTTransformer defines visit_BoolOp, TileLens leaves BoolOps to it.
• _tilelens_bool_op is bound like _tilelens_loop_iter_wrapper, as a hidden keyword-only default, so no name is added to the user's globals; _hidden_kernel_args lists both for the function-def rewrite, _compile_and_exec, and _prepare_grid_launch.

As for loops, a device function that is called as a plain JITFunction (not traced, not interpreted) is not rewritten.

Tests

• New in tests/end_to_end/test_tracer.py: test_and_or_on_tensors_follow_compiled_triton, 14 cases (and, or, chains, nesting, and constexpr operands, on vectors and scalars) that store the result under the tracer and compare it with the logical operators. The 5 vector cases fail on main. Native runs of the same kernel on an RTX 4090 give the expected values with Triton 3.8.0 and 3.7.1.
• New in tests/end_to_end/test_sanitizer.py: test_sanitizer_and_or_on_tensors_follow_compiled_triton, the kernel above. It fails on main with a spurious out-of-bounds report on x_ptr.
• pytest tests/: 366 passed, 6 skipped on main; 381 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.100s	0.099s	-0.9%	20 / 20
gemm_oob	0.112s	0.111s	-0.6%	20 / 20
indirect_load	0.021s	0.021s	-0.2%	20 / 20
nested_loop	0.222s	0.219s	-1.1%	20 / 20
block_pointer_loop_advance	0.120s	0.118s	-1.5%	20 / 20
liger_jsd	0.133s	0.131s	-0.9%	20 / 20
flaggems_layernorm	0.375s	0.374s	-0.4%	20 / 20
swiglu	0.162s	0.161s	-0.7%	20 / 20
cross_entropy	0.929s	0.925s	-0.4%	20 / 20
fused_linear_jsd	0.201s	0.200s	-0.5%	20 / 20
Total	2.374s	2.359s	-0.6%	N/A

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