[mark14wu]

Summary

With fake tensors (SANITIZER_ENABLE_FAKE_TENSOR=1), the profiler fails on two kinds of kernels.

Atomics. The profiler skips loads, stores, and dots, but tl.atomic_add / tl.atomic_cas and the other atomics still run in Triton's interpreter. A FakeTensor carries the original tensor's address but no storage, so the atomic dereferences that address on the host: for a GPU tensor the process segfaults (exit -11), and for a CPU tensor the caller's tensor is written.

@triton.jit
def kernel(x_ptr, counter_ptr, BLOCK: tl.constexpr):
    tl.load(x_ptr + tl.program_id(0) * BLOCK + tl.arange(0, BLOCK))
    tl.atomic_add(counter_ptr, 1.0)   # segfaults under tile-profiler with fake tensors on a GPU tensor

Reinterpreted tensors. triton.reinterpret(x, tl.float8e5) gives a TensorWrapper. The interpreter unwraps it, turns the base into a FakeTensor, and rewraps it; TensorWrapper.__init__ reads base.data, which FakeTensor does not have:

AttributeError: 'FakeTensor' object has no attribute 'data'
  triton/runtime/jit.py, in TensorWrapper.__init__: self.data = base.data

Changes

• Profiler.register_op_callback: under fake tensors (cfg.virtual_memory), AtomicRMW and AtomicCas get an overrider that returns zeros of the pointer's element type without touching memory, as the load overrider does. With real tensors atomics still run, so test_block_sampling, which counts executed blocks with tl.atomic_add, is unchanged.
• FakeTensor.data returns the FakeTensor itself.

Tests

tests/end_to_end/test_profiler.py: test_fake_tensor_atomics_do_not_touch_memory (a kernel with atomic_add and atomic_cas on CPU tensors under fake tensors; before the fix the caller's counter became 4.0) and test_fake_tensor_accepts_reinterpreted_tensor (a triton.reinterpreted float8 argument; before the fix AttributeError). Both fail without the fix. Full suite on Triton 3.8.0: 368 passed, 6 skipped (366 passed before the two new tests).

Found on TritonBench scripts that the profiler could not run with fake tensors: int4_matmul, int8_matmul_quantization, and bgmv_shrink_kernel (atomics) and f8_conversion_utils (reinterpret).

🤖 Generated with Claude Code

https://claude.ai/code/session_0169QRxKtsivkLDCS5f7rrsm

[github-actions]

Performance Benchmark

Benchmark	main (min)	PR (min)	Change	Samples
gemm	0.100s	0.101s	+0.9%	20 / 20
gemm_oob	0.112s	0.113s	+0.4%	20 / 20
indirect_load	0.021s	0.021s	-0.5%	20 / 20
nested_loop	0.224s	0.224s	-0.2%	20 / 20
block_pointer_loop_advance	0.120s	0.119s	-1.1%	20 / 20
liger_jsd	0.133s	0.133s	-0.2%	20 / 20
flaggems_layernorm	0.378s	0.376s	-0.6%	20 / 20
swiglu	0.163s	0.162s	-0.3%	20 / 20
cross_entropy	0.943s	0.932s	-1.2%	20 / 20
fused_linear_jsd	0.201s	0.201s	-0.1%	20 / 20
Total	2.395s	2.380s	-0.6%	N/A

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