diff --git a/README.md b/README.md
index 69d2ede66..ff031dcc0 100644
--- a/README.md
+++ b/README.md
@@ -123,7 +123,7 @@ uv sync --extra test # tests but no NKI support
 ### Testing
 * To run core TileLens tests, run `pytest tests/`.
 * (if NKI installed) To run NKI-specific tests, run `pytest tests/ -m nki`.
-* To run cuTile-specific tests, run `pytest tests/unit/test_cutile.py tests/end_to_end/test_cutile.py tests/cutile/test_cutile_tracer.py`.
+* To run cuTile-specific tests, run `pytest tests/unit/test_cutile.py tests/end_to_end/test_cutile.py tests/cutile/`.
 * To run all tests (Triton + NKI), run `pytest tests/ -m ""`.
 * To run visualizer web UI tests, run `npm run test:frontend`.
 
@@ -187,7 +187,9 @@ simulation runtimes live under `tilelens/core/simulation/`.
 
 ### cuTile tracing
 
-cuTile kernels run on the CPU with NumPy arrays and support tracing and visualization:
+cuTile kernels run on the CPU with NumPy arrays or Torch tensors and support tracing
+and visualization. Like Triton's interpreter, device tensors are copied to the host
+for the run and copied back afterwards:
 
 ```sh
 uv sync --extra test
@@ -205,7 +207,10 @@ kernel[grid](lhs, rhs, result)
 ```
 
 Use contiguous NumPy arrays for visualization. cuTile analysis currently supports
-only the tracer client; numerical behavior may differ from GPU execution.
+the tracer and profiler clients; numerical behavior may differ from GPU execution.
+The profiler reports load/store bytes, out-of-bounds tile elements, and for-loop
+lengths, and `tile-profiler` runs `ct.launch` calls through the interpreter.
+It skips the buffer-load check, which only applies to Triton on AMD GPUs.
 
 ## Analysis Clients
 
diff --git a/tests/cutile/test_cutile_profiler.py b/tests/cutile/test_cutile_profiler.py
new file mode 100644
index 000000000..cfb9fa6ce
--- /dev/null
+++ b/tests/cutile/test_cutile_profiler.py
@@ -0,0 +1,488 @@
+"""Tests for running the profiler on cuTile kernels."""
+# ruff: noqa: E402
+
+import importlib.util
+import inspect
+import os
+from pathlib import Path
+import shutil
+import subprocess
+import textwrap
+
+import numpy as np
+import pytest
+import torch
+
+ct = pytest.importorskip("cuda.tile")
+
+import tilelens
+from tilelens.clients import Profiler
+from tilelens.clients.profiler.data import LoadStoreBytes
+from tilelens.core.config import config as cfg
+from tilelens.core.frontend.cutile import (
+    cutile_skip_store,
+    cutile_zero_dot,
+    cutile_zero_load,
+)
+from tilelens.core.simulation.cutile import Array, Tile, cutile_builder
+from tilelens.core.trace import launches
+from tilelens.wrapper import PROFILER_COMMAND, SANITIZER_COMMAND
+
+example_path = Path(__file__).resolve().parents[2] / "examples" / "cutile" / "matmul.py"
+spec = importlib.util.spec_from_file_location("cutile_matmul_example", example_path)
+assert spec is not None
+assert spec.loader is not None
+matmul_example = importlib.util.module_from_spec(spec)
+spec.loader.exec_module(matmul_example)
+matmul_kernel = matmul_example.matmul_kernel
+
+
+@pytest.fixture(autouse=True)
+def clear_traces():
+    tilelens.clear()
+    yield
+    tilelens.clear()
+
+
+@pytest.fixture
+def _isolate_profiler_cfg():
+    """Save and restore profiler-related config values around every test."""
+    saved = (
+        cfg.profiler_enable_block_sampling,
+        cfg.profiler_enable_load_store_skipping,
+        cfg.profiler_disable_buffer_load_check,
+    )
+    yield
+    (
+        cfg.profiler_enable_block_sampling,
+        cfg.profiler_enable_load_store_skipping,
+        cfg.profiler_disable_buffer_load_check,
+    ) = saved
+
+
+def copy_kernel(src, dst):
+    block = ct.bid(0)
+    tile = ct.load(src, index=(block,), shape=(4,), padding_mode=ct.PaddingMode.ZERO)
+    ct.store(dst, index=(block,), tile=tile)
+
+
+def copy_pair(src, dst, index):
+    for offset in range(2):
+        tile = ct.load(
+            src, index=(index + offset,), shape=(2,), padding_mode=ct.PaddingMode.ZERO
+        )
+        ct.store(dst, index=(index + offset,), tile=tile)
+
+
+def looped_copy_kernel(src, dst):
+    block = ct.bid(0)
+    for step in range(3):
+        copy_pair(src, dst, block * 6 + step * 2)
+    # Like Triton, loops over other iterables are reported without a length;
+    # comprehensions are not for statements.
+    for _ in tuple(range(4)):
+        pass
+    _ = [i for i in range(4)]
+
+
+def _bytes_records():
+    records = [r for r in launches[-1].records if isinstance(r, LoadStoreBytes)]
+    return {r.type: r for r in records}
+
+
+def _line_of(fn, text):
+    source, start = inspect.getsourcelines(fn)
+    return start + next(i for i, line in enumerate(source) if text in line)
+
+
+def test_profiler_counts_masked_load_store_bytes(_isolate_profiler_cfg):
+    cfg.profiler_enable_block_sampling = False
+    cfg.profiler_enable_load_store_skipping = False
+    src = np.arange(1, 11, dtype=np.float32)
+    dst = np.zeros(10, dtype=np.float32)
+    profiler = Profiler()
+    tilelens.trace(profiler, frontend="cutile")(copy_kernel)[(3,)](src, dst)
+
+    np.testing.assert_array_equal(dst, src)
+    records = _bytes_records()
+    for kind in ("load", "store"):
+        assert records[kind].total_bytes_attempted == 12 * 4
+        assert records[kind].total_bytes_true == 10 * 4
+
+    assert profiler.load_mask_total_count == 12
+    assert profiler.load_mask_false_count == 2
+    assert profiler.store_mask_false_count == 2
+    load_stats = [s for s in profiler.mask_op_stats if s.op_type == "load"]
+    assert load_stats and all("ct.load" in s.code_line for s in load_stats)
+
+
+def test_profiler_counts_indexed_access_bytes(_isolate_profiler_cfg):
+    cfg.profiler_enable_block_sampling = False
+    cfg.profiler_enable_load_store_skipping = False
+
+    @ct.kernel
+    def indexed_kernel(src, dst):
+        indices = ct.arange(4, dtype=ct.int32) - 1
+        values = ct.gather(src, indices, padding_value=0)
+        ct.scatter(dst, ct.arange(4, dtype=ct.int32), values)
+
+    src = np.array([10, 20, 30, 40], np.int64)
+    dst = np.zeros(4, np.int64)
+    tilelens.trace("profiler", frontend="cutile")(indexed_kernel)[(1,)](src, dst)
+
+    np.testing.assert_array_equal(dst, [0, 10, 20, 30])
+    records = _bytes_records()
+    assert records["load"].total_bytes_attempted == 4 * 8
+    assert records["load"].total_bytes_true == 3 * 8
+    assert records["store"].total_bytes_true == 4 * 8
+
+
+def test_profiler_skips_memory_and_matmul(_isolate_profiler_cfg):
+    cfg.profiler_enable_block_sampling = False
+    cfg.profiler_enable_load_store_skipping = True
+    lhs = np.ones((4, 8), dtype=np.float32)
+    rhs = np.ones((8, 16), dtype=np.float32)
+    result = np.full((4, 16), -1, dtype=np.float32)
+    traced = tilelens.trace("profiler", frontend="cutile")(matmul_kernel)
+    traced[(2, 2)](lhs, rhs, result)
+
+    # Skipped stores leave the output untouched, but bytes are still counted.
+    np.testing.assert_array_equal(result, -1)
+    records = _bytes_records()
+    assert records["load"].total_bytes_true == (4 * 8 + 8 * 16) * 2 * 4
+    assert records["store"].total_bytes_true == 4 * 16 * 4
+
+
+@pytest.mark.parametrize(
+    "shape, index, load_shape, kwargs",
+    [
+        ((10,), (2,), (4,), {}),
+        ((4, 6), (1, 1), (2, 4), {"order": "F"}),
+        ((4, 8), (1, 0), (2, 8), {"traversal_steps": (1, 8)}),
+    ],
+)
+def test_zero_load_matches_tiled_load(shape, index, load_shape, kwargs):
+    array = Array(np.arange(np.prod(shape), dtype=np.float16).reshape(shape))
+    expected = cutile_builder.load(
+        array, index, load_shape, padding_mode=ct.PaddingMode.ZERO, **kwargs
+    )
+    zeros = cutile_zero_load(
+        array, index, load_shape, padding_mode=ct.PaddingMode.ZERO, **kwargs
+    )
+    assert zeros.shape == expected.shape
+    assert zeros.dtype == expected.dtype
+    assert not zeros.data.any()
+
+
+def test_zero_load_matches_gather():
+    array = Array(np.arange(12, dtype=np.int32).reshape(3, 4))
+    indices = (Tile(np.array([[0], [2]], np.int32)), Tile(np.arange(4, dtype=np.int32)))
+    expected = cutile_builder.gather(array, indices)
+    zeros = cutile_zero_load(array, indices, None, mask=None, check_bounds=True)
+    assert zeros.shape == expected.shape == (2, 4)
+    assert zeros.dtype == expected.dtype
+    assert not zeros.data.any()
+    assert cutile_skip_store(array, indices, zeros, _indexed=True) is None
+    np.testing.assert_array_equal(array.data, np.arange(12).reshape(3, 4))
+
+
+@pytest.mark.parametrize(
+    "lhs, rhs",
+    [
+        (np.ones((2, 3), np.float32), np.ones((3, 4), np.float32)),
+        (np.ones((2, 2, 3), np.float16), np.ones((2, 3, 5), np.float16)),
+        (np.ones((2, 3), np.float16), np.ones((3, 4), np.float32)),
+    ],
+)
+def test_zero_dot_matches_matmul(lhs, rhs):
+    x, y = Tile(lhs), Tile(rhs)
+    expected = cutile_builder.matmul(x, y)
+    zeros = cutile_zero_dot(x, y)
+    assert zeros.shape == expected.shape
+    assert zeros.dtype == expected.dtype
+    assert not zeros.data.any()
+
+
+def test_zero_dot_matches_mma_accumulators():
+    x = Tile(np.ones((2, 3), np.float16))
+    y = Tile(np.ones((3, 4), np.float16))
+    acc = Tile(np.full((2, 4), 7, np.float32))
+    expected = cutile_builder.mma(x, y, acc)
+    for zeros in (cutile_zero_dot(x, y, acc), cutile_zero_dot(x, y, acc=acc)):
+        assert zeros.shape == expected.shape
+        assert zeros.dtype == expected.dtype
+        assert not zeros.data.any()
+
+    # mma_scaled passes (x, x_scale, y, y_scale, acc); only acc decides the result.
+    scale = Tile(np.ones((2, 1), np.float32))
+    for zeros in (
+        cutile_zero_dot(x, scale, y, scale, acc),
+        cutile_zero_dot(x, scale, y, scale, acc=acc),
+    ):
+        assert zeros.shape == acc.shape
+        assert zeros.dtype == ct.float32
+        assert not zeros.data.any()
+
+
+def test_profiler_reports_range_loops(_isolate_profiler_cfg):
+    cfg.profiler_enable_block_sampling = False
+    cfg.profiler_enable_load_store_skipping = False
+    src = np.arange(1, 13, dtype=np.float32)
+    dst = np.zeros(12, dtype=np.float32)
+    profiler = Profiler()
+    tilelens.trace(profiler, frontend="cutile")(looped_copy_kernel)[(1,)](src, dst)
+
+    np.testing.assert_array_equal(dst, src)
+    kernel_loop = _line_of(looped_copy_kernel, "range(3)")
+    helper_loop = _line_of(copy_pair, "range(2)")
+    tuple_loop = _line_of(looped_copy_kernel, "tuple(range(4))")
+    assert set(profiler.loop_info) == {kernel_loop, helper_loop, tuple_loop}
+    assert profiler.loop_info[kernel_loop].length == 3
+    assert profiler.loop_info[helper_loop].length == 2
+    assert profiler.loop_info[tuple_loop].length is None
+    assert profiler.loop_info[kernel_loop].range_type == "python_range"
+    assert profiler.loop_info[helper_loop].range_type == "python_range"
+    assert profiler.loop_info[tuple_loop].range_type == "unknown"
+    # Rewritten helpers are only swapped in during traced runs.
+    assert "_tilelens_loop_iter_wrapper" not in inspect.signature(copy_pair).parameters
+
+
+def test_profiler_reports_loops_in_closures_and_tile_functions(_isolate_profiler_cfg):
+    cfg.profiler_enable_block_sampling = False
+    cfg.profiler_enable_load_store_skipping = False
+    steps = 2
+
+    @ct.function
+    def copy_steps(src, dst):
+        for step in range(steps):
+            ct.store(dst, (step,), ct.load(src, (step,), (2,)))
+
+    # Reached only at call time, through cuTile's tile-function dispatch.
+    helpers = {"copy": copy_steps}
+
+    def kernel(src, dst):
+        helpers["copy"](src, dst)
+
+    src = np.arange(1, 5, dtype=np.float32)
+    dst = np.zeros(4, dtype=np.float32)
+    profiler = Profiler()
+    tilelens.trace(profiler, frontend="cutile")(kernel)[(1,)](src, dst)
+
+    np.testing.assert_array_equal(dst, src)
+    helper_loop = _line_of(copy_steps, "range(steps)")
+    assert set(profiler.loop_info) == {helper_loop}
+    assert profiler.loop_info[helper_loop].length == 2
+
+
+def test_profiler_runs_kernels_without_source(_isolate_profiler_cfg):
+    cfg.profiler_enable_block_sampling = False
+    cfg.profiler_enable_load_store_skipping = False
+    namespace = {"ct": ct}
+    exec(
+        "def kernel(src, dst):\n"
+        "    for i in range(2):\n"
+        "        ct.store(dst, (i,), ct.load(src, (i,), (2,)))\n",
+        namespace,
+    )
+    src = np.arange(1, 5, dtype=np.float32)
+    dst = np.zeros(4, dtype=np.float32)
+    profiler = Profiler()
+    tilelens.trace(profiler, frontend="cutile")(namespace["kernel"])[(1,)](src, dst)
+
+    # Without source there is nothing to rewrite, so loops go unreported.
+    np.testing.assert_array_equal(dst, src)
+    assert profiler.loop_info == {}
+
+
+def nested_copy_kernel(src, dst):
+    for row in range(2):
+        for col in range(1, 4):
+            tile = ct.load(src, index=(row, col - 1), shape=(2, 2))
+            ct.store(dst, index=(row, col - 1), tile=tile)
+
+
+def test_profiler_reports_nested_range_loops(_isolate_profiler_cfg):
+    cfg.profiler_enable_block_sampling = False
+    cfg.profiler_enable_load_store_skipping = False
+    src = np.arange(24, dtype=np.float32).reshape(4, 6)
+    dst = np.zeros((4, 6), dtype=np.float32)
+    profiler = Profiler()
+    tilelens.trace(profiler, frontend="cutile")(nested_copy_kernel)[(1,)](src, dst)
+
+    np.testing.assert_array_equal(dst, src)
+    outer_loop = _line_of(nested_copy_kernel, "range(2)")
+    inner_loop = _line_of(nested_copy_kernel, "range(1, 4)")
+    # The inner loop runs once per outer iteration but is one loop site.
+    assert set(profiler.loop_info) == {outer_loop, inner_loop}
+    assert profiler.loop_info[outer_loop].length == 2
+    assert profiler.loop_info[inner_loop].length == 3
+    # All 2 x 3 tile copies are counted.
+    assert profiler.load_mask_total_count == 6 * 4
+
+
+def test_profiler_block_sampling_skips_unsampled_blocks(_isolate_profiler_cfg):
+    cfg.profiler_enable_block_sampling = True
+    cfg.profiler_enable_load_store_skipping = False
+    src = np.arange(1, 17, dtype=np.float32)
+    dst = np.zeros(16, dtype=np.float32)
+    profiler = Profiler(k=2)
+    tilelens.trace(profiler, frontend="cutile")(copy_kernel)[(4,)](src, dst)
+
+    written_blocks = {
+        block for block in range(4) if np.any(dst[block * 4 : block * 4 + 4])
+    }
+    assert profiler.sampled_blocks is not None
+    assert written_blocks == {idx[0] for idx in profiler.sampled_blocks}
+    assert _bytes_records()["load"].total_bytes_true == 2 * 4 * 4
+
+
+CLI_COPY_KERNEL = """
+import cuda.tile as ct
+import numpy as np
+
+
+@ct.kernel
+def copy_kernel(src, dst):
+    for step in range(2):
+        tile = ct.load(
+            src, index=(step,), shape=(4,), padding_mode=ct.PaddingMode.ZERO
+        )
+        ct.store(dst, index=(step,), tile=tile)
+"""
+
+
+def _run_profiler_cli(tmp_path: Path, launch_code: str, command=PROFILER_COMMAND):
+    exe = shutil.which(command)
+    assert exe is not None, f"{command} command not found"
+    script = tmp_path / "cutile_program.py"
+    script.write_text(CLI_COPY_KERNEL + textwrap.dedent(launch_code), encoding="utf-8")
+    env = os.environ.copy()
+    env["PROFILER_ENABLE_LOAD_STORE_SKIPPING"] = "0"
+    return subprocess.run(
+        [exe, str(script)],
+        cwd=str(tmp_path),
+        env=env,
+        capture_output=True,
+        text=True,
+    )
+
+
+def test_tile_profiler_cli_runs_cutile_launches(tmp_path: Path):
+    proc = _run_profiler_cli(
+        tmp_path,
+        """
+        src = np.arange(1, 7, dtype=np.float32)
+        dst = np.zeros(6, dtype=np.float32)
+        ct.launch(None, (1,), copy_kernel, (src, dst))
+        print("DST", dst.tolist())
+        """,
+    )
+
+    assert proc.returncode == 0, proc.stderr
+    assert "DST [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]" in proc.stdout
+    assert "Total steps:    2" in proc.stdout
+    assert "Masked percentage:       25.00%" in proc.stdout
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_tile_profiler_cli_runs_device_tensor_launches(tmp_path: Path):
+    proc = _run_profiler_cli(
+        tmp_path,
+        """
+        import torch
+
+        src = torch.arange(1, 7, dtype=torch.float32, device="cuda")
+        dst = torch.zeros(6, dtype=torch.float32, device="cuda")
+        ct.launch(torch.cuda.current_stream(), (1,), copy_kernel, (src, dst))
+        print("DST", dst.tolist())
+        """,
+    )
+
+    assert proc.returncode == 0, proc.stderr
+    assert "DST [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]" in proc.stdout
+
+
+def test_tile_profiler_cli_rejects_interpreted_launches(tmp_path: Path):
+    proc = _run_profiler_cli(
+        tmp_path,
+        """
+        from tilelens.core.frontend.cutile import interpreted_launches
+
+        with interpreted_launches():
+            pass
+        """,
+    )
+
+    assert proc.returncode != 0
+    assert "cannot be used together with CLI wrapper" in proc.stderr
+
+
+def test_tile_profiler_cli_rejects_traced_kernels(tmp_path: Path):
+    proc = _run_profiler_cli(
+        tmp_path,
+        """
+        import tilelens
+
+        traced = tilelens.trace("tracer", frontend="cutile")(copy_kernel)
+        src = np.zeros(4, dtype=np.float32)
+        ct.launch(None, (1,), traced, (src, src))
+        """,
+    )
+
+    assert proc.returncode != 0
+    assert "cannot be used together with CLI wrapper" in proc.stderr
+
+
+def test_tile_profiler_cli_accumulates_across_launches(tmp_path: Path):
+    proc = _run_profiler_cli(
+        tmp_path,
+        """
+        src = np.arange(1, 7, dtype=np.float32)
+        dst = np.zeros(6, dtype=np.float32)
+        ct.launch(None, (1,), copy_kernel, (src, dst))
+        ct.launch(None, (1,), copy_kernel, (src, dst))
+        """,
+    )
+
+    assert proc.returncode == 0, proc.stderr
+    # Like the Triton CLI, one profiler per kernel reports after every launch.
+    assert proc.stdout.count("Profiler Issues Summary Ends") == 2
+    assert "Total mask elements:     8\n" in proc.stdout
+    assert "Total mask elements:     16\n" in proc.stdout
+
+
+def test_other_clis_keep_nested_interpreted_launches(tmp_path: Path):
+    proc = _run_profiler_cli(
+        tmp_path,
+        """
+        from tilelens.core.frontend.cutile import interpreted_launches
+
+        src = np.arange(1, 7, dtype=np.float32)
+        dst = np.zeros(6, dtype=np.float32)
+        with interpreted_launches(), interpreted_launches():
+            ct.launch(None, (1,), copy_kernel, (src, dst))
+        print("DST", dst.tolist())
+        """,
+        command=SANITIZER_COMMAND,
+    )
+
+    assert proc.returncode == 0, proc.stderr
+    assert "DST [1.0, 2.0, 3.0, 4.0, 5.0, 6.0]" in proc.stdout
+
+
+def test_profiler_skipping_keeps_atomics(_isolate_profiler_cfg):
+    cfg.profiler_enable_block_sampling = False
+    cfg.profiler_enable_load_store_skipping = True
+
+    @ct.kernel
+    def atomic_kernel(counter, old):
+        previous = ct.atomic_add(counter, ct.full((4,), 0, ct.int32), 1)
+        ct.store(old, (0,), previous)
+
+    counter = np.array([10], np.int32)
+    old = np.full(4, -1, np.int32)
+    tilelens.trace("profiler", frontend="cutile")(atomic_kernel)[(1,)](counter, old)
+
+    # Like Triton, skipping drops plain stores but still applies atomics.
+    np.testing.assert_array_equal(counter, [14])
+    np.testing.assert_array_equal(old, -1)
diff --git a/tests/cutile/test_cutile_tracer.py b/tests/cutile/test_cutile_tracer.py
index 2298f249c..b80ce3aeb 100644
--- a/tests/cutile/test_cutile_tracer.py
+++ b/tests/cutile/test_cutile_tracer.py
@@ -7,6 +7,7 @@
 
 import numpy as np
 import pytest
+import torch
 
 ct = pytest.importorskip("cuda.tile")
 
@@ -215,6 +216,44 @@ def reduce_kernel(src, dst):
     assert [type(r) for r in launches[-1].records] == [Grid, Load, ReduceSum, Store]
 
 
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("trace", [False, True])
+def test_device_tensors_run_on_host_copies(trace):
+    from tilelens.core.simulation.cutile import CuTileInterpretedFunction
+
+    def kernel(src, dst):
+        block = ct.bid(0)
+        tile = ct.load(src, (block,), (2,))
+        ct.store(dst, (block,), tile + 1)
+
+    buffer = torch.zeros(8, dtype=torch.bfloat16, device="cuda")
+    buffer[:4] = torch.arange(4, device="cuda")
+    runner = (
+        tilelens.trace(frontend="cutile")(kernel)
+        if trace
+        else CuTileInterpretedFunction(kernel)
+    )
+    # Aliased views share one host copy, and writes return to the device.
+    runner.run(buffer[:4], buffer[4:], grid=(2,))
+    assert buffer.device.type == "cuda"
+    torch.testing.assert_close(
+        buffer.cpu(), torch.tensor([0, 1, 2, 3, 1, 2, 3, 4], dtype=torch.bfloat16)
+    )
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_device_tensors_are_restored_after_early_stop():
+    class StopAfterFirstBlock(Tracer):
+        def post_run_callback(self, fn):
+            return False
+
+    dst = torch.zeros(4, dtype=torch.float32, device="cuda")
+    tilelens.trace(StopAfterFirstBlock(), frontend="cutile")(copy_kernel).run(
+        torch.ones(8, dtype=torch.float32, device="cuda"), dst, grid=(2,)
+    )
+    torch.testing.assert_close(dst.cpu(), torch.ones(4))
+
+
 @pytest.mark.parametrize("grid", [2, (2,), (2, 2), (2, 2, 2)])
 @pytest.mark.parametrize("trace", [False, True])
 def test_grid_axes(grid, trace):
@@ -239,8 +278,8 @@ def kernel(dst):
 
 
 def test_unsupported_client_is_rejected():
-    with pytest.raises(ValueError, match="only the tracer"):
-        tilelens.trace("profiler", frontend="cutile")(copy_kernel)
+    with pytest.raises(ValueError, match="only the tracer and profiler"):
+        tilelens.trace("sanitizer", frontend="cutile")(copy_kernel)
 
 
 def test_trace_tiled_views(tmp_path):
diff --git a/tilelens/clients/profiler/profiler.py b/tilelens/clients/profiler/profiler.py
index e95723549..3b2f4648f 100644
--- a/tilelens/clients/profiler/profiler.py
+++ b/tilelens/clients/profiler/profiler.py
@@ -15,6 +15,13 @@
 from dataclasses import dataclass, replace
 
 
+def _element_size(ptr) -> int:
+    if isinstance(ptr, TensorHandle):
+        return _get_np_dtype(ptr.get_element_ty()).itemsize
+    # Non-Triton frontends (e.g. cuTile) pass arrays backed by NumPy data.
+    return ptr.data.dtype.itemsize
+
+
 @dataclass(frozen=False)
 class LoopInfo:
     length: int | None = None
@@ -93,6 +100,9 @@ def __init__(
         # Config has enable_load_store_skipping, but profiler uses disable_load_store_skipping
         self.disable_load_store_skipping = not cfg.profiler_enable_load_store_skipping
 
+        # Frontend whose op signatures the skipping overriders follow
+        self.frontend_name = "triton"
+
     def pre_run_callback(self, fn: Callable) -> bool:
         # If block sampling is enabled, check if current block is selected
         if self.block_sampling and self.sampled_blocks is not None:
@@ -159,12 +169,11 @@ def grid_callback(self, grid: tuple[int, ...]):
     def _report_load_store_bytes(
         self, type, ptr: TensorHandle, mask: TensorHandle
     ):  # internal methods assumed to be called under the lock
-        dtype_tt = ptr.get_element_ty()
-        dtype_np: np.dtype = _get_np_dtype(dtype_tt)
+        itemsize = _element_size(ptr)
         mask_true = np.count_nonzero(mask.data)
         mask_false = np.count_nonzero(np.logical_not(mask.data))
-        total_bytes_true = mask_true * dtype_np.itemsize
-        total_bytes_attempted = (mask_true + mask_false) * dtype_np.itemsize
+        total_bytes_true = mask_true * itemsize
+        total_bytes_attempted = (mask_true + mask_false) * itemsize
         if type == "load":
             self.load_bytes.total_bytes_attempted += total_bytes_attempted
             self.load_bytes.total_bytes_true += total_bytes_true
@@ -305,25 +314,43 @@ def pre_addptr_callback(ptr, offset):
             if not self.disable_buffer_load_check:
                 self._check_32bit_range(byte_offset, element_bytewidth, offset_data)
 
+        overriders: dict[type[Op], Callable] = {
+            Load: load_overrider,
+            Store: store_overrider,
+            Dot: dot_overrider,
+        }
+        if self.frontend_name == "cutile":
+            from ...core.frontend.cutile import (
+                cutile_skip_store,
+                cutile_zero_dot,
+                cutile_zero_load,
+            )
+
+            overriders = {
+                Load: self.lock_fn(cutile_zero_load),
+                Store: self.lock_fn(cutile_skip_store),
+                Dot: self.lock_fn(cutile_zero_dot),
+            }
+
         if op_type is Load:
             if self.disable_load_store_skipping:
                 return OpCallbacks(before_callback=pre_load_callback)
             else:
                 return OpCallbacks(
-                    before_callback=pre_load_callback, op_overrider=load_overrider
+                    before_callback=pre_load_callback, op_overrider=overriders[Load]
                 )
         elif op_type is Store:
             if self.disable_load_store_skipping:
                 return OpCallbacks(before_callback=pre_store_callback)
             else:
                 return OpCallbacks(
-                    before_callback=pre_store_callback, op_overrider=store_overrider
+                    before_callback=pre_store_callback, op_overrider=overriders[Store]
                 )
         elif op_type is Dot:
             if self.disable_load_store_skipping:
                 return OpCallbacks()
             else:
-                return OpCallbacks(op_overrider=dot_overrider)
+                return OpCallbacks(op_overrider=overriders[Dot])
         elif op_type is AddPtr:
             return OpCallbacks(before_callback=pre_addptr_callback)
 
diff --git a/tilelens/core/frontend/cutile.py b/tilelens/core/frontend/cutile.py
index 1226b2ff5..a8e55514e 100644
--- a/tilelens/core/frontend/cutile.py
+++ b/tilelens/core/frontend/cutile.py
@@ -2,6 +2,7 @@
 
 from collections.abc import Callable
 from contextlib import contextmanager
+import sys
 from threading import Lock
 from typing import Any
 
@@ -9,6 +10,7 @@
 
 from tilelens.core.data import Dot, Load, Op, ProgramId, ReduceSum, Store
 
+from ..patch import LoopSite, loop_file_token
 from .base import AdapterResult, Frontend, _LangPatchScope, register_frontend
 
 
@@ -19,6 +21,8 @@
         Array,
         Tile,
         _indexed_access,
+        _numpy_dtype,
+        _promotion_type,
         _tile_access,
         _tile_value,
         cutile_builder,
@@ -138,6 +142,49 @@ def _cutile_reduce_sum_adapter(
     return AdapterResult(input_tensor, axis, keepdims)
 
 
+def cutile_zero_load(
+    array: Any,
+    index: Any,
+    shape: Any = None,
+    *,
+    order: Any = "C",
+    traversal_steps: Any = None,
+    **kwargs: Any,
+) -> Any:
+    """Return zeros shaped like ct.load or ct.gather without reading memory."""
+    assert HAS_CUTILE
+    if shape is None:
+        _, valid = _indexed_access(
+            array, index, kwargs.get("mask"), kwargs.get("check_bounds", True)
+        )
+        return Tile(np.zeros(valid.shape, dtype=array.data.dtype))
+    shape, *_ = _tile_access(array, index, shape, order, traversal_steps)
+    return Tile(np.zeros(shape, dtype=array.data.dtype))
+
+
+def cutile_skip_store(*_args: Any, **_kwargs: Any) -> None:
+    """Drop ct.store and ct.scatter writes."""
+    return None
+
+
+def cutile_zero_dot(x: Any, y: Any, *args: Any, **kwargs: Any) -> Any:
+    """Return zeros shaped like matmul, mma, or mma_scaled without computing."""
+    assert HAS_CUTILE
+    # mma takes (acc,); mma_scaled takes (y, y_scale[, acc]) with acc= allowed.
+    acc = kwargs.get("acc", args[-1] if args else None)
+    if acc is not None:
+        return Tile(np.zeros(acc.shape, dtype=_tile_value(acc).dtype))
+
+    from cuda.tile._ir.ops_utils import promote_dtypes
+
+    common = promote_dtypes(_promotion_type(x), _promotion_type(y))
+    lhs, rhs = _tile_value(x), _tile_value(y)
+    # An empty contraction axis yields matmul's output shape with no arithmetic.
+    lhs = np.zeros(lhs.shape[:-1] + (0,))
+    rhs = np.zeros(rhs.shape[:-2] + (0,) + rhs.shape[-1:] if rhs.ndim > 1 else (0,))
+    return Tile(np.matmul(lhs, rhs).astype(_numpy_dtype(common)), common)
+
+
 CUTILE_ADAPTERS: dict[type[Op], Callable[..., AdapterResult]] = {}
 CUTILE_NAMESPACES: dict[Any, dict[str, type[Op]]] = {}
 if HAS_CUTILE:
@@ -179,12 +226,64 @@ def __init__(self) -> None:
             adapters=definition.adapters,
             namespaces=definition.namespaces,
         )
+        self._loop_patched = False
+        self._current_client_manager: Any = None
+
+    def run_op_overrider(
+        self,
+        op: Callable,
+        op_type: type[Op],
+        op_overrider: Callable,
+        args: tuple[Any, ...],
+        kwargs: dict[str, Any],
+    ) -> Any:
+        # Atomics access memory through load and store. Triton keeps atomics as
+        # separate ops that overriders never replace, so run them unchanged.
+        if kwargs.get("_atomic", False):
+            return op(*args, **kwargs)
+        return op_overrider(*args, **kwargs)
+
+    def patch_for_loop(self) -> None:
+        self._loop_patched = True
+
+    def unpatch_for_loop(self) -> None:
+        self._loop_patched = False
+
+    def loop_iter_wrapper(
+        self,
+        iterable_callable: Callable,
+        iter_args: Any,
+        iter_kwargs: Any,
+        lineno: int,
+        range_type: str,
+    ) -> Any:
+        """Report a rewritten for-loop to clients, like Triton's loop wrapper."""
+        client_manager = self._current_client_manager
+        if client_manager is None:
+            # A rewritten helper can outlive its traced launch.
+            return iterable_callable(*iter_args, **iter_kwargs)
+        caller_file = sys._getframe(1).f_code.co_filename
+        loop_site = LoopSite(lineno, loop_file_token(caller_file))
+        return client_manager.loop_iter_wrapper(
+            iterable_callable, iter_args, iter_kwargs, loop_site, range_type
+        )
+
+    def rewrite_loops(self, fn: Any) -> Any:
+        """Return fn with its for-loops routed through loop_iter_wrapper."""
+        from .cutile_transform import rewrite_loops
+
+        return rewrite_loops(fn, self.loop_iter_wrapper)
 
     def patch_lang(self, fn, client_manager: Any = None) -> _LangPatchScope:
-        from tilelens.core.simulation.cutile import cutile_patch_lang
+        from tilelens.core.simulation import cutile as cutile_simulation
 
         scope = _LangPatchScope()
-        cutile_patch_lang(scope)
+        cutile_simulation.cutile_patch_lang(scope)
+        # Like Triton and NKI, traced runs execute kernels and helpers with
+        # their for-loops rewritten; standalone runs execute the originals.
+        if self._loop_patched and client_manager is not None:
+            scope.set_attr(self, "_current_client_manager", client_manager)
+            scope.set_attr(cutile_simulation, "_rewrite_loops", self.rewrite_loops)
         return scope
 
 
@@ -195,29 +294,47 @@ def patch_lang(self, fn, client_manager: Any = None) -> _LangPatchScope:
 _launch_scope = None
 _launch_options = None
 _launch_users = 0
+_launch_cli = False
 
 
 @contextmanager
-def interpreted_launches(client="tracer"):
-    """Run ct.launch on CPU; overlapping scopes must use identical settings."""
-    global _launch_scope, _launch_options, _launch_users
+def interpreted_launches(client="tracer", *, cli=False):
+    """Run ct.launch on CPU; overlapping scopes must use identical settings.
+
+    CLI wrappers pass cli=True. Like the Triton CLI, that scope traces each
+    kernel once so client state accumulates across launches, and it rejects
+    user-traced kernels and nested scopes.
+    """
+    global _launch_scope, _launch_options, _launch_users, _launch_cli
 
     import cuda.tile as ct
     import tilelens
     from tilelens.core.trace import CuTileTrace
 
+    traces: dict[Any, CuTileTrace] = {}
+
     def launch(stream, grid, kernel, args):
-        runner = (
-            kernel
-            if isinstance(kernel, CuTileTrace)
-            else tilelens.trace(client, frontend="cutile")(kernel)
-        )
+        if cli:
+            # tilelens.trace() rejects already-traced kernels under the CLI.
+            if kernel not in traces:
+                traces[kernel] = tilelens.trace(client, frontend="cutile")(kernel)
+            runner = traces[kernel]
+        elif isinstance(kernel, CuTileTrace):
+            runner = kernel
+        else:
+            runner = tilelens.trace(client, frontend="cutile")(kernel)
         return runner.run(*args, grid=grid)
 
     # The global entry point can be shared by worker threads. Hold this guard
     # only while changing ownership; each runner serializes its own execution.
     with _launch_scope_lock:
         options = client
+        if _launch_users and (cli or _launch_cli):
+            raise RuntimeError(
+                "interpreted_launches() cannot be used together with CLI wrapper "
+                "(e.g., tile-profiler). Please remove interpreted_launches() from "
+                "your code when using CLI tools."
+            )
         if _launch_users and options != _launch_options:
             raise RuntimeError(
                 "Overlapping interpreted_launches contexts must use the same "
@@ -227,6 +344,7 @@ def launch(stream, grid, kernel, args):
             _launch_scope = _LangPatchScope()
             _launch_scope.set_attr(ct, "launch", launch)
             _launch_options = options
+            _launch_cli = cli
         _launch_users += 1
     try:
         yield
@@ -237,3 +355,4 @@ def launch(stream, grid, kernel, args):
                 _launch_scope.restore()
                 _launch_scope = None
                 _launch_options = None
+                _launch_cli = False
diff --git a/tilelens/core/frontend/cutile_transform.py b/tilelens/core/frontend/cutile_transform.py
new file mode 100644
index 000000000..aedefa8f3
--- /dev/null
+++ b/tilelens/core/frontend/cutile_transform.py
@@ -0,0 +1,209 @@
+"""Rewrite cuTile kernel and helper for-loops to report loop iterations."""
+
+import ast
+import inspect
+import os
+from pathlib import Path
+import sysconfig
+import textwrap
+import types
+from collections.abc import Callable
+from typing import Any, cast
+from weakref import WeakKeyDictionary
+
+LOOP_WRAPPER_ARG = "_tilelens_loop_iter_wrapper"
+_FACTORY_NAME = "_tilelens_loop_factory"
+
+
+class CuTileLoopTransformer(ast.NodeTransformer):
+    """
+    AST transformer for cuTile loop tracing, matching Triton's loop rewrite.
+
+    Transformations:
+    - `for x in f(*args, **kwargs)` -> `for x in wrapper(f, args, kwargs, lineno, range_type)`
+    - `for x in expr` -> `for x in wrapper(lambda: expr, (), {}, lineno, range_type)`
+    - adds the wrapper as a hidden keyword-only argument of every function
+    """
+
+    def visit_For(self, node: ast.For) -> ast.AST:
+        self.generic_visit(node)
+
+        # Preserve which spelling the user wrote so clients can tell loop kinds apart.
+        range_type = "unknown"
+        if isinstance(node.iter, ast.Call):
+            func = node.iter.func
+            if isinstance(func, ast.Name) and func.id == "range":
+                range_type = "python_range"
+
+        # Pass the callable and its arguments separately so range-wrapper
+        # clients can replace the iterable before it is created.
+        if isinstance(node.iter, ast.Call):
+            iter_callable = node.iter.func
+            iter_args = ast.Tuple(elts=node.iter.args, ctx=ast.Load())
+            kw_keys: list[ast.expr | None] = []
+            kw_vals: list[ast.expr] = []
+            for kw in node.iter.keywords:
+                if kw.arg is None:
+                    continue
+                kw_keys.append(ast.Constant(value=kw.arg))
+                kw_vals.append(kw.value)
+            iter_kwargs = ast.Dict(keys=kw_keys, values=kw_vals)
+        else:
+            iter_callable = ast.Lambda(
+                args=ast.arguments(
+                    posonlyargs=[],
+                    args=[],
+                    kwonlyargs=[],
+                    kw_defaults=[],
+                    defaults=[],
+                ),
+                body=node.iter,
+            )
+            iter_args = ast.Tuple(elts=[], ctx=ast.Load())
+            iter_kwargs = ast.Dict(keys=[], values=[])
+
+        loop = ast.For(
+            target=node.target,
+            iter=ast.Call(
+                func=ast.Name(id=LOOP_WRAPPER_ARG, ctx=ast.Load()),
+                args=[
+                    iter_callable,
+                    iter_args,
+                    iter_kwargs,
+                    ast.Constant(value=node.lineno),
+                    ast.Constant(value=range_type),
+                ],
+                keywords=[],
+            ),
+            body=node.body,
+            orelse=node.orelse,
+            type_comment=node.type_comment,
+        )
+        return ast.fix_missing_locations(ast.copy_location(loop, node))
+
+    def visit_FunctionDef(self, node: ast.FunctionDef) -> ast.AST:
+        node = cast(ast.FunctionDef, self.generic_visit(node))
+        # A defaulted keyword-only argument lets rewritten loops call the
+        # wrapper without adding a name to the user's module globals.
+        if LOOP_WRAPPER_ARG in {
+            arg.arg
+            for arg in (node.args.posonlyargs + node.args.args + node.args.kwonlyargs)
+        }:
+            return node
+        node.args.kwonlyargs.append(ast.arg(arg=LOOP_WRAPPER_ARG))
+        node.args.kw_defaults.append(ast.Name(id=LOOP_WRAPPER_ARG, ctx=ast.Load()))
+        return ast.fix_missing_locations(node)
+
+
+# Like Triton rewriting only @jit functions, leave library code untouched.
+_LIBRARY_DIRS = tuple(
+    os.path.join(os.path.realpath(path), "")
+    for key in ("stdlib", "platstdlib", "purelib", "platlib")
+    if (path := sysconfig.get_paths().get(key))
+) + (os.path.join(str(Path(__file__).resolve().parents[2]), ""),)
+
+
+def _is_library_code(fn: types.FunctionType) -> bool:
+    return os.path.realpath(fn.__code__.co_filename).startswith(_LIBRARY_DIRS)
+
+
+# Rewritten functions keyed by original function; None marks functions that
+# are left unchanged (no for-loops or no available source).
+_REWRITTEN: WeakKeyDictionary[Callable, Callable | None] = WeakKeyDictionary()
+
+
+def _function_node(fn: types.FunctionType) -> ast.FunctionDef | None:
+    if fn.__code__.co_name == "<lambda>":
+        return None
+    try:
+        source = textwrap.dedent(inspect.getsource(fn))
+        tree = ast.parse(source)
+    except (OSError, TypeError, SyntaxError):
+        return None
+    if not tree.body or not isinstance(tree.body[0], ast.FunctionDef):
+        return None
+    node = tree.body[0]
+    if node.name != fn.__name__:
+        return None
+    if not any(isinstance(child, ast.For) for child in ast.walk(node)):
+        return None
+    # Line 1 of the source is co_firstlineno, the first decorator if any.
+    # Keep the original file line numbers in reports and tracebacks.
+    ast.increment_lineno(node, fn.__code__.co_firstlineno - 1)
+    return node
+
+
+def _compile_rewritten(
+    fn: types.FunctionType, wrapper: Callable
+) -> types.FunctionType | None:
+    node = _function_node(fn)
+    if node is None:
+        return None
+    node = cast(ast.FunctionDef, CuTileLoopTransformer().visit(node))
+
+    # Compile inside a factory whose parameters are the original free
+    # variables, so closure references stay closure references.
+    freevars = fn.__code__.co_freevars
+    factory = ast.FunctionDef(
+        name=_FACTORY_NAME,
+        args=ast.arguments(
+            posonlyargs=[],
+            args=[ast.arg(arg=name) for name in freevars],
+            kwonlyargs=[],
+            kw_defaults=[],
+            defaults=[],
+        ),
+        body=[node, ast.Return(value=ast.Name(id=node.name, ctx=ast.Load()))],
+        decorator_list=[],
+        returns=None,
+        type_params=[],
+    )
+    module = ast.Module(body=[ast.copy_location(factory, node)], type_ignores=[])
+    ast.fix_missing_locations(module)
+    try:
+        module_code = compile(module, fn.__code__.co_filename, "exec")
+    except SyntaxError:
+        return None
+
+    factory_code = next(
+        const
+        for const in module_code.co_consts
+        if isinstance(const, types.CodeType) and const.co_name == _FACTORY_NAME
+    )
+    code = next(
+        (
+            const
+            for const in factory_code.co_consts
+            if isinstance(const, types.CodeType) and const.co_name == fn.__name__
+        ),
+        None,
+    )
+    if code is None or set(code.co_freevars) - set(freevars):
+        return None
+    # The qualified name identifies user frames in tracebacks and the code view.
+    code = code.replace(co_qualname=fn.__code__.co_qualname)
+    cells = dict(zip(freevars, fn.__closure__ or ()))
+    rewritten = types.FunctionType(
+        code,
+        fn.__globals__,
+        fn.__name__,
+        fn.__defaults__,
+        tuple(cells[name] for name in code.co_freevars) or None,
+    )
+    rewritten.__kwdefaults__ = {**(fn.__kwdefaults__ or {}), LOOP_WRAPPER_ARG: wrapper}
+    rewritten.__qualname__ = fn.__qualname__
+    rewritten.__module__ = fn.__module__
+    rewritten.__doc__ = fn.__doc__
+    rewritten.__dict__.update(fn.__dict__)
+    return rewritten
+
+
+def rewrite_loops(fn: Any, wrapper: Callable) -> Any:
+    """Return fn with for-loops routed through wrapper, or fn unchanged."""
+    if not isinstance(fn, types.FunctionType) or _is_library_code(fn):
+        return fn
+    if fn in _REWRITTEN:
+        return _REWRITTEN[fn] or fn
+    rewritten = _compile_rewritten(fn, wrapper)
+    _REWRITTEN[fn] = rewritten
+    return rewritten or fn
diff --git a/tilelens/core/simulation/cutile.py b/tilelens/core/simulation/cutile.py
index f2aba9d90..6937fe50e 100644
--- a/tilelens/core/simulation/cutile.py
+++ b/tilelens/core/simulation/cutile.py
@@ -35,6 +35,10 @@
 
 cutile_execution_lock = threading.RLock()
 
+# Rewrites a function's for-loops to report them to clients. The cuTile
+# frontend installs it through its patch scope for traced runs only.
+_rewrite_loops: Any = None
+
 
 DTYPES = {
     ct.bool_: np.dtype(np.bool_),
@@ -1087,7 +1091,8 @@ def _atomic(
             )
         keys, valid = _indexed_access(array, indices, check_bounds=check_bounds)
         updates = np.broadcast_to(_implicit_value(update, array.dtype), valid.shape)
-        previous = self.gather(array, indices, check_bounds=check_bounds)
+        # _atomic marks accesses that op overriders must not skip.
+        previous = self.gather(array, indices, check_bounds=check_bounds, _atomic=True)
         old = previous.data.copy()
         final = previous.data.copy()
         positions = tuple(key[valid] for key in keys)
@@ -1143,7 +1148,12 @@ def _atomic(
         old[valid], final[valid] = old_values, new_values
         # One vector access per operation, even when addresses repeat.
         self.scatter(
-            array, indices, Tile(final), mask=Tile(valid), check_bounds=check_bounds
+            array,
+            indices,
+            Tile(final),
+            mask=Tile(valid),
+            check_bounds=check_bounds,
+            _atomic=True,
         )
         return Tile(old)
 
@@ -1928,13 +1938,16 @@ def set_global(namespace, name, value):
             scope.callback(operator.setitem, namespace, name, previous)
         namespace[name] = value
 
+    def rewritten(function):
+        return function if _rewrite_loops is None else _rewrite_loops(function)
+
     def patched_value(value):
         operation = originals.get(id(value))
         if operation is not None:
             wrapper = _language_operation(operation)
             visited.add(id(wrapper))
             return wrapper
-        return inspect.unwrap(value) if _is_tile_function(value) else value
+        return rewritten(inspect.unwrap(value) if _is_tile_function(value) else value)
 
     def tile_minmax(original, operation):
         def execute(*args, **kwargs):
@@ -1991,13 +2004,14 @@ def patch_function(function):
             # through values the static scan cannot see, patching them on call.
             class InterpreterMode(DispatchMode):
                 def call_tile_function_from_host(self, func, args, kwargs):
-                    function = inspect.unwrap(func)
+                    function = rewritten(inspect.unwrap(func))
                     patch_function(function)
                     return function(*args, **kwargs)
 
             scope.enter_context(InterpreterMode().as_current())
-        patch_function(fn)
-        yield fn
+        execution_fn = rewritten(fn)
+        patch_function(execution_fn)
+        yield execution_fn
 
 
 def _parameter_annotations(fn):
@@ -2034,7 +2048,25 @@ def _check_integer_scalar(value, dtype):
         raise ValueError(f"Integer argument is out of range of {dtype}")
 
 
-def _wrap_argument(arg, scalar_dtype=None):
+def _to_host(arg, storages):
+    """Copy a device tensor to the CPU once per storage, like Triton's interpreter."""
+    storage = arg.untyped_storage()
+    if storage.data_ptr() not in storages:
+        storages[storage.data_ptr()] = (storage, storage.cpu())
+    host = arg.new_empty(0, device="cpu")
+    host.set_(
+        storages[storage.data_ptr()][1], arg.storage_offset(), arg.size(), arg.stride()
+    )
+    return host
+
+
+def _restore_args_dev(storages):
+    """Copy host storages back to their device tensors to propagate writes."""
+    for device_storage, host_storage in storages.values():
+        device_storage.copy_(host_storage)
+
+
+def _wrap_argument(arg, scalar_dtype, storages):
     if scalar_dtype is not None:
         _check_integer_scalar(arg, scalar_dtype)
         return Tile(arg, scalar_dtype)
@@ -2047,7 +2079,7 @@ def _wrap_argument(arg, scalar_dtype=None):
     torch = sys.modules.get("torch")
     if torch is not None and isinstance(arg, torch.Tensor):
         if arg.device.type != "cpu":
-            raise TypeError("cuTile interpretation requires CPU tensors")
+            arg = _to_host(arg, storages)
         arg = arg.detach()
         storage = {
             torch.bfloat16: torch.int16,
@@ -2095,19 +2127,25 @@ def _run(self, *args, grid=(1,), client_manager=None, **kwargs) -> None:
         bound = signature.bind(*args, **kwargs)
         bound.apply_defaults()
         annotations = _parameter_annotations(self.fn)
+        # Device tensors run on host copies of their storages, keyed by device
+        # pointer so aliased arguments keep sharing memory.
+        storages: dict[int, tuple[Any, Any]] = {}
         for name, arg in bound.arguments.items():
             is_constant, scalar_dtype = annotations.get(name, (False, None))
             if is_constant:
                 continue
             kind = signature.parameters[name].kind
             if kind is inspect.Parameter.VAR_POSITIONAL:
-                arg = tuple(_wrap_argument(item, scalar_dtype) for item in arg)
+                arg = tuple(
+                    _wrap_argument(item, scalar_dtype, storages) for item in arg
+                )
             elif kind is inspect.Parameter.VAR_KEYWORD:
                 arg = {
-                    key: _wrap_argument(item, scalar_dtype) for key, item in arg.items()
+                    key: _wrap_argument(item, scalar_dtype, storages)
+                    for key, item in arg.items()
                 }
             else:
-                arg = _wrap_argument(arg, scalar_dtype)
+                arg = _wrap_argument(arg, scalar_dtype, storages)
             bound.arguments[name] = arg
         args, kwargs = bound.args, bound.kwargs
 
@@ -2142,10 +2180,12 @@ def _run(self, *args, grid=(1,), client_manager=None, **kwargs) -> None:
                     if client_manager is not None:
                         client_manager.grid_idx_callback(grid_idx)
                         if not client_manager.pre_run_callback(self.fn):
-                            return
+                            continue
                     execution_fn(*args, **kwargs)
                     if client_manager is not None:
                         if not client_manager.post_run_callback(self.fn):
-                            return
+                            break
         finally:
             cutile_unpatch_lang(scope)
+        # Like Triton, copy arguments back to propagate side effects.
+        _restore_args_dev(storages)
diff --git a/tilelens/core/trace.py b/tilelens/core/trace.py
index ff57846ff..c531a9b12 100644
--- a/tilelens/core/trace.py
+++ b/tilelens/core/trace.py
@@ -241,8 +241,15 @@ def __init__(self, kernel, client: str | Client) -> None:
 
     def add_client(self, new_client: str | Client) -> None:
         selected = self._normalize_client(new_client)
-        if not isinstance(selected, Tracer):
-            raise ValueError("cuTile currently supports only the tracer client")
+        if not isinstance(selected, (Tracer, Profiler)):
+            raise ValueError(
+                "cuTile currently supports only the tracer and profiler clients"
+            )
+        if isinstance(selected, Profiler):
+            # Skip cuTile ops with cuTile signatures; the buffer-load check
+            # inspects AMD assembly, which interpreted cuTile never produces.
+            selected.frontend_name = self.frontend_name
+            selected.disable_buffer_load_check = True
         super().add_client(selected)
 
     def run(self, *args, grid=(1,), **kwargs):
diff --git a/tilelens/wrapper.py b/tilelens/wrapper.py
index d2759da85..b6cfe2d42 100644
--- a/tilelens/wrapper.py
+++ b/tilelens/wrapper.py
@@ -2,6 +2,8 @@
 import shutil
 import runpy
 import sys
+from contextlib import nullcontext
+from typing import Any
 import triton
 import tilelens
 from triton.experimental import gluon as _gluon
@@ -78,9 +80,12 @@ def _decorator(f):
     return _patched_autotune
 
 
-def _apply_wrapper(wrapper_func, command_name, usage_msg):
+def _apply_wrapper(wrapper_func, command_name, usage_msg, cutile_client=None):
     """
     Generic function to apply a wrapper to triton.jit and run the user script.
+
+    When cutile_client is given, ct.launch also runs kernels through the cuTile
+    interpreter with that client.
     """
     legacy_command = LEGACY_COMMANDS[command_name]
     if os.path.basename(sys.argv[0]) not in (command_name, legacy_command):
@@ -113,7 +118,18 @@ def _apply_wrapper(wrapper_func, command_name, usage_msg):
         frontend="gluon",
     )
 
-    # run user script
+    # Patch cuTile launches when the client supports the cuTile frontend.
+    cutile_launches: Any = nullcontext()
+    if cutile_client is not None:
+        from tilelens.core.frontend.cutile import HAS_CUTILE, interpreted_launches
+
+        if HAS_CUTILE:
+            cutile_launches = interpreted_launches(cutile_client, cli=True)
+    with cutile_launches:
+        _run_user_script(usage_msg)
+
+
+def _run_user_script(usage_msg):
     if len(sys.argv) < 2:
         print(usage_msg)
         sys.exit(1)
@@ -166,6 +182,7 @@ def apply_profiler():
         profiler_wrapper,
         PROFILER_COMMAND,
         f"Usage: {PROFILER_COMMAND} <script.py> [args...]",
+        cutile_client="profiler",
     )