From 7fe0d46c82d57a396d613389e8cf1bdf69651c9b Mon Sep 17 00:00:00 2001 From: Michael Lazos Date: Fri, 20 Mar 2026 17:15:58 -0700 Subject: [PATCH 0001/1172] [user-streams] fix bug where certain APIs don't take cuda events (#178027) Fixes issues where torch.Stream.record_stream doesn't work with torch.cuda.event even though it should. We just call torch.cuda.event.record(torch.Stream) and it works. lol Pull Request resolved: https://github.com/pytorch/pytorch/pull/178027 Approved by: https://github.com/xmfan --- test/dynamo/test_streams.py | 24 ++++++++++++++++++++++++ torch/_dynamo/variables/streams.py | 4 ++-- 2 files changed, 26 insertions(+), 2 deletions(-) diff --git a/test/dynamo/test_streams.py b/test/dynamo/test_streams.py index 2784c3655bd99..dcca23b5546f5 100644 --- a/test/dynamo/test_streams.py +++ b/test/dynamo/test_streams.py @@ -1702,6 +1702,30 @@ def fn(x): torch.ones(2, 2, device="cuda") ) + @requires_cuda + def test_cuda_event_record_on_stream(self): + """torch.cuda.Event should be accepted by torch.Stream.record_event (C++ type check).""" + s = torch.Stream(device="cuda") + e = torch.cuda.Event() + # This hits THPStream_record_event in Stream.cpp which does a type check + s.record_event(e) + + @requires_cuda + def test_event_record_wait_on_default_stream(self): + e = torch.cuda.Event() + + def f(x): + y = x + 1 + e.record() + e.wait() + return y + 1 + + f_compiled = torch.compile(f) + x = torch.randn(10, device="cuda") + eager_result = f(x) + compiled_result = f_compiled(x) + self.assertEqual(eager_result, compiled_result) + if __name__ == "__main__": from torch._dynamo.test_case import run_tests diff --git a/torch/_dynamo/variables/streams.py b/torch/_dynamo/variables/streams.py index 7a861d3d7caf9..e9be887c31ead 100644 --- a/torch/_dynamo/variables/streams.py +++ b/torch/_dynamo/variables/streams.py @@ -125,7 +125,7 @@ def _( def record_event(event_index: int, stream_index: int) -> None: event = _get_event_by_index(event_index) stream = _get_stream_by_index(stream_index) - stream.record_event(event) + event.record(stream) @record_event.register_fake @@ -143,7 +143,7 @@ def _( def wait_event(event_index: int, stream_index: int) -> None: event = _get_event_by_index(event_index) stream = _get_stream_by_index(stream_index) - stream.wait_event(event) + event.wait(stream) @wait_event.register_fake From 0aae4e35f59ceb2fa8d5d1991152d8762fdf9c43 Mon Sep 17 00:00:00 2001 From: Michael Lazos Date: Fri, 20 Mar 2026 17:15:58 -0700 Subject: [PATCH 0002/1172] [dynamo] Add synchronize_event custom op for Event.synchronize() (#177613) Previously, Event.synchronize() was traced as a call_method proxy node which would get silently dropped during AOTAutograd re-tracing since it's not a dispatched operator. This adds a proper torch.ops.streams.synchronize_event custom op (following the pattern of record_event/wait_event) and updates EventVariable.call_method to emit it as a call_function node. Authored with Claude. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177613 Approved by: https://github.com/williamwen42 ghstack dependencies: #178027 --- test/dynamo/test_streams.py | 38 ++++++++++++++++++++++++++++++ torch/_dynamo/variables/streams.py | 19 ++++++++++++++- torch/_inductor/lowering.py | 1 + 3 files changed, 57 insertions(+), 1 deletion(-) diff --git a/test/dynamo/test_streams.py b/test/dynamo/test_streams.py index dcca23b5546f5..96f78f6bef5ac 100644 --- a/test/dynamo/test_streams.py +++ b/test/dynamo/test_streams.py @@ -1501,6 +1501,10 @@ def test_is_marked_side_effectful(self): torch.ops.streams.record_event.default, torch.fx.node._side_effectful_functions, ) + self.assertIn( + torch.ops.streams.synchronize_event.default, + torch.fx.node._side_effectful_functions, + ) @requires_cuda def test_backward_sync_control_deps_e2e(self) -> None: @@ -1710,6 +1714,40 @@ def test_cuda_event_record_on_stream(self): # This hits THPStream_record_event in Stream.cpp which does a type check s.record_event(e) + @requires_cuda + def test_event_synchronize_tracing(self): + def fn(x): + e = torch.Event() + e.record() + x = x + 1 + e.synchronize() + return x + + inp = (torch.ones(2, 2, device="cuda"),) + ( + _, + _, + fw_graphs, + _, + ) = extract_graph(fn, *inp) + + self.assertExpectedInline( + print_graph(fw_graphs[0]), + """\ +class (torch.nn.Module): + def forward(self, arg0_1: "f32[2, 2]"): + # + record_event = torch.ops.streams.record_event.default(0, 1); record_event = None + + # + add: "f32[2, 2]" = torch.ops.aten.add.Tensor(arg0_1, 1); arg0_1 = None + + # + synchronize_event = torch.ops.streams.synchronize_event.default(0); synchronize_event = None + return (add,) +""", # noqa: B950 + ) + @requires_cuda def test_event_record_wait_on_default_stream(self): e = torch.cuda.Event() diff --git a/torch/_dynamo/variables/streams.py b/torch/_dynamo/variables/streams.py index e9be887c31ead..7c443d81664c8 100644 --- a/torch/_dynamo/variables/streams.py +++ b/torch/_dynamo/variables/streams.py @@ -157,6 +157,20 @@ def _( has_side_effect(torch.ops.streams.wait_event.default) +@custom_op("streams::synchronize_event", mutates_args=()) +def synchronize_event(event_index: int) -> None: + event = _get_event_by_index(event_index) + event.synchronize() + + +@synchronize_event.register_fake +def _(event_index: int) -> None: + pass + + +has_side_effect(torch.ops.streams.synchronize_event.default) + + @custom_op("streams::wait_stream", mutates_args=()) def wait_stream(waiting_stream_index: int, waited_on_stream_index: int) -> None: waiting = _get_stream_by_index(waiting_stream_index) @@ -515,7 +529,10 @@ def call_method( return CONSTANT_VARIABLE_NONE elif name == "synchronize": tx.output.create_proxy( - "call_method", name, *proxy_args_kwargs([self] + args, kwargs) + "call_function", + torch.ops.streams.synchronize_event, + (self.user_object_index,), + {}, ) return CONSTANT_VARIABLE_NONE elif name == "query": diff --git a/torch/_inductor/lowering.py b/torch/_inductor/lowering.py index 062cbb9408d27..a4a36534e2152 100644 --- a/torch/_inductor/lowering.py +++ b/torch/_inductor/lowering.py @@ -2593,6 +2593,7 @@ def warn_triton_random(): # TODO: mlazos reevaluate if we want to codegen something different make_fallback(torch.ops.streams.record_event.default) make_fallback(torch.ops.streams.wait_event.default) +make_fallback(torch.ops.streams.synchronize_event.default) @register_lowering(aten.rand) From 7040cd29f6714a0044743b9a73c56dd48eee3d5e Mon Sep 17 00:00:00 2001 From: Michael Lazos Date: Fri, 20 Mar 2026 17:15:59 -0700 Subject: [PATCH 0003/1172] [inductor] Add lowering and control_deps support for synchronize_event (#177614) Add make_fallback for synchronize_event so Inductor can lower it, and update wrap_all_sync_nodes_with_control_deps to handle the new op. Since synchronize_event only takes an event_index (no stream arg), the stream is inferred from the matching record_event via a new event_to_stream map. In this case, we thread through dependencies that were recorded before the event and used after the syncrhonize to prevent reordering correctness around the synchronize. Authored with Claude. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177614 Approved by: https://github.com/aorenste ghstack dependencies: #178027, #177613 --- test/dynamo/test_streams.py | 223 ++++++++++++++++++++++ torch/_functorch/_aot_autograd/streams.py | 82 ++++++-- 2 files changed, 288 insertions(+), 17 deletions(-) diff --git a/test/dynamo/test_streams.py b/test/dynamo/test_streams.py index 96f78f6bef5ac..f6daeeeef26ff 100644 --- a/test/dynamo/test_streams.py +++ b/test/dynamo/test_streams.py @@ -1748,6 +1748,229 @@ def forward(self, arg0_1: "f32[2, 2]"): """, # noqa: B950 ) + @requires_cuda + def test_event_synchronize_inductor_lowering(self): + with patch("torch._inductor.config.implicit_fallbacks", False): + + @torch.compile() + def fn(x): + e = torch.Event() + x = x + 1 + e.record() + e.synchronize() + return x + + inp = (torch.ones(2, 2, device="cuda"),) + fn(*inp) + + @requires_cuda + def test_control_deps_wrapping_synchronize_event(self) -> None: + """Test that synchronize_event threads recorded ops' values through. + + After record_event wraps ops in control_deps and produces getitem + pass-throughs, synchronize_event must also thread those through so + that subsequent consumers depend on the synchronize. + """ + + def fn(x) -> torch.Tensor: + e = torch.Event() + y = x + 1 + e.record() + e.synchronize() + # z uses y which was produced before the record — its value must + # be threaded through both record and synchronize control_deps. + z = y * 2 + return z + + inp = (torch.ones(2, 2, device="cuda"),) + ( + _, + _, + fw_graphs, + _, + ) = extract_graph(fn, *inp) + + gm = fw_graphs[0] + graph = gm.graph + + import operator + + from torch._functorch._aot_autograd.streams import ( + set_stream, + wrap_all_sync_nodes_with_control_deps, + ) + from torch._inductor.fx_passes.control_dependencies import control_deps + + # extract_graph doesn't annotate streams, so set stream metadata on + # compute nodes to match the record_event's stream index. + record_node = next( + n + for n in graph.nodes + if n.op == "call_function" + and n.target is torch.ops.streams.record_event.default + ) + stream_idx = record_node.args[1] + for n in graph.nodes: + if ( + n.op == "call_function" + and "val" in n.meta + and n.target + not in ( + torch.ops.streams.record_event.default, + torch.ops.streams.synchronize_event.default, + ) + ): + set_stream(n, stream_idx) + + wrap_all_sync_nodes_with_control_deps(gm) + + ctrl_nodes = list(graph.find_nodes(op="call_function", target=control_deps)) + # record_event + synchronize_event = 2 control_deps nodes + self.assertEqual(len(ctrl_nodes), 2) + record_ctrl = ctrl_nodes[0] + sync_ctrl = ctrl_nodes[1] + + # synchronize_event's control_deps should depend on record's ctrl + self.assertIn(record_ctrl, sync_ctrl.args[0]) + + # The record should thread through the add (y = x + 1) + record_getitems = [ + n + for n in graph.nodes + if n.op == "call_function" + and n.target == operator.getitem + and n.args[0] is record_ctrl + ] + self.assertGreaterEqual(len(record_getitems), 1) + + # Those getitems should be passed through synchronize's control_deps + # as additional args (the passthrough deps) + sync_passthrough_args = sync_ctrl.args[2:] # skip (deps_tuple, subgraph) + for getitem in record_getitems: + self.assertIn( + getitem, + sync_passthrough_args, + "record_event's getitem should be threaded through synchronize_event", + ) + + # The mul (z = y * 2) should consume a getitem from synchronize's + # control_deps, not directly from record's. + sync_getitems = [ + n + for n in graph.nodes + if n.op == "call_function" + and n.target == operator.getitem + and n.args[0] is sync_ctrl + ] + self.assertGreaterEqual(len(sync_getitems), 1) + + # Find the mul node and verify it uses a sync getitem + mul_nodes = [ + n + for n in graph.nodes + if n.op == "call_function" and n.target == torch.ops.aten.mul.Tensor + ] + self.assertEqual(len(mul_nodes), 1) + mul_args = set(mul_nodes[0].args) + self.assertTrue( + mul_args & set(sync_getitems), + "mul should depend on synchronize_event's getitem, not record_event's", + ) + + @requires_cuda + def test_external_event_synchronize_threads_inputs(self) -> None: + """When the event was recorded externally, synchronize threads graph inputs through.""" + + def fn(x): + e = torch.Event() + y = x + 1 + e.record() + e.synchronize() + z = y * 2 + return z + + inp = (torch.ones(2, 2, device="cuda"),) + ( + _, + _, + fw_graphs, + _, + ) = extract_graph(fn, *inp) + + gm = fw_graphs[0] + graph = gm.graph + + from torch._functorch._aot_autograd.streams import ( + set_stream, + wrap_all_sync_nodes_with_control_deps, + ) + + # Remove the record_event to simulate an externally-recorded event. + record_node = next( + n + for n in graph.nodes + if n.op == "call_function" + and n.target is torch.ops.streams.record_event.default + ) + stream_idx = record_node.args[1] + graph.erase_node(record_node) + + # Set stream metadata on compute nodes. + for n in graph.nodes: + if ( + n.op == "call_function" + and "val" in n.meta + and n.target is not torch.ops.streams.synchronize_event.default + ): + set_stream(n, stream_idx) + + wrap_all_sync_nodes_with_control_deps(gm) + gm.recompile() + + self.assertExpectedInline( + print_graph(gm), + """\ +class (torch.nn.Module): + def forward(self, arg0_1: "f32[2, 2]"): + # Annotation: {'stream': 1} + add: "f32[2, 2]" = torch.ops.aten.add.Tensor(arg0_1, 1) + + # No stacktrace found for following nodes + subgraph_synchronize_event = self.subgraph_synchronize_event + control_deps = torch.ops.higher_order.control_deps((arg0_1, add), subgraph_synchronize_event, add); arg0_1 = add = subgraph_synchronize_event = None + + # Annotation: {'stream': 1} + getitem: "f32[2, 2]" = control_deps[1]; control_deps = None + + # Annotation: {'stream': 1} + mul: "f32[2, 2]" = torch.ops.aten.mul.Tensor(getitem, 2); getitem = None + return (mul,) + + class subgraph_synchronize_event(torch.nn.Module): + def forward(self, dep_0: "f32[2, 2]"): + # + synchronize_event_default = torch.ops.streams.synchronize_event.default(0) + return (synchronize_event_default, dep_0) +""", # noqa: B950 + ) + + @requires_cuda + def test_event_synchronize_control_deps_e2e(self): + """E2E: compute → record → synchronize → use result through torch.compile.""" + + def f(x): + e = torch.Event() + y = x + 1 + e.record() + e.synchronize() + z = y * 2 + return z + + inp = torch.ones(2, 2, device="cuda") + eager_result = f(inp) + compiled_result = torch.compile(f)(inp) + self.assertEqual(eager_result, compiled_result) + @requires_cuda def test_event_record_wait_on_default_stream(self): e = torch.cuda.Event() diff --git a/torch/_functorch/_aot_autograd/streams.py b/torch/_functorch/_aot_autograd/streams.py index ed2b62d57b26d..d7134ed51fbf1 100644 --- a/torch/_functorch/_aot_autograd/streams.py +++ b/torch/_functorch/_aot_autograd/streams.py @@ -28,6 +28,7 @@ _SYNC_OPS = ( torch.ops.streams.record_event.default, torch.ops.streams.wait_event.default, + torch.ops.streams.synchronize_event.default, ) @@ -363,11 +364,12 @@ def _wrap_sync_node( sync_node: Node, deps_before_sync: list[Node], visited: set[Node], -) -> Node: +) -> tuple[Node, list[Node]]: """ Core logic: wrap a single sync node in control_deps. - Returns the control_deps node that replaced the sync node. + Returns (control_deps_node, passthrough_getitems) where passthrough_getitems + are the getitem nodes that thread dependencies through the control_deps node. ``visited`` is the set of nodes at or before the sync node in graph order, used to distinguish pre-sync vs post-sync users. """ @@ -440,7 +442,7 @@ def _wrap_sync_node( # Remove original sync node sync_node.replace_all_uses_with(control_deps_node) graph.erase_node(sync_node) - return control_deps_node + return control_deps_node, list(replacements.values()) def wrap_all_sync_nodes_with_control_deps(gm: torch.fx.GraphModule) -> None: @@ -457,8 +459,14 @@ def wrap_all_sync_nodes_with_control_deps(gm: torch.fx.GraphModule) -> None: raise RuntimeError("Expected a non-empty graph") stream_to_nodes: dict[int | None, list[Node]] = {} # Maps event_index -> control_deps node that wrapped its record_event, - # so the corresponding wait_event can depend on the record. + # so the corresponding wait_event/synchronize_event can depend on the record. event_to_ctrl: dict[int, Node] = {} + # Maps event_index -> getitem nodes threaded through record_event's control_deps, + # so synchronize_event can thread them through to subsequent ops. + event_to_passthrough: dict[int, list[Node]] = {} + # Maps event_index -> stream that the event was recorded on, + # so synchronize_event can infer its stream. + event_to_stream: dict[int, int | None] = {} visited: set[Node] = set() found_sync = False @@ -472,14 +480,35 @@ def wrap_all_sync_nodes_with_control_deps(gm: torch.fx.GraphModule) -> None: if node.op == "call_function": if node.target in _SYNC_OPS: event_index: int = node.args[0] # type: ignore[assignment] - sync_stream: int | None = node.args[1] # type: ignore[assignment] - deps_before_sync = stream_to_nodes.get(sync_stream, []) - # For wait_events, add a cross-event dependency on the - # matching record_event's control_deps node so the wait - # cannot be reordered before the record. + # synchronize_event blocks the CPU thread, so it acts + # as a barrier across all streams. Collect deps from every + # stream and reset them all afterward. If the event was + # recorded externally, thread the graph inputs through so + # that any post-sync uses depend on the synchronize. + if node.target is torch.ops.streams.synchronize_event.default: + sync_stream: int | None = event_to_stream.get(event_index) + all_stream_deps: list[Node] = [ + n for nodes in stream_to_nodes.values() for n in nodes + ] + if event_index not in event_to_stream: + placeholders = [n for n in graph.nodes if n.op == "placeholder"] + deps_before_sync = [*placeholders, *all_stream_deps] + else: + deps_before_sync = all_stream_deps + else: + sync_stream = node.args[1] # type: ignore[assignment] + deps_before_sync = list(stream_to_nodes.get(sync_stream, ())) + + # For wait_event and synchronize_event, add a cross-event + # dependency on the matching record_event's control_deps node + # so they cannot be reordered before the record. if ( - node.target is torch.ops.streams.wait_event.default + node.target + in ( + torch.ops.streams.wait_event.default, + torch.ops.streams.synchronize_event.default, + ) and event_index in event_to_ctrl ): deps_before_sync = [ @@ -487,22 +516,41 @@ def wrap_all_sync_nodes_with_control_deps(gm: torch.fx.GraphModule) -> None: *deps_before_sync, ] + # For synchronize_event, also include the getitem nodes + # threaded through record_event's control_deps. This ensures + # subsequent ops that depend on recorded values get rewired + # through synchronize_event. + if ( + node.target is torch.ops.streams.synchronize_event.default + and event_index in event_to_passthrough + ): + deps_before_sync = [ + *deps_before_sync, + *event_to_passthrough[event_index], + ] + if deps_before_sync: found_sync = True - ctrl_node = _wrap_sync_node(gm, node, deps_before_sync, visited) + ctrl_node, passthrough = _wrap_sync_node( + gm, node, deps_before_sync, visited + ) else: ctrl_node = None + passthrough: list[torch.fx.Node] = [] - if ( - node.target is torch.ops.streams.record_event.default - and ctrl_node is not None - ): - event_to_ctrl[event_index] = ctrl_node + if node.target is torch.ops.streams.record_event.default: + event_to_stream[event_index] = sync_stream + if ctrl_node is not None: + event_to_ctrl[event_index] = ctrl_node + event_to_passthrough[event_index] = passthrough # Reset: ops between this sync and the next will accumulate # fresh. Ordering with prior ops is already enforced because # their uses were rewired through getitems from control_deps. - stream_to_nodes[sync_stream] = [] + if node.target is torch.ops.streams.synchronize_event.default: + stream_to_nodes.clear() + else: + stream_to_nodes[sync_stream] = [] elif "val" in node.meta: stream = get_stream(node) stream_to_nodes.setdefault(stream, []).append(node) From d90ad8efea989a1ac833bb901f992ae70e5b3f65 Mon Sep 17 00:00:00 2001 From: Michael Lazos Date: Fri, 20 Mar 2026 17:16:00 -0700 Subject: [PATCH 0004/1172] [dynamo] Add end-to-end test for Event.synchronize() with D2H copies (#177615) Tests the original repro pattern: non-blocking D2H copies with event record + synchronize, verifying compiled output matches eager. Authored with Claude. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177615 Approved by: https://github.com/williamwen42 ghstack dependencies: #178027, #177613, #177614 --- test/dynamo/test_streams.py | 25 +++++++++++++++++++++++++ 1 file changed, 25 insertions(+) diff --git a/test/dynamo/test_streams.py b/test/dynamo/test_streams.py index f6daeeeef26ff..8fad34a9407dc 100644 --- a/test/dynamo/test_streams.py +++ b/test/dynamo/test_streams.py @@ -1971,6 +1971,31 @@ def f(x): compiled_result = torch.compile(f)(inp) self.assertEqual(eager_result, compiled_result) + @requires_cuda + def test_event_synchronize_e2e(self): + def f(a_list): + a_cpu_list = [] + a_to_cpu_event_list = [] + for a in a_list: + a_cpu = a.to(device="cpu", non_blocking=True) + e = torch.Event() + e.record() + a_cpu_list.append(a_cpu) + a_to_cpu_event_list.append(e) + + for e in a_to_cpu_event_list: + e.synchronize() + + return torch.cat(a_cpu_list) + + f_compiled = torch.compile(f) + inputs = [ + torch.rand(100, dtype=torch.float16, device="cuda") for _ in range(10) + ] + eager_result = f(inputs) + compiled_result = f_compiled(inputs) + self.assertEqual(eager_result, compiled_result) + @requires_cuda def test_event_record_wait_on_default_stream(self): e = torch.cuda.Event() From 28f2520bc523dd488d303ee2641d216092c52aa7 Mon Sep 17 00:00:00 2001 From: Corbin Robeck Date: Mon, 23 Mar 2026 09:15:15 +0000 Subject: [PATCH 0005/1172] [Triton] [Inductor] Add non-TMA persistent MM Triton template for max-autotune (#177781) (#177781) Summary: Add a non-TMA persistent matmul Triton template (`persistent_mm_template`) for use in max-autotune exhaustive search. This enables persistent-kernel-style matmul on platforms that lack TMA support (e.g. AMD GPUs). Initial benchmarking shows ~10% improvement over standard matmul for non-square shapes. Test Plan: Tested on a AMD 350 machine with test_max_autotune.py and confirmed the new tests pass. Reviewed By: PaulZhang12 Differential Revision: D96971956 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177781 Approved by: https://github.com/PaulZhang12 --- test/inductor/test_max_autotune.py | 49 ++++++++++++ torch/_inductor/config.py | 5 +- torch/_inductor/kernel/mm.py | 14 +++- .../templates/triton_persistent_mm.py.jinja | 75 +++++++++++++++++++ torch/_inductor/template_heuristics/triton.py | 28 +++++++ 5 files changed, 167 insertions(+), 4 deletions(-) create mode 100644 torch/_inductor/kernel/templates/triton_persistent_mm.py.jinja diff --git a/test/inductor/test_max_autotune.py b/test/inductor/test_max_autotune.py index 75e25b526f138..2db856dc67e17 100644 --- a/test/inductor/test_max_autotune.py +++ b/test/inductor/test_max_autotune.py @@ -253,6 +253,55 @@ def mm(a, b): torch.testing.assert_close(c_actual, c_expected, atol=1e-2, rtol=1e-2) + @unittest.skipIf(not torch.version.hip, "ROCM only") + @parametrize("a_transposed", (False, True)) + @parametrize("b_transposed", (False, True)) + @parametrize("dynamic", (False, True)) + def test_max_autotune_regular_mm_persistent( + self, + a_transposed: bool, + b_transposed: bool, + dynamic: bool, + ): + def mm(a, b): + a = a.repeat(8, 8) + b = b.repeat(8, 8) + + if a_transposed: + a = a.T + if b_transposed: + b = b.T + + return torch.mm(a, b) + + M, N, K = 21, 31, 11 + a = ( + torch.randn(*((K, M) if a_transposed else (M, K))) + .to(torch.float16) + .to(GPU_TYPE) + ) + b = ( + torch.randn(*((N, K) if b_transposed else (K, N))) + .to(torch.float16) + .to(GPU_TYPE) + ) + + with config.patch( + { + "max_autotune": True, + "triton.enable_persistent_tma_matmul": "1", + "triton.native_matmul": False, + "test_configs.autotune_choice_name_regex": "mm_persistent", + } + ): + c_actual, code = run_and_get_code(torch.compile(mm, dynamic=dynamic), a, b) + c_expected = mm(a, b) + + # Verify that we are using the non-TMA persistent implementation + FileCheck().check("triton_tem_fused_mm").check("NUM_SMS").run(code[0]) + + torch.testing.assert_close(c_actual, c_expected, atol=1e-2, rtol=1e-2) + @unittest.skipIf( not has_triton_tma_device(), "Need device-side TMA support in Triton" ) diff --git a/torch/_inductor/config.py b/torch/_inductor/config.py index 5c92dd60f69bc..936e2d090efcb 100644 --- a/torch/_inductor/config.py +++ b/torch/_inductor/config.py @@ -1810,8 +1810,9 @@ class triton: # Whether to upcast float16 / bfloat16 to float32 in triton codegen (Experimental) codegen_upcast_to_fp32 = True - # Whether persistent matmul kernels should be enabled this flag only has effect when on h100 - # with a version of triton new enough to support TMA + # Whether persistent matmul kernels should be enabled. On NVIDIA H100+ with TMA support, + # this enables TMA persistent kernels. On AMD GPUs without TMA, this enables + # non-TMA persistent kernels as a fallback. enable_persistent_tma_matmul = ( os.environ.get("ENABLE_PERSISTENT_TMA_MATMUL", "0") == "1" ) diff --git a/torch/_inductor/kernel/mm.py b/torch/_inductor/kernel/mm.py index 0031cacee95e7..357a25cba2ad8 100644 --- a/torch/_inductor/kernel/mm.py +++ b/torch/_inductor/kernel/mm.py @@ -100,6 +100,14 @@ source=load_kernel_template("triton_persistent_tma_mm"), ) +# Non-TMA Triton template for persistent MM +# used on AMD +persistent_mm_template = TritonTemplate( + name="mm_persistent", + grid=persistent_mm_grid, + source=load_kernel_template("triton_persistent_mm"), +) + scaled_mm_device_tma_epilogue_scaling_template = TritonTemplate( name="scaled_mm_device_tma_epilogue_scaling", @@ -359,7 +367,6 @@ def _to_dtype(x): return ops.to_dtype(x, mat1.dtype, use_compute_types=False) args = [make_pointwise(_to_dtype)(x) for x in args] - mul_pointwise = make_pointwise(ops.dot)(*args) dot_reduction = make_reduction("dot")(mul_pointwise, 1) @@ -422,7 +429,10 @@ def _to_dtype(x): if use_triton_blackwell_tma_template(mat1, mat2, output_layout=layout): templates_to_use.append(blackwell_ws_persistent_device_tma_mm_template) elif use_triton_tma_template(mat1, mat2, output_layout=layout): - templates_to_use.append(persistent_tma_mm_template) + if torch.version.hip is None: + templates_to_use.append(persistent_tma_mm_template) + else: + templates_to_use.append(persistent_mm_template) templates_to_use.append(mm_contiguous_subgraph_template) diff --git a/torch/_inductor/kernel/templates/triton_persistent_mm.py.jinja b/torch/_inductor/kernel/templates/triton_persistent_mm.py.jinja new file mode 100644 index 0000000000000..7450533be5adf --- /dev/null +++ b/torch/_inductor/kernel/templates/triton_persistent_mm.py.jinja @@ -0,0 +1,75 @@ +{{def_kernel("A", "B")}} + M = {{size("A", 0)}} + N = {{size("B", 1)}} + K = {{size("A", 1)}} + if M * N == 0: + # early exit due to zero-size input(s) + return + stride_am = {{stride("A", 0)}} + stride_ak = {{stride("A", 1)}} + stride_bk = {{stride("B", 0)}} + stride_bn = {{stride("B", 1)}} + + # persistent kernel: each CTA processes multiple tiles + start_pid = tl.program_id(0).to(INDEX_DTYPE) + grid_m = tl.cdiv(M, BLOCK_M) + grid_n = tl.cdiv(N, BLOCK_N) + num_tiles = grid_m * grid_n + width = GROUP_M * grid_n + + for tile_id in tl.range(start_pid, num_tiles, NUM_SMS): + + # re-order program ID for better L2 performance + group_id = tile_id // width + group_size = min(grid_m - group_id * GROUP_M, GROUP_M) + pid_m = group_id * GROUP_M + (tile_id % group_size) + pid_n = (tile_id % width) // (group_size) + tl.assume(pid_m >= 0) + tl.assume(pid_n >= 0) + + rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M) + rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N) + if ((stride_am == 1 and stride_ak == M) or (stride_am == K and stride_ak == 1)) and (M >= BLOCK_M and K > 1): + offs_a_m = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M) + else: + offs_a_m = rm % M + if ((stride_bk == 1 and stride_bn == K) or (stride_bk == N and stride_bn == 1)) and (N >= BLOCK_N and K > 1): + offs_b_n = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N) + else: + offs_b_n = rn % N + offs_k = tl.arange(0, BLOCK_K) + acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=ACC_TYPE) + + for k_idx in range(0, tl.cdiv(K, BLOCK_K)): + {% if not EVEN_K %} + a_mask = offs_k[None, :] < (K - k_idx * BLOCK_K) + b_mask = offs_k[:, None] < (K - k_idx * BLOCK_K) + {% endif %} + a_k_idx_vals = offs_k[None, :] + (k_idx * BLOCK_K) + b_k_idx_vals = offs_k[:, None] + (k_idx * BLOCK_K) + + idx_m = offs_a_m[:, None] + idx_n = a_k_idx_vals + {{load_input("A", "a", ("idx_m", "idx_n"), mask=None if EVEN_K else "a_mask", + indent_width=12, index_shape=("BLOCK_M", "BLOCK_K"))}} + + idx_m = b_k_idx_vals + idx_n = offs_b_n[None, :] + {{load_input("B", "b", ("idx_m", "idx_n"), mask=None if EVEN_K else "b_mask", + indent_width=12, index_shape=("BLOCK_K", "BLOCK_N"))}} + + {% if USE_FAST_ACCUM %} + acc = tl.dot(a, b, acc, allow_tf32=ALLOW_TF32, out_dtype=ACC_TYPE) + {% else %} + acc += tl.dot(a, b, allow_tf32=ALLOW_TF32, out_dtype=ACC_TYPE) + {% endif %} + + # rematerialize rm and rn to save registers + rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M) + rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N) + idx_m = rm[:, None] + idx_n = rn[None, :] + mask = (idx_m < M) & (idx_n < N) + + # inductor generates a suffix + {{store_output(("idx_m", "idx_n"), "acc", "mask", indent_width=8, val_shape=("BLOCK_M", "BLOCK_N"))}} diff --git a/torch/_inductor/template_heuristics/triton.py b/torch/_inductor/template_heuristics/triton.py index 13c65ab25c520..793a6ae49a3aa 100644 --- a/torch/_inductor/template_heuristics/triton.py +++ b/torch/_inductor/template_heuristics/triton.py @@ -23,6 +23,7 @@ get_scaling_options, get_tile_size, mm_template, + persistent_mm_template, persistent_tma_mm_template, scaled_mm_device_tma_epilogue_scaling_template, scaled_mm_device_tma_main_loop_scaling_template, @@ -2531,6 +2532,33 @@ def __init__(self) -> None: self.mm_configs = self.persistent_mm_configs +@register_template_heuristic( + persistent_mm_template.uid, + "cuda", + register=torch.version.hip is not None, +) +class PersistentMMTemplateConfigHeuristic( + MMTemplateConfigMixin, + ROCmConfigHeuristic, # type: ignore[misc] +): + """Persistent MM template heuristic (no TMA, standard pointer loads)""" + + def __init__(self) -> None: + super().__init__() + self.mm_configs = self.persistent_mm_configs + + def _get_template_configs_impl( + self, + kernel_inputs: KernelInputs, + op_name: str, + **kwargs, + ) -> Generator[dict[str, Any], None, None]: + for template_kwargs in super()._get_template_configs_impl( + kernel_inputs, op_name, **kwargs + ): + yield {**template_kwargs, "NUM_SMS": get_num_sms()} + + @register_template_heuristic( blackwell_ws_persistent_device_tma_mm_template.uid, "cuda", From 1a4108cabed3ce40c7d5c7210812bf099a6790f2 Mon Sep 17 00:00:00 2001 From: Nicolas Hug Date: Mon, 23 Mar 2026 10:07:16 +0000 Subject: [PATCH 0006/1172] F.interpolate cleanups - again (#177877) This PR combines: - https://github.com/pytorch/pytorch/pull/176713 (reverted) - https://github.com/pytorch/pytorch/pull/176714 (reverted) - https://github.com/pytorch/pytorch/pull/177835 forward fix to both PRs above, but it also ended up being reverted (because it was merged on GH at a time where both PRs were already removed on fbcode). https://github.com/pytorch/pytorch/pull/176713 and https://github.com/pytorch/pytorch/pull/176714 had to be reverted for internal builds reasons. It was validated internally that https://github.com/pytorch/pytorch/pull/177835 **is the correct forward fix**. So this PR should be safe to merge. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177877 Approved by: https://github.com/Skylion007 --- aten/src/ATen/native/cpu/UpSampleKernel.cpp | 279 ++++++++++-------- .../native/cpu/UpSampleKernelAVXAntialias.h | 6 +- 2 files changed, 166 insertions(+), 119 deletions(-) diff --git a/aten/src/ATen/native/cpu/UpSampleKernel.cpp b/aten/src/ATen/native/cpu/UpSampleKernel.cpp index 39fe91ccc06e4..ab98daba2fc02 100644 --- a/aten/src/ATen/native/cpu/UpSampleKernel.cpp +++ b/aten/src/ATen/native/cpu/UpSampleKernel.cpp @@ -24,21 +24,54 @@ namespace { using scale_t = std::vector>; +// Naming conventions used in this file: +// +// - "non_separable": All spatial dimensions are interpolated in a single +// TensorIterator pass, using the recursive InterpolateNonSeparable struct. +// Non-separable is always multi-dimensional (Nd). +// +// - "separable": The multi-dimensional interpolation is decomposed into a +// sequence of 1d passes (one per spatial dimension). The entry point is Nd +// (it loops over dims), and inner functions are 1d. +// +// - "1d": Refers to processing a single spatial dimension within the +// separable approach. Not to be confused with 1d interpolation (e.g. +// linear); it means one dimension of a multi-dimensional separable +// decomposition. +// +// - "Nd": Templated on out_ndims (1, 2, or 3) for the separable approach. + + +// ---- Non-separable interpolation ---- +// +// Used by: nearest, linear, bilinear (float), cubic (float), trilinear. +// Processes all spatial dims in a single TensorIterator pass via the recursive +// InterpolateNonSeparable struct. +// +// Call chain: +// upsample_non_separable_Nd_kernel_impl +// -> upsample_non_separable +// -> basic_loop_non_separable +// -> interpolate_non_separable +// -> InterpolateNonSeparable (recursive struct) +// +// Helper structs and methods for upsample_non_separable +// // Interpolation structure to compute output value in n-dimensional case. // - recursively compute interpolated output for each dimension // - we rely a lot on compiler's code optimization such that implemented operations // can be automatically factorized and vectorized using SSE and AVX2 template -struct Interpolate { +struct InterpolateNonSeparable { static inline opmath_t eval(char* src, char** data, const int64_t* strides, int64_t i) { index_t ids = *(index_t*)&data[0][i * strides[0]]; opmath_t wts = *(scalar_t*)&data[1][i * strides[1]]; - opmath_t t = Interpolate::eval(src + ids, &data[2 * interp_size], &strides[2 * interp_size], i); + opmath_t t = InterpolateNonSeparable::eval(src + ids, &data[2 * interp_size], &strides[2 * interp_size], i); opmath_t output = t * wts; for (const auto j : c10::irange(1, interp_size)) { ids = *(index_t*)&data[2 * j + 0][i * strides[2 * j + 0]]; wts = *(scalar_t*)&data[2 * j + 1][i * strides[2 * j + 1]]; - t = Interpolate::eval(src + ids, &data[2 * interp_size], &strides[2 * interp_size], i); + t = InterpolateNonSeparable::eval(src + ids, &data[2 * interp_size], &strides[2 * interp_size], i); output += t * wts; } return output; @@ -46,7 +79,7 @@ struct Interpolate { }; template -struct Interpolate<1, scalar_t, opmath_t, index_t, interp_size> { +struct InterpolateNonSeparable<1, scalar_t, opmath_t, index_t, interp_size> { static inline opmath_t eval(char* src, char** data, const int64_t* strides, int64_t i) { index_t ids = *(index_t*)&data[0][i * strides[0]]; opmath_t wts = *(scalar_t*)&data[1][i * strides[1]]; @@ -63,15 +96,15 @@ struct Interpolate<1, scalar_t, opmath_t, index_t, interp_size> { }; template -struct Interpolate { +struct InterpolateNonSeparable { static inline opmath_t eval(char* src, char** data, const int64_t* strides, int64_t i) { index_t ids = *(index_t*)&data[0][i * strides[0]]; - return Interpolate::eval(src + ids, &data[2], &strides[2], i); + return InterpolateNonSeparable::eval(src + ids, &data[2], &strides[2], i); } }; template -struct Interpolate<1, scalar_t, opmath_t, index_t, 1> { +struct InterpolateNonSeparable<1, scalar_t, opmath_t, index_t, 1> { static inline opmath_t eval(char* src, char** data, const int64_t* strides, int64_t i) { index_t ids = *(index_t*)&data[0][i * strides[0]]; return *(scalar_t *)&src[ids]; @@ -81,25 +114,25 @@ struct Interpolate<1, scalar_t, opmath_t, index_t, 1> { // There is an unexpected 2x slowdown for upsample_trilinear3d channels_first // for both 1 and 6 threads. We have to specialize this case as below: // Once the issue is fixed we can keep generic implementation and remove: -// struct Interpolate and -// struct Interpolate<1, scalar_t, index_t, 2> +// struct InterpolateNonSeparable and +// struct InterpolateNonSeparable<1, scalar_t, index_t, 2> template -struct Interpolate { +struct InterpolateNonSeparable { static inline opmath_t eval(char* src, char** data, const int64_t* strides, int64_t i) { index_t i0 = *(index_t*)&data[0][i * strides[0]]; index_t i1 = *(index_t*)&data[2][i * strides[2]]; opmath_t w0 = *(scalar_t *)&data[1][i * strides[1]]; opmath_t w1 = *(scalar_t *)&data[3][i * strides[3]]; - opmath_t t0 = Interpolate::eval(src + i0, &data[4], &strides[4], i); - opmath_t t1 = Interpolate::eval(src + i1, &data[4], &strides[4], i); + opmath_t t0 = InterpolateNonSeparable::eval(src + i0, &data[4], &strides[4], i); + opmath_t t1 = InterpolateNonSeparable::eval(src + i1, &data[4], &strides[4], i); return t0 * w0 + t1 * w1; } }; template -struct Interpolate<1, scalar_t, opmath_t, index_t, 2> { +struct InterpolateNonSeparable<1, scalar_t, opmath_t, index_t, 2> { static inline opmath_t eval(char* src, char** data, const int64_t* strides, int64_t i) { index_t i0 = *(index_t*)&data[0][i * strides[0]]; index_t i1 = *(index_t*)&data[2][i * strides[2]]; @@ -112,13 +145,28 @@ struct Interpolate<1, scalar_t, opmath_t, index_t, 2> { }; template -inline scalar_t interpolate(char* src, char** data, const int64_t* strides, int64_t i) { +inline scalar_t interpolate_non_separable(char* src, char** data, const int64_t* strides, int64_t i) { using opmath_t = at::opmath_type; - return Interpolate::eval(src, data, strides, i); + return InterpolateNonSeparable::eval(src, data, strides, i); } +// ---- Separable interpolation ---- +// +// Used by: bilinear/bicubic with antialias=True, and bilinear/bicubic uint8 +// (as fallback when AVX isn't supported). +// Processes one spatial dimension at a time. The outer loop over dimensions +// is in upsample_separable_Nd_kernel_impl. +// +// Call chain: +// upsample_separable_Nd_kernel_impl (loops over dims) +// -> upsample_separable_1d +// -> basic_loop_separable_1d_horizontal (for last spatial dim, i.e. W) +// -> interpolate_separable_1d +// -> basic_loop_separable_1d_vertical (for other spatial dims, e.g. H, D) +// -> interpolate_separable_1d_zero_strides + template -inline scalar_t interpolate_aa_single_dim_zero_strides( +inline scalar_t interpolate_separable_1d_zero_strides( char* src, char** data, const index_t ids_stride) { @@ -142,7 +190,7 @@ inline scalar_t interpolate_aa_single_dim_zero_strides( } template -inline scalar_t interpolate_aa_single_dim( +inline scalar_t interpolate_separable_1d( char* src, char** data, const int64_t* strides, @@ -215,7 +263,7 @@ inline bool is_contiguous_stride(const int64_t* strides) { // strides=(0, 0, 0, 0, 0, 0, 0, 0, ..., 0, 0, 0, 0) // // Using these methods we can hint the compiler to factorize constant indices and weights -// in cpu_upsample_linear method +// in upsample_non_separable template struct CheckAlmostAllZeroStrides { static inline bool eval(const int64_t* strides) { @@ -247,17 +295,17 @@ inline bool check_almost_all_zero_stride(const int64_t* strides) { // Helper method to compute interpolation for nearest, linear, cubic modes template -inline void basic_loop(char** data, const int64_t* strides, int64_t n) { +inline void basic_loop_non_separable(char** data, const int64_t* strides, int64_t n) { char* dst = data[0]; char* src = data[1]; for (const auto i : c10::irange(n)) { - *(scalar_t*)&dst[i * strides[0]] = interpolate( + *(scalar_t*)&dst[i * strides[0]] = interpolate_non_separable( src + i * strides[1], &data[2], &strides[2], i); } } template -inline void basic_loop_aa_vertical( +inline void basic_loop_separable_1d_vertical( char** data, const int64_t* strides, int64_t n, @@ -269,13 +317,13 @@ inline void basic_loop_aa_vertical( for (const auto i : c10::irange(n)) { *(scalar_t*)&dst[i * strides[0]] = - interpolate_aa_single_dim_zero_strides( + interpolate_separable_1d_zero_strides( src + i * strides[1], &data[2], ids_stride); } } template <> -inline void basic_loop_aa_vertical( +inline void basic_loop_separable_1d_vertical( char** data, const int64_t* strides, int64_t n, @@ -313,7 +361,7 @@ inline void basic_loop_aa_vertical( } template -inline void basic_loop_aa_horizontal( +inline void basic_loop_separable_1d_horizontal( char** data, const int64_t* strides, int64_t n, @@ -326,20 +374,20 @@ inline void basic_loop_aa_horizontal( if (strides[1] == 0) { for (const auto i : c10::irange(n)) { *(scalar_t*)&dst[i * strides[0]] = - interpolate_aa_single_dim( + interpolate_separable_1d( src, &data[2], &strides[2], i, ids_stride); } } else { for (const auto i : c10::irange(n)) { *(scalar_t*)&dst[i * strides[0]] = - interpolate_aa_single_dim( + interpolate_separable_1d( src + i * strides[1], &data[2], &strides[2], i, ids_stride); } } } template <> -inline void basic_loop_aa_horizontal( +inline void basic_loop_separable_1d_horizontal( char** data, const int64_t* strides, int64_t n, @@ -390,10 +438,10 @@ inline void basic_loop_aa_horizontal( // output_DN[a] = interpolate(input_DN[a], w_DN[a], input_DN[a+1], w_DN[a+1], ...) // and i - dimension index and a - linear index for spatial coordinates // -// The recursive call is implemented with InterpLinear struct using template for +// The recursive call is implemented with the InterpolateNonSeparable struct using template for // the loop unrolling on compile time. template -void cpu_upsample_generic(at::TensorIterator& iter) +void upsample_non_separable(at::TensorIterator& iter) { auto loop = [&](char** data, const int64_t* strides, int64_t n) { // special-cases to let the compiler apply compile-time input-specific optimizations @@ -401,21 +449,21 @@ void cpu_upsample_generic(at::TensorIterator& iter) // NOLINTNEXTLINE(bugprone-branch-clone) check_almost_all_zero_stride(&strides[2]))) { // contiguous channels-first case - basic_loop(data, strides, n); + basic_loop_non_separable(data, strides, n); } else if ((strides[0] == sizeof(scalar_t) && (strides[1] == sizeof(scalar_t)) && check_almost_all_zero_stride(&strides[2]))) { // contiguous channels-last case - basic_loop(data, strides, n); + basic_loop_non_separable(data, strides, n); } else { // fallback - basic_loop(data, strides, n); + basic_loop_non_separable(data, strides, n); } }; iter.for_each(loop); } template -void cpu_upsample_nearest_channels_last( +void upsample_nearest_channels_last( const Tensor& output_, const Tensor& input_, const scale_type& scales) { @@ -520,7 +568,7 @@ inline VecType interpolate(const scalar_t* t, accscalar_t w, Args... a } template -void cpu_upsample_linear_channels_last( +void upsample_linear_channels_last( const Tensor& output_, const Tensor& input_, bool align_corners, @@ -680,7 +728,7 @@ void cpu_upsample_linear_channels_last( } } -// Helper structs to use with upsample_generic_Nd_kernel_impl +// Helper structs to use with upsample_non_separable_Nd_kernel_impl struct HelperInterpBase { static inline void init_indices_weights( @@ -746,7 +794,7 @@ struct HelperInterpBase { // for interpolation with antialiasing=false mode. It returns the maximal weights value. // This function is templated with scalar_t for type of scale and weights but is only used for // bilinear/bicubic modes on uint8 input and antialiasing=false (in this case scalar_t is double). - // For float input types we are using upsample_generic_Nd_kernel_impl and compute_indices_weights methods + // For float input types we are using upsample_non_separable_Nd_kernel_impl and compute_indices_weights methods template static inline scalar_t _compute_indices_min_size_weights( const int64_t i, const int64_t input_size, const scalar_t scale, @@ -906,7 +954,7 @@ struct HelperInterpBase { weights as double, but then convert them to int16 via some conversion logic detailed below. This allows us to compute all interpolation operation (sum of multiplications) as ints instead of floats. The result is converted back into - uint8 in basic_loop_aa_horizontal (and vertical) + uint8 in basic_loop_separable_1d_horizontal (and vertical) In essence the idea is to avoid a multiplication between a float (the weight) and an int (the pixel value) and instead run a multiplication between @@ -1377,7 +1425,7 @@ struct HelperInterpCubic : public HelperInterpBase { // - scale_type is template type for scales, typically std::optional // - template class F is one of the above structs to compute indices and weights template -void upsample_generic_Nd_kernel_impl( +void upsample_non_separable_Nd_kernel_impl( const Tensor& output, const Tensor& input, bool align_corners, @@ -1437,43 +1485,27 @@ void upsample_generic_Nd_kernel_impl( if (interp_size > 1) { // Nearest also supports uint8 tensor, so need to handle it separately + // Dispatch name should be "upsample_non_separable" but we keep the old + // name for internal BC. AT_DISPATCH_FLOATING_TYPES_AND2( kBFloat16, kHalf, iter.dtype(), "upsample_generic_Nd", [&] { // MSVC can not catch constexpr int interp_size here constexpr int mode = F::interp_size; - cpu_upsample_generic(iter); + upsample_non_separable(iter); }); } else { + // Dispatch name should be "upsample_non_separable" but we keep the old + // name for internal BC. AT_DISPATCH_FLOATING_TYPES_AND3(kByte, kBFloat16, kHalf, iter.dtype(), "upsample_generic_Nd", [&] { constexpr int mode = F::interp_size; - cpu_upsample_generic(iter); + upsample_non_separable(iter); }); } } -template -void cpu_upsample_generic_aa(at::TensorIterator& iter, unsigned int weights_precision) { - - auto loop = [&](char** data, const int64_t* strides, int64_t n) { - if constexpr (is_horizontal) { - - // Strides are : X 0 | 8 8 8 0 8 (Channels first) - // Strides are : X X | 0 0 0 0 0 (Channels last) - basic_loop_aa_horizontal(data, strides, n, weights_precision); - } else { - // Strides are : X Y | 0 0 0 0 0 (Channels first) - // Strides are : X X | 0 0 0 0 0 (Channels last) - // upsampling data between contiguous dimensions (aka vertical resampling) - basic_loop_aa_vertical(data, strides, n, weights_precision); - } - }; - - iter.for_each(loop); -} - template -void _separable_upsample_generic_Nd_kernel_impl_single_dim( +void upsample_separable_1d( const Tensor& output, const Tensor& input, int interp_dim, @@ -1533,9 +1565,22 @@ void _separable_upsample_generic_Nd_kernel_impl_single_dim( auto iter = config.build(); + // Dispatch name should be "upsample_separable_1d" but we keep the old + // name for internal BC. AT_DISPATCH_FLOATING_TYPES_AND( at::ScalarType::Byte, iter.dtype(), "upsample_generic_Nd_aa", [&] { - cpu_upsample_generic_aa(iter, weights_precision); + auto loop = [&](char** data, const int64_t* strides, int64_t n) { + if constexpr (is_horizontal) { + // Strides are : X 0 | 8 8 8 0 8 (Channels first) + // Strides are : X X | 0 0 0 0 0 (Channels last) + basic_loop_separable_1d_horizontal(data, strides, n, weights_precision); + } else { + // Strides are : X Y | 0 0 0 0 0 (Channels first) + // Strides are : X X | 0 0 0 0 0 (Channels last) + basic_loop_separable_1d_vertical(data, strides, n, weights_precision); + } + }; + iter.for_each(loop); }); } @@ -1544,7 +1589,7 @@ void _separable_upsample_generic_Nd_kernel_impl_single_dim( // (dtype == uint8 and mode in ("bilinear", "bicubic")): this is used as // fallback in these settings when AVX isn't supported. template -void separable_upsample_generic_Nd_kernel_impl( +void upsample_separable_Nd_kernel_impl( const Tensor& output, const Tensor& input, bool align_corners, @@ -1563,29 +1608,29 @@ void separable_upsample_generic_Nd_kernel_impl( at::Tensor temp_output, temp_input = input; int interp_dim = 0; - // Precompute the number of single dim resize method invocations + // Precompute the number of 1d resize ops // to avoid copying temporary buffer to output - int num_single_dim_ops = 0; + int num_1d_ops = 0; for (const auto i : c10::irange(out_ndims)) { interp_dim = 2 + out_ndims - 1 - i; if (output_shape[interp_dim] != input_shape[interp_dim]) { - num_single_dim_ops += 1; + num_1d_ops += 1; } } - // upsampling data within the contiguous dimension (aka horizontal resampling) + // Horizontal resampling (last spatial dim, i.e. W) interp_dim = 2 + out_ndims - 1; if (output_shape[interp_dim] != input_shape[interp_dim]) { - num_single_dim_ops -= 1; - if (num_single_dim_ops > 0) { + num_1d_ops -= 1; + if (num_1d_ops > 0) { temp_oshape[interp_dim] = output_shape[interp_dim]; temp_output = at::empty(temp_oshape, input.options()); } else { temp_output = output; } - _separable_upsample_generic_Nd_kernel_impl_single_dim< + upsample_separable_1d< out_ndims, scale_t, F, @@ -1594,20 +1639,20 @@ void separable_upsample_generic_Nd_kernel_impl( temp_input = temp_output; } - // upsampling data between contiguous dimensions (aka vertical resampling) + // Vertical resampling (remaining spatial dims, e.g. H, D) for (const auto i : c10::irange(1, out_ndims)) { interp_dim = 2 + out_ndims - 1 - i; if (output_shape[interp_dim] != input_shape[interp_dim]) { - num_single_dim_ops -= 1; - if (num_single_dim_ops > 0) { + num_1d_ops -= 1; + if (num_1d_ops > 0) { temp_oshape[interp_dim] = output_shape[interp_dim]; temp_output = at::empty(temp_oshape, input.options()); } else { temp_output = output; } - _separable_upsample_generic_Nd_kernel_impl_single_dim< + upsample_separable_1d< out_ndims, scale_t, F, @@ -1622,7 +1667,7 @@ void upsample_nearest1d_kernel_impl( const Tensor& output, const Tensor& input, std::optional scales_w) { - upsample_generic_Nd_kernel_impl<1, scale_t, HelperInterpNearest>( + upsample_non_separable_Nd_kernel_impl<1, scale_t, HelperInterpNearest>( output, input, false, {scales_w}); } @@ -1630,26 +1675,28 @@ void _upsample_nearest_exact1d_kernel_impl( const Tensor& output, const Tensor& input, std::optional scales_w) { - upsample_generic_Nd_kernel_impl<1, scale_t, HelperInterpNearestExact>( + upsample_non_separable_Nd_kernel_impl<1, scale_t, HelperInterpNearestExact>( output, input, false, {scales_w}); } -int _use_vectorized_kernel_cond_2d( +int _use_channels_last_kernel_2d( const Tensor& output, const Tensor& input) { - // This condition is used to know whether we should dispatch to a vectorized - // kernel, or to the more general upsample_generic_Nd_kernel_impl(). For now, - // the vectorized kernels are only optimized for channels_last and when C >= 4 - // (shape = NCHW). For a very wide range of use-cases (typically image or mask - // resizing where we have C < 4), using upsample_generic_Nd_kernel_impl() is - // actually faster. On top of that, benchmarks showed that this also depends on - // the *output* size (output_H + output_W), for both upsampling and - // downsampling. The current 128 threshold was determined through benchmarks. + // This condition is used to know whether we should dispatch to a + // channels-last-optimized kernel, or to the more general + // upsample_non_separable_Nd_kernel_impl(). For now, the channels-last kernels + // are only optimized for channels_last and when C >= 4 (shape = NCHW). + // For a very wide range of use-cases (typically image or mask resizing + // where we have C < 4), using upsample_non_separable_Nd_kernel_impl() is + // actually faster. On top of that, benchmarks showed that this also + // depends on the *output* size (output_H + output_W), for both + // upsampling and downsampling. The current 128 threshold was determined + // through benchmarks. return ((input.is_contiguous(at::MemoryFormat::ChannelsLast)) && (input.size(1) > 3)) || ((output.size(-2) + output.size(-1)) <= 128); } -int _use_vectorized_kernel_cond_3d( - // Similar to _use_vectorized_kernel_cond_2d() but for 3d resampling (e.g. videos) +int _use_channels_last_kernel_3d( + // Similar to _use_channels_last_kernel_2d() but for 3d resampling (e.g. videos) // Note that unlike the 2d case, this is not subject to small output size // overhead - hence the absence of the 128 threshold in the condition. const Tensor& output, @@ -1663,13 +1710,13 @@ void upsample_nearest2d_kernel_impl( const Tensor& input, std::optional scales_h, std::optional scales_w) { - if (_use_vectorized_kernel_cond_2d(output, input)) { + if (_use_channels_last_kernel_2d(output, input)) { AT_DISPATCH_FLOATING_TYPES_AND3(kByte, kBFloat16, kHalf, input.scalar_type(), "upsample_nearest2d_channels_last", [&] { - cpu_upsample_nearest_channels_last(output, input, {scales_h, scales_w}); + upsample_nearest_channels_last(output, input, {scales_h, scales_w}); }); } else { - upsample_generic_Nd_kernel_impl<2, scale_t, HelperInterpNearest>( + upsample_non_separable_Nd_kernel_impl<2, scale_t, HelperInterpNearest>( output, input, false, {scales_h, scales_w}); } } @@ -1679,12 +1726,12 @@ void _upsample_nearest_exact2d_kernel_impl( const Tensor& input, std::optional scales_h, std::optional scales_w) { - if (_use_vectorized_kernel_cond_2d(output, input)) { + if (_use_channels_last_kernel_2d(output, input)) { AT_DISPATCH_FLOATING_TYPES_AND3(kByte, kBFloat16, kHalf, input.scalar_type(), "upsample_nearest2d_channels_last", [&] { - cpu_upsample_nearest_channels_last(output, input, {scales_h, scales_w}); + upsample_nearest_channels_last(output, input, {scales_h, scales_w}); }); } else { - upsample_generic_Nd_kernel_impl<2, scale_t, HelperInterpNearestExact>( + upsample_non_separable_Nd_kernel_impl<2, scale_t, HelperInterpNearestExact>( output, input, false, {scales_h, scales_w}); } } @@ -1695,13 +1742,13 @@ void upsample_nearest3d_kernel_impl( std::optional scales_d, std::optional scales_h, std::optional scales_w) { - if (_use_vectorized_kernel_cond_3d(output, input)) { + if (_use_channels_last_kernel_3d(output, input)) { AT_DISPATCH_FLOATING_TYPES_AND3(kByte, kBFloat16, kHalf, input.scalar_type(), "upsample_nearest3d_channels_last", [&] { - cpu_upsample_nearest_channels_last(output, input, {scales_d, scales_h, scales_w}); + upsample_nearest_channels_last(output, input, {scales_d, scales_h, scales_w}); }); } else { - upsample_generic_Nd_kernel_impl<3, scale_t, HelperInterpNearest>( + upsample_non_separable_Nd_kernel_impl<3, scale_t, HelperInterpNearest>( output, input, false, {scales_d, scales_h, scales_w}); } } @@ -1712,12 +1759,12 @@ void _upsample_nearest_exact3d_kernel_impl( std::optional scales_d, std::optional scales_h, std::optional scales_w) { - if (_use_vectorized_kernel_cond_3d(output, input)) { + if (_use_channels_last_kernel_3d(output, input)) { AT_DISPATCH_FLOATING_TYPES_AND3(kByte, kBFloat16, kHalf, input.scalar_type(), "upsample_nearest3d_channels_last", [&] { - cpu_upsample_nearest_channels_last(output, input, {scales_d, scales_h, scales_w}); + upsample_nearest_channels_last(output, input, {scales_d, scales_h, scales_w}); }); } else { - upsample_generic_Nd_kernel_impl<3, scale_t, HelperInterpNearestExact>( + upsample_non_separable_Nd_kernel_impl<3, scale_t, HelperInterpNearestExact>( output, input, false, {scales_d, scales_h, scales_w}); } } @@ -1727,7 +1774,7 @@ void upsample_linear1d_kernel_impl( const Tensor& input, bool align_corners, std::optional scales_w) { - upsample_generic_Nd_kernel_impl<1, scale_t, HelperInterpLinear>( + upsample_non_separable_Nd_kernel_impl<1, scale_t, HelperInterpLinear>( output, input, align_corners, {scales_w}); } @@ -1739,17 +1786,17 @@ void upsample_bilinear2d_kernel_impl_float( std::optional scales_h, std::optional scales_w) { - // See note above about _use_vectorized_kernel_cond_2d(output, input). The extra cond is present + // See note above about _use_channels_last_kernel_2d(output, input). The extra cond is present // because benchmarks showed that with only 1 thread, images (C == 3) were - // slightly faster with the vectorized kernel than with the generic one. + // slightly faster with the channels-last kernel than with the generic one. // That's not the case for masks though (C == 1), which strongly benefit from // using the generic kernel. - if ((_use_vectorized_kernel_cond_2d(output, input)) || (at::get_num_threads() == 1 && input.size(1) == 3)) { + if ((_use_channels_last_kernel_2d(output, input)) || (at::get_num_threads() == 1 && input.size(1) == 3)) { AT_DISPATCH_FLOATING_TYPES_AND2(kBFloat16, kHalf, input.scalar_type(), "upsample_bilinear2d_channels_last", [&] { - cpu_upsample_linear_channels_last(output, input, align_corners, {scales_h, scales_w}); + upsample_linear_channels_last(output, input, align_corners, {scales_h, scales_w}); }); } else { - upsample_generic_Nd_kernel_impl<2, scale_t, HelperInterpLinear>( + upsample_non_separable_Nd_kernel_impl<2, scale_t, HelperInterpLinear>( output, input, align_corners, {scales_h, scales_w}); } } @@ -1777,7 +1824,7 @@ void upsample_bilinear2d_kernel_impl( /*antialias=*/false); } #endif // CPU_CAPABILITY_AVX2 - return separable_upsample_generic_Nd_kernel_impl<2, scale_t, HelperInterpLinear>( + return upsample_separable_Nd_kernel_impl<2, scale_t, HelperInterpLinear>( output, input, align_corners, {scales_h, scales_w}, /*antialias=*/false); } @@ -1808,7 +1855,7 @@ void upsample_bilinear2d_aa_kernel_impl( } #endif // CPU_CAPABILITY_AVX2 } - return separable_upsample_generic_Nd_kernel_impl<2, scale_t, HelperInterpLinear>( + return upsample_separable_Nd_kernel_impl<2, scale_t, HelperInterpLinear>( output, input, align_corners, {scales_h, scales_w}, /*antialias=*/true); } @@ -1820,12 +1867,12 @@ void upsample_trilinear3d_kernel_impl( std::optional scales_d, std::optional scales_h, std::optional scales_w) { - if ((_use_vectorized_kernel_cond_3d(output, input))) { + if ((_use_channels_last_kernel_3d(output, input))) { AT_DISPATCH_FLOATING_TYPES_AND2(kBFloat16, kHalf, input.scalar_type(), "upsample_trilinear3d_channels_last", [&] { - cpu_upsample_linear_channels_last(output, input, align_corners, {scales_d, scales_h, scales_w}); + upsample_linear_channels_last(output, input, align_corners, {scales_d, scales_h, scales_w}); }); } else { - upsample_generic_Nd_kernel_impl<3, scale_t, HelperInterpLinear>( + upsample_non_separable_Nd_kernel_impl<3, scale_t, HelperInterpLinear>( output, input, align_corners, {scales_d, scales_h, scales_w}); } } @@ -1853,11 +1900,11 @@ void upsample_bicubic2d_kernel_impl( /*antialias=*/false); } #endif // CPU_CAPABILITY_AVX2 - return separable_upsample_generic_Nd_kernel_impl<2, scale_t, HelperInterpCubic>( + return upsample_separable_Nd_kernel_impl<2, scale_t, HelperInterpCubic>( output, input, align_corners, {scales_h, scales_w}, /*antialias=*/false); } - return upsample_generic_Nd_kernel_impl<2, scale_t, HelperInterpCubic>( + return upsample_non_separable_Nd_kernel_impl<2, scale_t, HelperInterpCubic>( output, input, align_corners, {scales_h, scales_w}); } @@ -1885,7 +1932,7 @@ void upsample_bicubic2d_aa_kernel_impl( } #endif // CPU_CAPABILITY_AVX2 } - return separable_upsample_generic_Nd_kernel_impl<2, scale_t, HelperInterpCubic>( + return upsample_separable_Nd_kernel_impl<2, scale_t, HelperInterpCubic>( output, input, align_corners, {scales_h, scales_w}, /*antialias=*/true); } @@ -1894,7 +1941,7 @@ template < typename scalar_t, typename scale_type, class F> -void cpu_upsample_genNd_backward_aa( +void upsample_separable_Nd_backward_aa( const Tensor& grad_input_, const Tensor& grad_output_, bool align_corners, @@ -2012,7 +2059,7 @@ void upsample_bilinear2d_aa_backward_kernel_impl( std::optional scales_w) { AT_DISPATCH_FLOATING_TYPES( grad_output.scalar_type(), "upsample_bilinear2d_aa_backward_cpu", [&] { - cpu_upsample_genNd_backward_aa( + upsample_separable_Nd_backward_aa( grad_input, grad_output, align_corners, {scales_h, scales_w}); }); } @@ -2025,7 +2072,7 @@ void upsample_bicubic2d_aa_backward_kernel_impl( std::optional scales_w) { AT_DISPATCH_FLOATING_TYPES( grad_output.scalar_type(), "upsample_bicubic2d_aa_backward_cpu", [&] { - cpu_upsample_genNd_backward_aa( + upsample_separable_Nd_backward_aa( grad_input, grad_output, align_corners, {scales_h, scales_w}); }); } diff --git a/aten/src/ATen/native/cpu/UpSampleKernelAVXAntialias.h b/aten/src/ATen/native/cpu/UpSampleKernelAVXAntialias.h index c1bf79dfa44e6..3c351b05dc92f 100644 --- a/aten/src/ATen/native/cpu/UpSampleKernelAVXAntialias.h +++ b/aten/src/ATen/native/cpu/UpSampleKernelAVXAntialias.h @@ -176,7 +176,7 @@ void ImagingResampleHorizontal( // // TODO: we may want to merge that into the fallback code (currently called - // basic_loop_aa_horizontal) + // basic_loop_separable_1d_horizontal) // Although this may not be needed if / when we port all this code to use // Vec.h since this would potentially give us another fall-back implem @@ -252,7 +252,7 @@ void ImagingResampleVertical( // oB[xoffset + i] = b[xoffset + ymin[i]] * w[i, 0] + ... + b[xoffset + ymin[i] + (K-1) * xsize] * w[i, K-1] // TODO: we may want to merge that into the fallback code (currently called - // basic_loop_aa_vertical) + // basic_loop_separable_1d_vertical) // Although this may not be needed if / when we port all this code to use // Vec.h since this would potentially give us another fall-back implem const int16_t* kk = (int16_t*)(vert_indices_weights[3].const_data_ptr()); @@ -289,7 +289,7 @@ void ImagingResampleVertical( // mode for uint8 dtype when C <= 4, with or without antialias. The // implem is based on PIL-SIMD. // Its equivalent implementation (fallback) for when AVX isn't supported or when -// C > 4 is separable_upsample_generic_Nd_kernel_impl() There are a bunch of +// C > 4 is upsample_separable_Nd_kernel_impl() There are a bunch of // future improvement that can be done: look for the TODOs in this file. // For details on how the weights are computed and how the multiplications are // run on int (instead of float weights), see From ed60e5655203154ae48a06ba8174e64404f26dd2 Mon Sep 17 00:00:00 2001 From: "Deng, Daisy" Date: Mon, 23 Mar 2026 02:49:06 +0000 Subject: [PATCH 0007/1172] [XPU] Enable skipped inductor test on Intel GPU - generalize code and enable xpu in efficient attention and flash attention test (#174056) In order to enable xpu for efficient attention and fash attention test: 1. Enabled xpu in checkers: - evaluate_platform_supports_bf16 - evaluate_platform_supports_efficient_attention - evaluate_platform_supports_flash_attention - evaluate_platform_supports_fp8 2. generalize test code with torch.accelerator, requires_gpu_and_triton, 3. remove some skipIfXpu according to latest status 4. add skipIfXpu for xpu limitations 5. updated test/test_scaled_matmul_cuda.py to generalize new enabled cases by flag PLATFORM_SUPPORTS_FLASH_ATTENTION general 6. updated test/test_fake_tensor.py to generalize new enabled case by flag PLATFORM_SUPPORTS_FLASH_ATTENTION Pull Request resolved: https://github.com/pytorch/pytorch/pull/174056 Approved by: https://github.com/etaf, https://github.com/chuanqi129, https://github.com/jansel ghstack dependencies: #174053, #174054, #174055 --- test/dynamo/test_ctx_manager.py | 30 +- test/export/test_export.py | 14 +- test/inductor/test_aot_inductor.py | 35 +- test/inductor/test_cuda_repro.py | 518 +++++++++++++++---------- test/inductor/test_torchinductor.py | 11 +- test/test_fake_tensor.py | 16 +- test/test_scaled_matmul_cuda.py | 5 +- torch/testing/_internal/common_cuda.py | 10 +- 8 files changed, 375 insertions(+), 264 deletions(-) diff --git a/test/dynamo/test_ctx_manager.py b/test/dynamo/test_ctx_manager.py index f03759df31f93..88caf0223548e 100644 --- a/test/dynamo/test_ctx_manager.py +++ b/test/dynamo/test_ctx_manager.py @@ -26,6 +26,8 @@ ) +device_type = acc.type if (acc := torch.accelerator.current_accelerator()) else "cpu" + try: from . import test_functions except ImportError: @@ -729,7 +731,7 @@ def test_autocast_sdpa(self): class MyModule(torch.nn.Module): def forward(self, query, key, value): with torch.autocast("cpu"): - with torch.autocast("cuda", dtype=torch.float32): + with torch.autocast(device_type, dtype=torch.float32): out = F.scaled_dot_product_attention( query, key, value, None, 0.0, True ) @@ -740,13 +742,31 @@ def forward(self, query, key, value): seq_len_k = 1 head_dim = 8 query = torch.ones( - 1, 8, seq_len_q, head_dim, device="cuda", dtype=dtype, requires_grad=True + 1, + 8, + seq_len_q, + head_dim, + device=device_type, + dtype=dtype, + requires_grad=True, ) key = torch.ones( - 1, 8, seq_len_k, head_dim, device="cuda", dtype=dtype, requires_grad=True + 1, + 8, + seq_len_k, + head_dim, + device=device_type, + dtype=dtype, + requires_grad=True, ) value = torch.ones( - 1, 8, seq_len_k, head_dim, device="cuda", dtype=dtype, requires_grad=True + 1, + 8, + seq_len_k, + head_dim, + device=device_type, + dtype=dtype, + requires_grad=True, ) module = MyModule() @@ -760,7 +780,7 @@ def forward(self, query, key, value): self.assertEqual(compiled.device, real_device) self.assertEqual(compiled.dtype, real_dtype) - self.assertEqual(compiled.device.type, "cuda") + self.assertEqual(compiled.device.type, device_type) self.assertEqual(compiled.device.index, 0) self.assertEqual(compiled.dtype, torch.float32) diff --git a/test/export/test_export.py b/test/export/test_export.py index 3da0eb33bd53c..bcce9c1836b30 100755 --- a/test/export/test_export.py +++ b/test/export/test_export.py @@ -127,6 +127,8 @@ from torch.export import export +device_type = acc.type if (acc := torch.accelerator.current_accelerator()) else "cpu" + torch.library.define("testlib::returns_tensor_symint", "(Tensor x) -> (Tensor, SymInt)") torch.library.define( "testlib::foo", @@ -11076,7 +11078,6 @@ def forward(self, x): export_res = decomposed_ep.module()(x) self.assertTrue(export_res.size() == exp_res.size()) - @skipIfXpu def test_export_with_fake_tensor_inputs_on_cuda_devices(self): fake_mode = torch._subclasses.fake_tensor.FakeTensorMode() @@ -11095,9 +11096,9 @@ def forward(self, x): model = Model() # Manually set the fake_device of fake tensors. - x.fake_device = torch.device("cuda:0") + x.fake_device = torch.device(f"{device_type}:0") for n, p in model.named_parameters(): - p.fake_device = torch.device("cuda:0") + p.fake_device = torch.device(f"{device_type}:0") # Need to set all the requires_grad of tensors to False, because fake_tensor with CUDA device # doesn't quite work well with aot_autograd right now due to some logic fails @@ -17607,6 +17608,7 @@ def forward(self, q, k, v): ep = torch.export.export(ScaledDotProductAttention(), (q, k, v)) ep.run_decompositions() + @skipIfXpu(msg="scaled_dot_product_attention issue on xpu") @skipIfCrossRef @unittest.skipIf( not PLATFORM_SUPPORTS_FLASH_ATTENTION, @@ -17631,9 +17633,9 @@ def forward(self, q, k, v): ) return attn_output - q = torch.randn(1, 16, 16, 64, dtype=torch.bfloat16, device="cuda") - k = torch.randn(1, 16, 16, 64, dtype=torch.bfloat16, device="cuda") - v = torch.randn(1, 16, 16, 64, dtype=torch.bfloat16, device="cuda") + q = torch.randn(1, 16, 16, 64, dtype=torch.bfloat16, device=device_type) + k = torch.randn(1, 16, 16, 64, dtype=torch.bfloat16, device=device_type) + v = torch.randn(1, 16, 16, 64, dtype=torch.bfloat16, device=device_type) ep = torch.export.export( ScaledDotProductAttention(), (q, k, v) diff --git a/test/inductor/test_aot_inductor.py b/test/inductor/test_aot_inductor.py index ebfb52c57ff46..1b2bf752bcb5f 100644 --- a/test/inductor/test_aot_inductor.py +++ b/test/inductor/test_aot_inductor.py @@ -1332,12 +1332,11 @@ def forward(self, x, y): @unittest.skipIf( not PLATFORM_SUPPORTS_FP8, - "FP8 is only supported on H100+, SM 8.9 and MI300+ devices", + "FP8 is only supported on H100+, SM 8.9 and MI300+ devices or XPU", ) - @skipIfXpu def test_fp8(self): # cuda only - if self.device != "cuda": + if self.device not in ("cuda", "xpu"): return class Model(torch.nn.Module): @@ -1381,7 +1380,6 @@ def forward(self, x, weight, bias, scale_a, scale_b): not PLATFORM_SUPPORTS_FP8_GROUPED_GEMM, "scaled_grouped_mm is only supported on SM90 and MI300+ devices", ) - @skipIfXpu def test_scaled_grouped_mm(self): # Test torch._scaled_grouped_mm AOTI lowering # cuda only @@ -1445,9 +1443,8 @@ def forward(self, x, weight, scale_a, scale_b, offsets): @unittest.skipIf( not PLATFORM_SUPPORTS_FP8, - "FP8 is only supported on H100+, SM 8.9 and MI300+ devices", + "FP8 is only supported on H100+, SM 8.9 and MI300+ devices or XPU", ) - @skipIfXpu def test_fp8_view_of_param(self): # cuda only if self.device != GPU_TYPE: @@ -2083,7 +2080,7 @@ def forward(self, values, repeats, mask, embeddings, x, y, z, lst): } self.check_model(Repro(), example_inputs, dynamic_shapes=spec) - @skipIfXpu(msg="_scaled_dot_product_flash_attention is not supported on XPU yet") + @skipIfXpu(msg="FlashAttentionForward headdim limitation on xpu - xpu-ops: 2698") @unittest.skipIf( not PLATFORM_SUPPORTS_FLASH_ATTENTION, "Some archs don't support flash SDPA" ) @@ -2534,8 +2531,8 @@ def test_cond_cpu_predicate_cuda_operands(self, max_autotune): determined device from [predicate] + operands, causing CPU predicates to force CUDA outputs onto CPU during autotuning. """ - if self.device != "cuda": - raise unittest.SkipTest("requires CUDA") + if self.device != "cuda" and self.device != "xpu": + raise unittest.SkipTest("requires CUDA or XPU") class Model(torch.nn.Module): def __init__(self, input_dim=4, hidden_dim=8): @@ -5062,7 +5059,7 @@ def forward(self, inputs): @unittest.skipIf( not PLATFORM_SUPPORTS_FP8, - "FP8 is only supported on H100+, SM 8.9 and MI300+ devices", + "FP8 is only supported on H100+, SM 8.9 and MI300+ devices or XPU", ) @patch.dict(os.environ, {"AOTI_RUNTIME_CHECK_INPUTS": "1"}) def test_runtime_checks_fp8(self): @@ -5805,9 +5802,8 @@ def forward(self, x): @unittest.skipIf( not PLATFORM_SUPPORTS_FP8, - "FP8 is only supported on H100+, SM 8.9 and MI300+ devices", + "FP8 is only supported on H100+, SM 8.9 and MI300+ devices or XPU", ) - @skipIfXpu def test_aoti_debug_printer_fp8_dtype(self): if self.device != GPU_TYPE: raise unittest.SkipTest("requires GPU") @@ -6960,7 +6956,6 @@ def forward(self, x): example_inputs = (torch.randn(500, device=self.device),) self.check_model(model, example_inputs) - @skipIfXpu def test_conv3d(self): if self.device != GPU_TYPE or not is_big_gpu(): raise unittest.SkipTest("requires modern GPU to run max-autotune") @@ -7825,8 +7820,8 @@ def test_codegen_int_array_var_fix_memory_leak(self): """ Fix https://github.com/pytorch/pytorch/issues/167630 """ - if self.device != "cuda": - raise unittest.SkipTest("test is only for cuda") + if self.device not in ("cuda", "xpu"): + raise unittest.SkipTest("test is only for cuda or xpu") def make_mlp(in_dim=128, hidden=256, out_dim=64, depth=3): layers = [] @@ -7847,7 +7842,7 @@ def make_mlp(in_dim=128, hidden=256, out_dim=64, depth=3): allocated_memory = [] for _ in range(3): - torch.cuda.reset_peak_memory_stats() + torch.accelerator.reset_peak_memory_stats() model = make_mlp(in_dim, hidden, out_dim, depth).to(self.device) example_inputs = (torch.randn(batch, in_dim, device=self.device),) @@ -7858,10 +7853,10 @@ def make_mlp(in_dim=128, hidden=256, out_dim=64, depth=3): torch._inductor.aoti_compile_and_package(ep) del model, example_inputs, ep - torch.cuda.synchronize() - torch.cuda.empty_cache() + torch.accelerator.synchronize() + torch.accelerator.empty_cache() gc.collect() - allocated_memory.append(torch.cuda.memory_allocated()) + allocated_memory.append(torch.accelerator.memory_allocated()) self.assertTrue(allocated_memory[1] == allocated_memory[2]) @@ -8386,8 +8381,6 @@ def fail_gpu(suffixes: tuple[str, ...], is_skip=False): "test_quantized_linear": fail_gpu(("cuda", "xpu")), "test_quanatized_int8_linear": fail_gpu(("cuda", "xpu")), "test_quantized_linear_bias_none": fail_gpu(("cuda", "xpu")), - # No scaled_dot_product_efficient_attention implementation for XPU yet. - "test_scaled_dot_product_efficient_attention": fail_gpu(("xpu",)), } MPS_TEST_FAILURES = { diff --git a/test/inductor/test_cuda_repro.py b/test/inductor/test_cuda_repro.py index f43960f064e79..e728cbb80c85e 100644 --- a/test/inductor/test_cuda_repro.py +++ b/test/inductor/test_cuda_repro.py @@ -44,8 +44,11 @@ parametrize, skipIfRocm, skipIfRocmArch, + skipIfXpu, + TEST_CUDA, TEST_WITH_ASAN, TEST_WITH_ROCM, + TEST_XPU, xfailIfROCm, ) from torch.testing._internal.inductor_utils import IS_BIG_GPU @@ -89,14 +92,21 @@ aten = torch.ops.aten +device_type = ( + acc.type + if (acc := torch.accelerator.current_accelerator(check_available=True)) + else "cpu" +) + + @instantiate_parametrized_tests class CudaReproTests(TestCase): - device = "cuda" + device = device_type common = check_model_cuda def test_mm_out_dtype_compile(self): - a = torch.randn(1, 3, device="cuda", dtype=torch.float16) - b = torch.randn(3, 2, device="cuda", dtype=torch.float16) + a = torch.randn(1, 3, device=device_type, dtype=torch.float16) + b = torch.randn(3, 2, device=device_type, dtype=torch.float16) def fn(x, y): return torch.mm(x, y, out_dtype=torch.float32) @@ -135,7 +145,8 @@ def forward( (torch.Size([512, 768]), torch.float16), ] inps = [torch.zeros(())] + [ - torch.ones(shape, dtype=dtype, device="cuda") for (shape, dtype) in inps + torch.ones(shape, dtype=dtype, device=device_type) + for (shape, dtype) in inps ] mod = make_fx(forward)(*inps) compiled = compile_fx_inner(mod, inps) @@ -167,8 +178,8 @@ def forward(self, x: torch.Tensor): return {"ten0": out0, "ten1": out1} torch.manual_seed(0) - model = ReproModule().cuda() - inputs = torch.randn(36, 9, 7, 16, device="cuda", requires_grad=True) + model = ReproModule().to(device_type) + inputs = torch.randn(36, 9, 7, 16, device=device_type, requires_grad=True) eager_out = model(inputs) compiled_model = torch.compile( @@ -207,15 +218,19 @@ def fn( scale=None, ) - query = torch.randn(batch_size, num_heads, seq_len, head_dim, device="cuda") - key = torch.randn(batch_size, num_heads, seq_len, head_dim, device="cuda") - value = torch.randn(batch_size, num_heads, seq_len, head_dim, device="cuda") + query = torch.randn( + batch_size, num_heads, seq_len, head_dim, device=device_type + ) + key = torch.randn(batch_size, num_heads, seq_len, head_dim, device=device_type) + value = torch.randn( + batch_size, num_heads, seq_len, head_dim, device=device_type + ) - input_tensor = torch.rand([2, 1, seq_len, 1], device="cuda") + input_tensor = torch.rand([2, 1, seq_len, 1], device=device_type) out, code = run_and_get_code(torch.compile(fn), query, key, value, input_tensor) - input_tensor2 = torch.rand([2, 32, seq_len, seq_len], device="cuda").copy_( + input_tensor2 = torch.rand([2, 32, seq_len, seq_len], device=device_type).copy_( input_tensor ) # even though the last dim is broadcasted, needs stride 1 for alignment @@ -238,10 +253,18 @@ def test_effn_attn_bias_padding_misaligned(self): torch._dynamo.reset() bsz = 32 - q = torch.randn(bsz, 16, seqlen, 64, dtype=torch.bfloat16, device="cuda") - k = torch.randn(bsz, 16, seqlen, 64, dtype=torch.bfloat16, device="cuda") - v = torch.randn(bsz, 16, seqlen, 64, dtype=torch.bfloat16, device="cuda") - mask = torch.ones([bsz, 1, seqlen, seqlen], dtype=torch.bool, device="cuda") + q = torch.randn( + bsz, 16, seqlen, 64, dtype=torch.bfloat16, device=device_type + ) + k = torch.randn( + bsz, 16, seqlen, 64, dtype=torch.bfloat16, device=device_type + ) + v = torch.randn( + bsz, 16, seqlen, 64, dtype=torch.bfloat16, device=device_type + ) + mask = torch.ones( + [bsz, 1, seqlen, seqlen], dtype=torch.bool, device=device_type + ) inputs = [q, k, v, mask] def f(q, k, v, mask): @@ -266,8 +289,12 @@ def test_input_channels_last(self): m = torch.nn.Sequential( torch.nn.Conv2d(3, 3, 1, 1), ToTuple(), - ).cuda() - inp = torch.randn([2, 3, 16, 16]).to(memory_format=torch.channels_last).cuda() + ).to(device_type) + inp = ( + torch.randn([2, 3, 16, 16]) + .to(memory_format=torch.channels_last) + .to(device_type) + ) self.common( m, @@ -291,10 +318,12 @@ def forward(self, x, y): return [permute, add] inps = [ - rand_strided((12, 3, 512, 64), (64, 196608, 768, 1), torch.float32, "cuda"), + rand_strided( + (12, 3, 512, 64), (64, 196608, 768, 1), torch.float32, device_type + ), rand_strided((), (), torch.int64, "cpu"), ] - mod = make_fx(Repro().to(device="cuda"))(*inps) + mod = make_fx(Repro().to(device=device_type))(*inps) compiled = compile_fx_inner(mod, inps) compiled(inps) @@ -305,7 +334,7 @@ def test_backward_context(self): def fn(x): return x * 3 - x = torch.randn(4, device="cuda", requires_grad=True) + x = torch.randn(4, device=device_type, requires_grad=True) gO = torch.rand_like(x) opt_fn = torch.compile(fn) out = opt_fn(x) @@ -317,7 +346,7 @@ def forward(): randn = torch.ops.aten.randn.default( [12, 64, 1, 64], dtype=torch.float32, - device=torch.device(type="cuda", index=0), + device=torch.device(type=device_type, index=0), pin_memory=False, ) unsqueeze_default_2 = torch.ops.aten.unsqueeze.default(randn, -1) @@ -325,9 +354,9 @@ def forward(): mod = make_fx(forward)() compiled = compile_fx_inner(mod, ()) - if compiled([])[0].device.type != "cuda": + if compiled([])[0].device.type != device_type: raise AssertionError( - f"Expected device type 'cuda', got {compiled([])[0].device.type!r}" + f"Expected device type {device_type}, got {compiled([])[0].device.type!r}" ) @config.patch({"triton.cudagraphs": True}) @@ -343,7 +372,7 @@ def forward(self): 1, dtype=torch.float32, layout=torch.strided, - device=torch.device(type="cuda", index=0), + device=torch.device(type=device_type, index=0), pin_memory=False, ) full_1 = torch.ops.aten.full.default( @@ -351,7 +380,7 @@ def forward(self): 0, dtype=torch.int64, layout=torch.strided, - device=torch.device(type="cuda", index=0), + device=torch.device(type=device_type, index=0), pin_memory=False, ) return (full_1, full) @@ -366,8 +395,8 @@ def fn(x, y): return x + y inputs = ( - rand_strided((5, 5, 5, 5), (0, 5, 0, 1), device="cuda"), - rand_strided((5, 5, 5, 5), (0, 5, 0, 1), device="cuda"), + rand_strided((5, 5, 5, 5), (0, 5, 0, 1), device=device_type), + rand_strided((5, 5, 5, 5), (0, 5, 0, 1), device=device_type), ) self.assertTrue(same(fn(*inputs), inputs[0] + inputs[1])) @@ -386,14 +415,14 @@ def fn(x, y): return r, r.size(0) inputs = ( - torch.randn((5, 5), device="cuda"), - torch.randn((5, 5), device="cuda"), + torch.randn((5, 5), device=device_type), + torch.randn((5, 5), device=device_type), ) self.assertTrue(same(fn(*inputs), (inputs[0] + inputs[1], 5))) inputs = ( - torch.randn((6, 6), device="cuda"), - torch.randn((6, 6), device="cuda"), + torch.randn((6, 6), device=device_type), + torch.randn((6, 6), device=device_type), ) self.assertTrue(same(fn(*inputs), (inputs[0] + inputs[1], 6))) @@ -412,13 +441,13 @@ def max(x): print(f"compile {ms_c=:.03f}, eager {ms_eager=:.03f}") def test_split_reduction_transposed(self): - x = torch.randn(4096, 8192, dtype=torch.bfloat16, device="cuda") + x = torch.randn(4096, 8192, dtype=torch.bfloat16, device=device_type) x = x.t().contiguous().t() self._test_split_reduction_impl(x) def test_split_reduction_channels_last(self): - x = torch.randn(4096, 8192, dtype=torch.bfloat16, device="cuda") + x = torch.randn(4096, 8192, dtype=torch.bfloat16, device=device_type) x = x.reshape([256, 256, 256, 2]).to(memory_format=torch.channels_last) self._test_split_reduction_impl(x) @@ -439,7 +468,7 @@ def forward(): pin_memory=False, ) device_put_3 = torch.ops.prims.device_put.default( - full_1, device(type="cuda", index=0) + full_1, device(type=device_type, index=0) ) full_1 = None @@ -460,13 +489,13 @@ def forward(): view_15 = torch.ops.aten.reshape.default(clone, [1536, 1536]) clone = None scalar_tensor = torch.ops.aten.scalar_tensor.default( - -inf, dtype=torch.float16, device=device(type="cuda", index=0) + -inf, dtype=torch.float16, device=device(type=device_type, index=0) ) scalar_tensor_1 = torch.ops.aten.scalar_tensor.default( 0.0, dtype=torch.float16, layout=torch.strided, - device=device(type="cuda", index=0), + device=device(type=device_type, index=0), ) where = torch.ops.aten.where.self(view_15, scalar_tensor_1, scalar_tensor) view_15 = scalar_tensor_1 = scalar_tensor = None @@ -482,7 +511,9 @@ def test_emulate_low_precision(self): def foo(x): return torch.nn.functional.gelu(x) * 10.0 - inp = torch.rand([32], device="cuda", requires_grad=True, dtype=torch.bfloat16) + inp = torch.rand( + [32], device=device_type, requires_grad=True, dtype=torch.bfloat16 + ) out, codes = run_fw_bw_and_get_code(lambda: torch.compile(foo)(inp)) # fwd, backward @@ -525,8 +556,8 @@ def fn(x, y): return x + y inputs = ( - rand_strided((5, 5, 5, 5), (0, 5, 0, 1), device="cuda"), - rand_strided((5, 5, 5, 5), (0, 5, 0, 1), device="cuda"), + rand_strided((5, 5, 5, 5), (0, 5, 0, 1), device=device_type), + rand_strided((5, 5, 5, 5), (0, 5, 0, 1), device=device_type), ) self.assertTrue(same(fn(*inputs), inputs[0] + inputs[1])) @@ -547,11 +578,11 @@ def forward(self, x): from copy import deepcopy - model = Repro().cuda() + model = Repro().to(device_type) model_ref = deepcopy(model) model_opt = torch.compile(model, backend="inductor") - input = torch.randn(10, 10, device="cuda", requires_grad=True) + input = torch.randn(10, 10, device=device_type, requires_grad=True) for _ in range(2): output_ref = model_ref(input) @@ -574,7 +605,7 @@ def foo(x): foo_opt = torch.compile(foo, backend="inductor") - inpt = torch.randn(10, 10, device="cuda", requires_grad=True) + inpt = torch.randn(10, 10, device=device_type, requires_grad=True) # TODO: this is broken, fix later # out = foo_opt(inpt) # out.add_(2) @@ -602,14 +633,14 @@ def forward(self, start_positions: torch.Tensor, x: torch.Tensor): ) return cross_entropy - mod = Repro().cuda() + mod = Repro().to(device_type) opt_mod = torch.compile(mod, backend="inductor") mod.eval() opt_mod.eval() args = [ - ((1,), (1,), torch.int64, "cuda", False), - ((1, 128, 768), (98304, 768, 1), torch.float32, "cuda", True), + ((1,), (1,), torch.int64, device_type, False), + ((1, 128, 768), (98304, 768, 1), torch.float32, device_type, True), ] args = [ rand_strided(sh, st, dt, dev).requires_grad_(rg) @@ -628,7 +659,7 @@ def forward(add_1): getitem_1 = var_mean[1] return getitem_1 - x = torch.randn(1, 8, 768, device="cuda") + x = torch.randn(1, 8, 768, device=device_type) correct = forward(x) actual = torch.compile(forward, fullgraph=True)(x) self.assertEqual(actual, correct) @@ -640,7 +671,7 @@ def forward(x): 0, dtype=torch.float64, layout=torch.strided, - device="cuda", + device=device_type, pin_memory=False, ) return x + full_10.to("cpu") @@ -657,7 +688,11 @@ def test_autotune_inplace_kernel(self): https://github.com/triton-lang/triton/issues/781 https://github.com/pytorch/torchdynamo/issues/1670 """ - from torch._C import _cuda_getCurrentRawStream as get_cuda_stream + if TEST_XPU: + from torch._C import _xpu_getCurrentRawStream as get_gpu_stream + else: + from torch._C import _cuda_getCurrentRawStream as get_gpu_stream + from torch._inductor.runtime.hints import AttrsDescriptorWrapper, HeuristicType from torch._inductor.runtime.triton_heuristics import CachingAutotuner from torch._inductor.utils import triton_version_uses_attrs_dict @@ -695,7 +730,7 @@ def decorator(fn): "in_ptr0": "*fp32", "xnumel": "i32", }, - "device": DeviceProperties.create(torch.device("cuda")), + "device": DeviceProperties.create(torch.device(device_type)), "configs": [ AttrsDescriptorWrapper(divisible_by_16=(0, 1), equal_to_1=()) ], @@ -714,11 +749,11 @@ def kernel(in_out_ptr0, in_ptr0, xnumel, XBLOCK: tl.constexpr): tl.store(in_out_ptr0 + offsets, output, mask=mask) xnumel = 384 - in0 = rand_strided((xnumel,), (1,), device="cuda", dtype=torch.float32) - inout1 = rand_strided((xnumel,), (1,), device="cuda", dtype=torch.float32) + in0 = rand_strided((xnumel,), (1,), device=device_type, dtype=torch.float32) + inout1 = rand_strided((xnumel,), (1,), device=device_type, dtype=torch.float32) inout2 = inout1.clone() - stream0 = get_cuda_stream(0) + stream0 = get_gpu_stream(0) kernel.run(inout1, in0, xnumel, stream=stream0) kernel.run(inout2, in0, xnumel, stream=stream0) @@ -734,7 +769,7 @@ def forward(pred_objectness_logits_3_: torch.Tensor): getitem_12 = sort_3[0] return getitem_12 - args = [((1, 100), (0, 1), torch.float16, "cuda", False)] + args = [((1, 100), (0, 1), torch.float16, device_type, False)] args = [ rand_strided(sh, st, dt, dev).requires_grad_(rg) for (sh, st, dt, dev, rg) in args @@ -751,7 +786,7 @@ def fn(a): zero = torch.zeros((16,), device=a.device, dtype=torch.int64) return (a[zero],) - a = torch.randn((8,), dtype=torch.float32, device="cuda") + a = torch.randn((8,), dtype=torch.float32, device=device_type) fn_optimized = torch.compile(fn, backend="inductor") if not same(fn(a), fn_optimized(a)): @@ -768,8 +803,8 @@ def fn(x, y): out = torch.ops.aten.multiply(y, squeeze) return (out,) - a = torch.zeros((1, 128), dtype=torch.int64, device="cuda") - b = torch.zeros((1, 128), dtype=torch.int64, device="cuda") + a = torch.zeros((1, 128), dtype=torch.int64, device=device_type) + b = torch.zeros((1, 128), dtype=torch.int64, device=device_type) fn_optimized = torch.compile(fn, backend="inductor") if not same(fn(a, b), fn_optimized(a, b)): @@ -779,7 +814,9 @@ def test_simplify_dims(self): def fn(a): return (a + 1,) - self.common(fn, (torch.randn(2, 3, 10, 5, 6, device="cuda")[:, :, 2::2, :, :],)) + self.common( + fn, (torch.randn(2, 3, 10, 5, 6, device=device_type)[:, :, 2::2, :, :],) + ) @config.patch(permute_fusion=True) def test_permute_fusion(self): @@ -792,8 +829,8 @@ def forward(self, view, reshape_2): return (bmm,) args = [ - ((1024, 642, 160), (102720, 160, 1), torch.float32, "cuda", True), - ((1024, 642, 20), (12840, 20, 1), torch.float32, "cuda", True), + ((1024, 642, 160), (102720, 160, 1), torch.float32, device_type, True), + ((1024, 642, 20), (12840, 20, 1), torch.float32, device_type, True), ] args = [ rand_strided(sh, st, dt, dev).requires_grad_(rg) @@ -813,10 +850,10 @@ def fn(x, y): aten.add_.Tensor(x, y, alpha=0.55) return (x,) - x1 = torch.zeros(2, 3, 4, 10, device="cuda") - x2 = torch.zeros(2, 3, 4, 10, device="cuda") - x3 = torch.zeros(2, 3, 4, 10, device="cuda") - y = torch.randn(2, 3, 4, 10, device="cuda").to( + x1 = torch.zeros(2, 3, 4, 10, device=device_type) + x2 = torch.zeros(2, 3, 4, 10, device=device_type) + x3 = torch.zeros(2, 3, 4, 10, device=device_type) + y = torch.randn(2, 3, 4, 10, device=device_type).to( memory_format=torch.channels_last ) fn_fx = make_fx(fn)(x1, y) @@ -832,9 +869,11 @@ def foo(x, y, z): a = x @ y return a.unsqueeze(0).unsqueeze(0) + z - x = torch.zeros(5, 5, device="cuda") - y = torch.zeros(5, 5, device="cuda") - z = torch.zeros(1, 1, 5, 5, device="cuda").to(memory_format=torch.channels_last) + x = torch.zeros(5, 5, device=device_type) + y = torch.zeros(5, 5, device=device_type) + z = torch.zeros(1, 1, 5, 5, device=device_type).to( + memory_format=torch.channels_last + ) self.common( foo, (x, y, z), @@ -851,16 +890,26 @@ def fn(x, w, b): x = called_inside_compile(x, w, b) return called_inside_compile(x, w, b) - w = torch.rand(3, 3, device="cuda") - b = torch.rand(3, device="cuda") - x = torch.rand(3, device="cuda") + w = torch.rand(3, 3, device=device_type) + b = torch.rand(3, device=device_type) + x = torch.rand(3, device=device_type) + + def record_memory_history(value: bool): + if torch.xpu.is_available(): + torch.xpu._record_memory_history(value) + else: + torch.cuda.memory._record_memory_history(value) + try: - torch.cuda.memory.empty_cache() - torch.cuda.memory._record_memory_history(True) + torch.accelerator.memory.empty_cache() + record_memory_history(True) r = fn(x, w, b) finally: - torch.cuda.memory._record_memory_history(False) - snapshot = str(torch.cuda.memory._snapshot()) + record_memory_history(False) + if torch.xpu.is_available(): + snapshot = str(torch.xpu.memory._snapshot()) + else: + snapshot = str(torch.cuda.memory._snapshot()) self.assertTrue("called_inside_compile" in snapshot) def test_negative_arange_dynamic_shapes(self): @@ -893,10 +942,14 @@ def forward(self, enc_out: torch.Tensor, dec_in: torch.Tensor): padmask = dec_in == 0 dec_mask = padmask.unsqueeze(-1) == padmask.unsqueeze(-2) dec_mask = dec_mask.to(dtype=torch.float32) - dec_mask = dec_mask.tril(diagonal=0).cuda() + dec_mask = dec_mask.tril(diagonal=0).to(device_type) - q_pos = torch.arange(dec_in.size(1), dtype=torch.long, device="cuda") - k_pos = torch.arange(dec_in.size(1), dtype=torch.long, device="cuda") + q_pos = torch.arange( + dec_in.size(1), dtype=torch.long, device=device_type + ) + k_pos = torch.arange( + dec_in.size(1), dtype=torch.long, device=device_type + ) rel_pos = k_pos[None, :] - q_pos[:, None] values = rel_pos.abs().neg().unsqueeze(0).unsqueeze(0) dec_bias = values * self.scales @@ -907,14 +960,15 @@ def forward(self, enc_out: torch.Tensor, dec_in: torch.Tensor): out = self.dec_layer(out, enc_out, tgt_mask=dec_mask) return self.head(out) - mod = Repro().cuda() + mod = Repro().to(device_type) opt_mod = torch.compile(mod, backend="inductor", dynamic=True) mod.eval() opt_mod.eval() - enc_out = torch.rand(1, 512, 256).cuda() + enc_out = torch.rand(1, 512, 256).to(device_type) dec_inputs = [ - torch.randint(0, 512, (1, i + 1), dtype=torch.long).cuda() for i in range(8) + torch.randint(0, 512, (1, i + 1), dtype=torch.long).to(device_type) + for i in range(8) ] for dec_inp in dec_inputs: @@ -928,9 +982,9 @@ def fn(arg3_1, relu, permute_1): return (cat_2,) args = [ - ((96,), (1,), torch.float32, "cuda"), - ((10, 256), (256, 1), torch.float32, "cuda"), - ((256, 96), (1, 256), torch.float32, "cuda"), + ((96,), (1,), torch.float32, device_type), + ((10, 256), (256, 1), torch.float32, device_type), + ((256, 96), (1, 256), torch.float32, device_type), ] args = [rand_strided(sh, st, dt, dev) for (sh, st, dt, dev) in args] correct = fn(*args) @@ -958,7 +1012,7 @@ def forward(self, x): inp = x / y[..., None] return self.layer(inp) - x = torch.rand([4, 4], device="cuda") + x = torch.rand([4, 4], device=device_type) m = MyModule() opt_m = torch.compile(backend="inductor")(m) self.assertEqual(opt_m(x), m(x)) @@ -971,10 +1025,10 @@ def fn(arg3_1, arg3_2, relu, permute_1): return (cat_2,) args = [ - ((96,), (1,), torch.float32, "cuda"), - ((96,), (1,), torch.float32, "cuda"), - ((10, 256), (256, 1), torch.float32, "cuda"), - ((256, 96), (1, 256), torch.float32, "cuda"), + ((96,), (1,), torch.float32, device_type), + ((96,), (1,), torch.float32, device_type), + ((10, 256), (256, 1), torch.float32, device_type), + ((256, 96), (1, 256), torch.float32, device_type), ] args = [rand_strided(sh, st, dt, dev) for (sh, st, dt, dev) in args] correct = fn(*args) @@ -1003,7 +1057,9 @@ def test_normalize_norm_leq_one(self): def fn(x: torch.Tensor) -> torch.Tensor: return torch.nn.functional.normalize(x, dim=-1) - inp = torch.tensor([[3.799999, 0.0, 0.0]], device="cuda", dtype=torch.float32) + inp = torch.tensor( + [[3.799999, 0.0, 0.0]], device=device_type, dtype=torch.float32 + ) compiled = torch.compile(fn, backend="inductor", fullgraph=True) out = compiled(inp) norm = out.norm(dim=-1) @@ -1015,7 +1071,7 @@ def test_libdevice_routing(self): def foo(x): return x.exp() - inp = torch.ones(64, device="cuda").to(torch.float64) + inp = torch.ones(64, device=device_type).to(torch.float64) out, code = run_and_get_code(torch.compile(foo), inp) FileCheck().check("libdevice.exp").run(code[0]) @@ -1029,7 +1085,7 @@ def foo(x): def foo(x): return x.sigmoid() - inp = torch.ones(64, device="cuda").to(torch.float64) + inp = torch.ones(64, device=device_type).to(torch.float64) out, code = run_and_get_code(torch.compile(foo), inp) FileCheck().check("libdevice.exp").run(code[0]) self.assertEqual(foo(inp), out) @@ -1041,7 +1097,7 @@ def view_copy(target, source): assert source.dtype == torch.uint16 # noqa: S101 target.view(torch.uint16).copy_(source) - target = torch.ones(1024, dtype=torch.bfloat16, device="cuda") + target = torch.ones(1024, dtype=torch.bfloat16, device=device_type) source = torch.full_like(target, 4, dtype=torch.uint16) out = target.view(torch.uint16).copy_(source).clone() @@ -1055,7 +1111,7 @@ def forward(arg0_1): 1, dtype=torch.float32, layout=torch.strided, - device=torch.device(type="cuda", index=0), + device=torch.device(type=device_type, index=0), pin_memory=False, ) convert_element_type_1 = torch.ops.prims.convert_element_type.default( @@ -1073,23 +1129,23 @@ def forward(arg0_1): ) return [var_mean[0], var_mean[1], add_2] - emb = torch.randn([2050, 768], device="cuda") + emb = torch.randn([2050, 768], device=device_type) gm = make_fx(forward)(emb) opt = torch._inductor.compile_fx.compile_fx_inner(gm, [emb]) opt([emb]) - torch.cuda.synchronize() + torch.accelerator.synchronize() def test_deterministic_algorithms(self): N = 10000 @torch.compile def fn(idx, values): - x = torch.zeros(1, device="cuda") + x = torch.zeros(1, device=device_type) x[idx] += values return x - idx = torch.zeros(N, dtype=torch.int64, device="cuda") - values = torch.randn(N, device="cuda") + idx = torch.zeros(N, dtype=torch.int64, device=device_type) + values = torch.randn(N, device=device_type) r0 = fn(idx, values) with DeterministicGuard(True): @@ -1106,10 +1162,10 @@ def __init__(self) -> None: self.linear = nn.Linear(4, 4) def forward(self, data): - data = data.to("cuda") + data = data.to(device_type) return self.linear(data) - mod = Model().cuda().eval() + mod = Model().to(device_type).eval() with torch.no_grad(): self.common(mod, (torch.randn(4, 4),)) @@ -1125,7 +1181,7 @@ def forward(self, x): return self.dropout(y) mod = Repro() - x = torch.randn((512, 1, 4096), requires_grad=True, device="cuda") + x = torch.randn((512, 1, 4096), requires_grad=True, device=device_type) y = torch.compile(mod)(x) # Inductor claims the output layout of gelu's saved variable for # backwards will be (4096, 4096, 1) but in actuality it is (4096, @@ -1152,9 +1208,9 @@ def forward(inductor_seeds, mul_4, view_15): sub_3 = torch.ops.aten.sub.Tensor(add_5, getitem_3) return (sub_3,) - buf0 = torch.zeros((37,), dtype=torch.int64, device="cuda") - buf1 = torch.zeros((2, 512, 768), device="cuda") - buf2 = torch.zeros((2, 512, 768), device="cuda") + buf0 = torch.zeros((37,), dtype=torch.int64, device=device_type) + buf1 = torch.zeros((2, 512, 768), device=device_type) + buf2 = torch.zeros((2, 512, 768), device=device_type) forward(buf0, buf1, buf2) def test_issue100806(self): @@ -1174,7 +1230,7 @@ def forward(self, x): x = self.relu(x) return x - device = "cuda" + device = device_type batch_size = 2 x = torch.randn(batch_size, 10).to(device) func = Model().to(device) @@ -1194,8 +1250,8 @@ def fn(x, y): add = mean + y return add - x = torch.rand(4, 4, 4, 4, 4, 4, device="cuda") - y = torch.rand((), device="cuda") + x = torch.rand(4, 4, 4, 4, 4, 4, device=device_type) + y = torch.rand((), device=device_type) expect = fn(x, y) opt_fn = torch.compile(fn) @@ -1214,8 +1270,8 @@ def test_bucketize_dynamic_dense(self): def fn(values, offsets): return torch.bucketize(values, offsets) - values = torch.rand((64, 64), device="cuda") - offsets = torch.tensor([0.05, 0.1, 0.5, 0.8, 0.85, 0.95], device="cuda") + values = torch.rand((64, 64), device=device_type) + offsets = torch.tensor([0.05, 0.1, 0.5, 0.8, 0.85, 0.95], device=device_type) expect = fn(values, offsets) @@ -1253,9 +1309,9 @@ def fn(x: torch.Tensor, y: torch.Tensor, buckets: torch.Tensor) -> torch.Tensor: z = torch.mm(x, y) return torch.bucketize(z, buckets) - buckets = torch.arange(-100, 100, 10, device="cuda") - x = torch.randn(64, 64, device="cuda").clamp(-99, 99) - y = torch.randn(64, 64, device="cuda").clamp(-99, 99) + buckets = torch.arange(-100, 100, 10, device=device_type) + x = torch.randn(64, 64, device=device_type).clamp(-99, 99) + y = torch.randn(64, 64, device=device_type).clamp(-99, 99) opt_fn = torch.compile(fn, mode="max-autotune") @@ -1268,7 +1324,7 @@ def test_float64_constants(self): def fn(): # NOTE: tensors of all the same value are constant folded, so we # need a tensor with two distinct values - a = torch.tensor([1 / 10, 2 / 10], dtype=torch.float64, device="cuda") + a = torch.tensor([1 / 10, 2 / 10], dtype=torch.float64, device=device_type) return a * 2e50 cfn = torch.compile(fn) @@ -1308,10 +1364,15 @@ def fn(arg7_1, add_1, permute_2, select_scatter, slice_8): return (bmm,) args = [] - args.append(torch.randn((2, 1, 4, 1200, 4), dtype=torch.float16, device="cuda")) + args.append( + torch.randn((2, 1, 4, 1200, 4), dtype=torch.float16, device=device_type) + ) args.append( rand_strided( - (1, 4, 1000, 4), (16000, 4, 16, 1), dtype=torch.float16, device="cuda" + (1, 4, 1000, 4), + (16000, 4, 16, 1), + dtype=torch.float16, + device=device_type, ) ) args.append( @@ -1319,7 +1380,7 @@ def fn(arg7_1, add_1, permute_2, select_scatter, slice_8): (3, 1, 4, 1000, 4), (16, 48000, 4, 48, 1), dtype=torch.float16, - device="cuda", + device=device_type, ) ) args.append( @@ -1327,7 +1388,7 @@ def fn(arg7_1, add_1, permute_2, select_scatter, slice_8): (2, 1, 4, 1000, 4), (16, 48000, 4, 48, 1), dtype=torch.float16, - device="cuda", + device=device_type, ) ) args.append( @@ -1335,7 +1396,7 @@ def fn(arg7_1, add_1, permute_2, select_scatter, slice_8): (2, 1, 4, 1000, 4), (19200, 19200, 4800, 4, 1), dtype=torch.float16, - device="cuda", + device=device_type, ) ) @@ -1352,9 +1413,9 @@ def fn(x, y, z): x = torch.zeros_like(x) return x.index_put_([y], z, True) - x = torch.zeros((512, 512), device="cuda", dtype=torch.bool) - y = torch.zeros((512,), device="cuda", dtype=torch.int64) - z = torch.ones((512, 512), device="cuda", dtype=torch.bool) + x = torch.zeros((512, 512), device=device_type, dtype=torch.bool) + y = torch.zeros((512,), device=device_type, dtype=torch.int64) + z = torch.ones((512, 512), device=device_type, dtype=torch.bool) opt_fn = torch.compile(fn, backend="inductor") @@ -1375,9 +1436,9 @@ def fn(x, y, z): x = torch.zeros_like(x) return x.index_put([y], z, True) - x = torch.zeros((512, 512), device="cuda", dtype=torch.bool) - y = torch.zeros((512,), device="cuda", dtype=torch.int64) - z = torch.ones((512, 512), device="cuda", dtype=torch.bool) + x = torch.zeros((512, 512), device=device_type, dtype=torch.bool) + y = torch.zeros((512,), device=device_type, dtype=torch.int64) + z = torch.ones((512, 512), device=device_type, dtype=torch.bool) opt_fn = torch.compile(fn, backend="inductor") @@ -1420,7 +1481,7 @@ def forward(self, x): cnts = torch._dynamo.testing.CompileCounterWithBackend("inductor") - model = Model().cuda().half() + model = Model().to(device_type).half() model = torch.compile(model, backend=cnts, dynamic=True) with torch.backends.cuda.sdp_kernel( @@ -1429,9 +1490,9 @@ def forward(self, x): enable_mem_efficient=False, enable_cudnn=False, ): - input1 = torch.rand(5, 512, 1024, device="cuda", dtype=torch.float16) - input2 = torch.rand(5, 513, 1024, device="cuda", dtype=torch.float16) - input3 = torch.rand(5, 514, 1024, device="cuda", dtype=torch.float16) + input1 = torch.rand(5, 512, 1024, device=device_type, dtype=torch.float16) + input2 = torch.rand(5, 513, 1024, device=device_type, dtype=torch.float16) + input3 = torch.rand(5, 514, 1024, device=device_type, dtype=torch.float16) out1 = model(input1) out2 = model(input2) @@ -1445,9 +1506,9 @@ def fn(x, y, z): x = torch.zeros_like(x) return x.index_put([y], z, True) - x = torch.zeros((512, 512), device="cuda", dtype=torch.int32) - y = torch.zeros((512,), device="cuda", dtype=torch.int64) - z = torch.ones((512, 512), device="cuda", dtype=torch.int32) + x = torch.zeros((512, 512), device=device_type, dtype=torch.int32) + y = torch.zeros((512,), device=device_type, dtype=torch.int64) + z = torch.ones((512, 512), device=device_type, dtype=torch.int32) opt_fn = torch.compile(fn, backend="inductor") @@ -1483,8 +1544,8 @@ def test_emulate_precision_casts_norm_rounding(self): torch.manual_seed(0) torch.cuda.manual_seed_all(0) - x = torch.rand(1000, device="cuda", dtype=torch.bfloat16) - scalar = torch.rand([], device="cuda", dtype=torch.float32) + x = torch.rand(1000, device=device_type, dtype=torch.bfloat16) + scalar = torch.rand([], device=device_type, dtype=torch.float32) def fn(inp, scale): y = inp.norm() @@ -1500,7 +1561,7 @@ def fn(inp, scale): @torch._inductor.config.patch(emulate_precision_casts=True) def test_emulate_precision_casts_min_pow_chain(self): torch.manual_seed(0) - torch.cuda.manual_seed_all(0) + torch.cuda.manual_seed_all(0) if TEST_CUDA else torch.xpu.manual_seed_all(0) with dynamo_config.patch( capture_scalar_outputs=True, capture_dynamic_output_shape_ops=True @@ -1508,17 +1569,17 @@ def test_emulate_precision_casts_min_pow_chain(self): arg0 = torch.rand( [383, 55, 2, 3], dtype=torch.float16, - device="cuda", + device=device_type, requires_grad=True, ) arg1 = torch.rand( - [383, 55], dtype=torch.bfloat16, device="cuda", requires_grad=True + [383, 55], dtype=torch.bfloat16, device=device_type, requires_grad=True ) arg2 = torch.rand( - [383, 55], dtype=torch.float32, device="cuda", requires_grad=True + [383, 55], dtype=torch.float32, device=device_type, requires_grad=True ) arg3 = torch.rand( - [383, 55], dtype=torch.float32, device="cuda", requires_grad=True + [383, 55], dtype=torch.float32, device=device_type, requires_grad=True ) def fn(a0, a1, a2, a3): @@ -1557,22 +1618,37 @@ def test_emulate_precision_casts_mean_ratio_chain(self): capture_scalar_outputs=True, capture_dynamic_output_shape_ops=True ): arg0 = torch.rand( - [125070], dtype=torch.bfloat16, device="cuda", requires_grad=True + [125070], dtype=torch.bfloat16, device=device_type, requires_grad=True ) arg1 = torch.rand( - [1895, 3, 11], dtype=torch.float16, device="cuda", requires_grad=True + [1895, 3, 11], + dtype=torch.float16, + device=device_type, + requires_grad=True, ) arg2 = torch.rand( - [1895, 3, 11], dtype=torch.float32, device="cuda", requires_grad=True + [1895, 3, 11], + dtype=torch.float32, + device=device_type, + requires_grad=True, ) arg3 = torch.rand( - [1895, 3, 11], dtype=torch.float32, device="cuda", requires_grad=True + [1895, 3, 11], + dtype=torch.float32, + device=device_type, + requires_grad=True, ) arg4 = torch.rand( - [1895, 3, 11], dtype=torch.float32, device="cuda", requires_grad=True + [1895, 3, 11], + dtype=torch.float32, + device=device_type, + requires_grad=True, ) arg5 = torch.rand( - [5, 379, 165], dtype=torch.float32, device="cuda", requires_grad=True + [5, 379, 165], + dtype=torch.float32, + device=device_type, + requires_grad=True, ) def fn(a0, a1, a2, a3, a4, a5): @@ -1604,14 +1680,14 @@ def fn(a0, a1, a2, a3, a4, a5): def test_dont_inplace_disjoint_accesses(self): # TODO - would not need mms if we could annotate donated buffer.. def forward( # noqa: F821, F722 - arg0_1: "bf16[2048, 2048][2048, 1]cuda:0", # noqa: F821, F722 - arg1_1: "bf16[8, 4096, 2048][8388608, 2048, 1]cuda:0", # noqa: F821, F722 - arg2_1: "bf16[2048, 2048][2048, 1]cuda:0", # noqa: F821, F722 - arg3_1: "bf16[2048, 2048][2048, 1]cuda:0", # noqa: F821, F722 - arg4_1: "bf16[2048][1]cuda:0", # noqa: F821, F722 - arg5_1: "bf16[2048][1]cuda:0", # noqa: F821, F722 - arg6_1: "f32[4096, 128][128, 1]cuda:0", # noqa: F821, F722 - arg7_1: "f32[4096, 128][128, 1]cuda:0", # noqa: F821, F722 + arg0_1: f"bf16[2048, 2048][2048, 1]{device_type}:0", # noqa: F821, F722 + arg1_1: f"bf16[8, 4096, 2048][8388608, 2048, 1]{device_type}:0", # noqa: F821, F722 + arg2_1: f"bf16[2048, 2048][2048, 1]{device_type}:0", # noqa: F821, F722 + arg3_1: f"bf16[2048, 2048][2048, 1]{device_type}:0", # noqa: F821, F722 + arg4_1: f"bf16[2048][1]{device_type}:0", # noqa: F821, F722 + arg5_1: f"bf16[2048][1]{device_type}:0", # noqa: F821, F722 + arg6_1: f"f32[4096, 128][128, 1]{device_type}:0", # noqa: F821, F722 + arg7_1: f"f32[4096, 128][128, 1]{device_type}:0", # noqa: F821, F722 ): permute = torch.ops.aten.permute.default(arg0_1, [1, 0]) arg0_1 = None @@ -1773,7 +1849,7 @@ def forward( # noqa: F821, F722 from torch._dynamo.debug_utils import aot_graph_input_parser - kwargs = aot_graph_input_parser(forward) + kwargs = aot_graph_input_parser(forward, device=device_type) out, code = run_and_get_code(torch.compile(forward), **kwargs) # ignore tiny values.. prior to this fix absolute error was ~28 self.assertEqual(forward(**kwargs), out, atol=0.01, rtol=2) @@ -1781,8 +1857,8 @@ def forward( # noqa: F821, F722 # https://github.com/pytorch/pytorch/issues/104937 def test_linear_with_zero_infeature_size(self): - m = nn.Linear(in_features=0, out_features=0, bias=True).to("cuda") - x = torch.rand(1, 1, 0, device="cuda") + m = nn.Linear(in_features=0, out_features=0, bias=True).to(device_type) + x = torch.rand(1, 1, 0, device=device_type) expect = m(x) opt_fn = torch.compile(m) actual = opt_fn(x) @@ -1791,7 +1867,7 @@ def test_linear_with_zero_infeature_size(self): @config.patch(fallback_random=True) def test_multi_output_layout_fallback(self): mod = nn.RReLU(lower=3.2350976, upper=8.4220314, inplace=True) - inp = torch.rand([4, 4]).cuda() + inp = torch.rand([4, 4]).to(device_type) m = torch.compile(mod) with freeze_rng_state(): @@ -1806,8 +1882,8 @@ def test_sorted_masks(self): def foo(x, y): return (x + y).sum(dim=1) - x = torch.rand([255, 255], device="cuda") - y = torch.rand([255, 255], device="cuda") + x = torch.rand([255, 255], device=device_type) + y = torch.rand([255, 255], device=device_type) _, code = run_and_get_code(foo, x, y) FileCheck().check("tl.load").check_same("r0_mask").check_same("xmask").run( @@ -1821,7 +1897,8 @@ def cat(inps): for dtype in [torch.uint8, torch.int8]: inps = [ - torch.empty([256, 256], dtype=dtype, device="cuda") for _ in range(4) + torch.empty([256, 256], dtype=dtype, device=device_type) + for _ in range(4) ] out, code = run_and_get_code(cat, inps) @@ -1860,20 +1937,20 @@ def fn(arg207_1, arg208_1, convert_element_type_40, expand, full, mul_3): return (sub_2,) args = [ - torch.randn((8, 1024, 4, 4), device="cuda") > 0, # torch.bool tensor - torch.randn((1, 1024, 1, 1), device="cuda"), - torch.randn((8, 1024, 4, 4), device="cuda"), - torch.randn((8, 1024, 1, 1), dtype=torch.float16, device="cuda").expand( - (8, 1024, 4, 4) - ), - torch.randn((), device="cuda"), - torch.randn((1024,), device="cuda"), + torch.randn((8, 1024, 4, 4), device=device_type) > 0, # torch.bool tensor + torch.randn((1, 1024, 1, 1), device=device_type), + torch.randn((8, 1024, 4, 4), device=device_type), + torch.randn( + (8, 1024, 1, 1), dtype=torch.float16, device=device_type + ).expand((8, 1024, 4, 4)), + torch.randn((), device=device_type), + torch.randn((1024,), device=device_type), ] fn(*args) - torch.cuda.synchronize() # shake out Triton Error [CUDA]: misaligned address + torch.accelerator.synchronize() # shake out Triton Error [CUDA]: misaligned address def test_mutated_aligned_tensor(self): - t = torch.rand(4096, device="cuda", dtype=torch.float16) + t = torch.rand(4096, device=device_type, dtype=torch.float16) def foo(x): return x.add_(1) @@ -1893,8 +1970,8 @@ def foo(x): def test_non_commutative_scan_op(self): from torch._higher_order_ops.associative_scan import associative_scan - a = torch.randn(1024, 8192, dtype=torch.float64, device="cuda") - b = torch.randn(1024, 8192, dtype=torch.float64, device="cuda") + a = torch.randn(1024, 8192, dtype=torch.float64, device=device_type) + b = torch.randn(1024, 8192, dtype=torch.float64, device=device_type) def baseline(v, u): A = [] @@ -1922,7 +1999,7 @@ def inner_reduce(x): assert x.shape[1] <= 1024 # noqa: S101 return x.sum(1) - a = torch.randn(50, 600, device="cuda") + a = torch.randn(50, 600, device=device_type) out, code = run_and_get_code(inner_reduce, a) self.assertEqual(inner_reduce(a), out) self.assertTrue("for roffset" not in code) @@ -1979,7 +2056,7 @@ def forward(self, hidden_states: Tensor, attention_mask: Tensor) -> Tensor: ) return attn_output - device = torch.device("cuda") + device = torch.device(device_type) num_attention_heads = 8 hidden_size = 512 attention_probs_dropout_prob = 0.0 @@ -2005,30 +2082,34 @@ def forward(self, hidden_states: Tensor, attention_mask: Tensor) -> Tensor: def test_non_contiguous_unaligned_input_indices(self): from torch._inductor.compile_fx import remove_unaligned_input_idxs - inputs = [torch.ones(2, 2, device="cuda"), torch.ones(2, 2, device="cuda")[1:]] + inputs = [ + torch.ones(2, 2, device=device_type), + torch.ones(2, 2, device=device_type)[1:], + ] idxs = remove_unaligned_input_idxs(inputs, [1]) self.assertEqual(idxs, []) inputs = [ - torch.ones(2, 2, device="cuda"), - torch.ones(2, 2, device="cuda"), - torch.ones(2, 2, device="cuda")[1:], + torch.ones(2, 2, device=device_type), + torch.ones(2, 2, device=device_type), + torch.ones(2, 2, device=device_type)[1:], ] idxs = remove_unaligned_input_idxs(inputs, [0, 2]) self.assertEqual(idxs, [0]) + @skipIfXpu(msg="cudagraph is not supported on xpu") @config.patch("triton.cudagraphs", True) def test_unused_cpu_input_cudagraphs(self): def fn(x, y): return x.sin().sin().sin().sin().cos() + 1 fx_graph = torch.fx.symbolic_trace(fn) - inp = [torch.randn(64, device="cuda"), torch.randn(64, device="cpu")] + inp = [torch.randn(64, device=device_type), torch.randn(64, device="cpu")] compiled_fn, (graph,) = run_and_get_graph_lowering( torch._inductor.compile, fx_graph, inp ) self.assertEqual(graph.disable_cudagraphs_reason, None) - self.assertEqual(graph.device_types, {"cuda"}) + self.assertEqual(graph.device_types, {device_type}) self.assertEqual(compiled_fn(*inp), fn(*inp)) def test_epilogue_fusion_with_view(self): @@ -2044,8 +2125,8 @@ def forward(self, x): x = x.view(x.size(0), -1) return self.relu(self.linear(x)) - m = ToyModel().to(device="cuda:0") - input_tensor = torch.randn(32, 3, 64, 64).to(device="cuda:0") + m = ToyModel().to(device=f"{device_type}:0") + input_tensor = torch.randn(32, 3, 64, 64).to(device=f"{device_type}:0") from torch._inductor.utils import fresh_cache with fresh_cache(): @@ -2054,6 +2135,7 @@ def forward(self, x): out2 = m(input_tensor) self.assertEqual(out, out2, atol=1e-3, rtol=1e-3) + @skipIfXpu(msg="cudagraph is not supported on xpu") @config.patch("triton.cudagraphs", True) def test_cpu_index(self): @torch.compile(fullgraph=True) @@ -2061,25 +2143,25 @@ def fn(x): return x[torch.arange(32)] result, (graph,) = run_and_get_graph_lowering( - fn, torch.randn(64, device="cuda") + fn, torch.randn(64, device=device_type) ) self.assertEqual(graph.disable_cudagraphs_reason, None) - self.assertEqual(graph.device_types, {"cuda"}) + self.assertEqual(graph.device_types, {device_type}) - inp = torch.randn(64, device="cuda", requires_grad=True) + inp = torch.randn(64, device=device_type, requires_grad=True) result, (graph,) = run_and_get_graph_lowering(fn, inp) self.assertEqual(graph.disable_cudagraphs_reason, None) - self.assertEqual(graph.device_types, {"cuda"}) + self.assertEqual(graph.device_types, {device_type}) result, (graph,) = run_and_get_graph_lowering(lambda: result.sum().backward()) self.assertEqual(graph.disable_cudagraphs_reason, None) - self.assertEqual(graph.device_types, {"cuda"}) + self.assertEqual(graph.device_types, {device_type}) @unittest.skipIf(IS_FBCODE, "Not runnable in fbcode") def test_triton_interpret(self): import subprocess - script = """ + script = f""" import os os.environ["TRITON_INTERPRET"] = "1" import torch @@ -2089,8 +2171,9 @@ def foo(x): return x + 1 # somehow gives different results.. still, check that it doesn't error -foo(torch.rand([256], device="cuda")) +foo(torch.rand([256], device=\"{device_type}\")) """ + subprocess.run([sys.executable, "-c", script], check=True) def test_reflection_pad_loop_order(self): @@ -2100,8 +2183,8 @@ def fn(x, y): return a + b cfn = torch.compile(fn) - a = torch.rand((10, 10, 10), device="cuda") - b = torch.rand((10, 10, 10), device="cuda") + a = torch.rand((10, 10, 10), device=device_type) + b = torch.rand((10, 10, 10), device=device_type) expect = fn(a, b) actual, code = run_and_get_code(cfn, a, b) self.assertEqual(expect, actual) @@ -2162,7 +2245,7 @@ def foo(inp): return cat_1 for mark_dynamic in [False, True]: - inp = torch.rand((65536, 8192), dtype=torch.bfloat16, device="cuda") + inp = torch.rand((65536, 8192), dtype=torch.bfloat16, device=device_type) if mark_dynamic: torch._dynamo.mark_dynamic(inp, 0) foo_c = torch.compile(foo) @@ -2172,7 +2255,7 @@ def foo(inp): not SM90OrLater, "uses bfloat16 atomic add instrs which requires SM >= 90" ) def test_float8_e8m0fnu(self): - device = "cuda" + device = device_type dtype = torch.float8_e8m0fnu hp_dtype = torch.float32 # and torch.bfloat16 @@ -2219,9 +2302,9 @@ def f(x, y): return torch.index_select(x, 0, y) x = torch.randn( - 2000, 384, dtype=torch.bfloat16, device="cuda", requires_grad=True + 2000, 384, dtype=torch.bfloat16, device=device_type, requires_grad=True ) - y = torch.ones(713268, dtype=torch.int64, device="cuda") + y = torch.ones(713268, dtype=torch.int64, device=device_type) x_ref = x.clone().detach().requires_grad_(True) y_ref = y.clone().detach() @@ -2254,15 +2337,15 @@ def f(arg0_1, arg1_1): start=0, step=1, dtype=torch.int64, - device="cuda", + device=device_type, requires_grad=False, ) view_3 = torch.ops.aten.view.default(iota, [1, 36]) max_1 = torch.ops.aten.max.default(view_3) return (max_1,) - x = torch.ones(1, 64, device="cuda", dtype=torch.int64) - y = torch.randn(64, 3072, device="cuda", dtype=torch.bfloat16) + x = torch.ones(1, 64, device=device_type, dtype=torch.int64) + y = torch.randn(64, 3072, device=device_type, dtype=torch.bfloat16) out = f(x, y) self.assertEqual(torch.compile(f)(x, y), out) @@ -2278,9 +2361,9 @@ def f(x, y): return torch.index_select(x, 0, y) x = torch.randn( - 2000, 384, dtype=torch.bfloat16, device="cuda", requires_grad=True + 2000, 384, dtype=torch.bfloat16, device=device_type, requires_grad=True ) - y = torch.ones(713268, dtype=torch.int64, device="cuda") + y = torch.ones(713268, dtype=torch.int64, device=device_type) x_ref = x.clone().detach().requires_grad_(True) y_ref = y.clone().detach() @@ -2322,7 +2405,7 @@ def test_3d_tiling(self): 1, 2, ) - GPU_TYPE = "cuda" + GPU_TYPE = device_type def get_input() -> torch.Tensor: device = torch.device(GPU_TYPE) @@ -2344,7 +2427,7 @@ def test_repeated_masked_load(self): mem_eff_temporal_upsampling_interp_chunks = 2 from functorch.einops import rearrange - x = torch.randn(1, 8, 12, 12, 4, dtype=torch.float16, device="cuda") + x = torch.randn(1, 8, 12, 12, 4, dtype=torch.float16, device=device_type) x = x.permute(0, 1, 4, 2, 3) # make non-contiguous x = rearrange(x, "b c t h w -> b c t (h w)") @@ -2406,8 +2489,8 @@ def forward(self, x): return x - model = ToyModel().to("cuda") - input_tensor = torch.randn((2, 4)).to("cuda") + model = ToyModel().to(device_type) + input_tensor = torch.randn((2, 4)).to(device_type) compile_default = torch.compile(model, mode="default") compile_max_autotune = torch.compile(model, mode="max-autotune") @@ -2434,7 +2517,7 @@ def forward(self, x): x = self.adaptive_pool(x) return x - model = Model().cuda() + model = Model().to(device_type) model.eval() test_cases = [ (1, 3, 8, 8), @@ -2446,7 +2529,7 @@ def forward(self, x): for batch, channels, h, w in test_cases: with self.subTest(input_shape=(batch, channels, h, w)): - input_tensor = torch.randn(batch, channels, h, w, device="cuda") + input_tensor = torch.randn(batch, channels, h, w, device=device_type) # Test eager mode with torch.no_grad(): @@ -2499,7 +2582,7 @@ def forward(self, x: torch.Tensor) -> torch.Tensor: numel = 1 for dim in quantiles_shape: numel *= dim - data = torch.randn(numel, dtype=torch.float32, device="cuda") + data = torch.randn(numel, dtype=torch.float32, device=device_type) # Create tensor with specified shape and strides quantiles = torch.as_strided( @@ -2509,7 +2592,7 @@ def forward(self, x: torch.Tensor) -> torch.Tensor: quantiles = torch.sort(quantiles, dim=0)[0] x_shape = (batch_size,) + quantiles_shape[1:] - x = torch.randn(*x_shape, dtype=torch.float32, device="cuda") + x = torch.randn(*x_shape, dtype=torch.float32, device=device_type) foo = Foo(quantiles) foo_compiled = torch.compile(Foo(quantiles), fullgraph=True) @@ -2522,7 +2605,7 @@ def forward(self, x: torch.Tensor) -> torch.Tensor: self.assertEqual(eager, compiled) def test_identity_load(self): - device = "cuda" + device = device_type def f(x, y): y2 = torch.cat( @@ -2557,9 +2640,12 @@ def f(x, y): self.assertEqual(eager_out, compile_out) + @skipIfXpu( + msg="Explicit attn_mask should not be set when is_causal=True - xpu-ops: 2802" + ) def test_qwen2_7b_sdpa_input_alignment_requires_recompile(self): # SDPA constraints ensures inputs have alignment (8). - device = "cuda" + device = device_type def forward(q_proj, k_proj, attn_mask): scale = 0.08838834764831845 # 1/sqrt(128) @@ -2677,8 +2763,8 @@ def compiled_divide(x, y): torch.float32, torch.float64, ]: - y_ten = torch.tensor([y], dtype=y_dtype, device="cuda") - x_ten = torch.tensor([x], dtype=x_dtype, device="cuda") + y_ten = torch.tensor([y], dtype=y_dtype, device=device_type) + x_ten = torch.tensor([x], dtype=x_dtype, device=device_type) torch._dynamo.reset() compiled_div = Decimal(compiled_divide(x_ten, y_ten).item()) @@ -2686,6 +2772,7 @@ def compiled_divide(x, y): self.assertEqual(eager_div, compiled_div) + @skipIfXpu(msg="triton dependency - xpu-ops: 2554") @config.patch({"eager_numerics.division_rounding": False}) @xfailIfROCm def test_truediv_base_not_bitwise_equivalent(self): @@ -2693,8 +2780,8 @@ def test_truediv_base_not_bitwise_equivalent(self): y, x = 7.0, 11.0 - y_ten = torch.tensor([y], dtype=torch.float32, device="cuda") - x_ten = torch.tensor([x], dtype=torch.float32, device="cuda") + y_ten = torch.tensor([y], dtype=torch.float32, device=device_type) + x_ten = torch.tensor([x], dtype=torch.float32, device=device_type) compile_out, code = run_and_get_code( torch.compile(lambda x, y: x / y), @@ -2722,13 +2809,14 @@ def fn(x): self.assertTrue(compile_decimal > Decimal(0)) + @skipIfXpu(msg="Decimal object comparison failed - xpu-ops: 2810") @skipIfRocm(msg="ROCm preserves subnormals by default") @config.patch({"eager_numerics.disable_ftz": False}) def test_not_disabling_ftz_yields_zero(self): from decimal import Decimal x = -128.0 - x_ten = torch.tensor([x], dtype=torch.float32, device="cuda") + x_ten = torch.tensor([x], dtype=torch.float32, device=device_type) def fn(x): return 2.0**x @@ -2766,7 +2854,7 @@ def fn(x): if __name__ == "__main__": from torch._inductor.test_case import run_tests - from torch.testing._internal.inductor_utils import HAS_CUDA_AND_TRITON + from torch.testing._internal.inductor_utils import HAS_GPU_AND_TRITON - if HAS_CUDA_AND_TRITON and not TEST_WITH_ASAN: + if HAS_GPU_AND_TRITON and not TEST_WITH_ASAN: run_tests(needs="filelock") diff --git a/test/inductor/test_torchinductor.py b/test/inductor/test_torchinductor.py index 4f6ee7c71ffc3..d9a9c17c7f490 100644 --- a/test/inductor/test_torchinductor.py +++ b/test/inductor/test_torchinductor.py @@ -10931,7 +10931,6 @@ def fn(x): ], ) - @skipIfXpu(msg="Incorrect XPU reference") def test_argmax_argmin2(self): def fn(x): return ( @@ -10943,7 +10942,6 @@ def fn(x): self.common(fn, (torch.randn([144, 144]),)) - @skipIfXpu(msg="Incorrect XPU reference") def test_argmax_argmin_with_duplicates(self): def fn(x): return ( @@ -10965,7 +10963,6 @@ def fn(x): t1 = torch.randint(8, size=(1028, 1028)) self.common(fn, (t1,)) - @skipIfXpu(msg="# Incorrect XPU reference ") @xfail_if_mps # eager nan is wrong, see https://github.com/pytorch/pytorch/issues/130295 @skip_if_halide # nan behavior def test_argmax_argmin_with_nan(self): @@ -11090,7 +11087,6 @@ def shrink_rank(x, rank): [rank4_inps, rank3_inps, rank5_inps], ) - @skipIfXpu(msg="Incorrect XPU reference") def test_argmax_argmin3(self): def fn(x): return ( @@ -12829,7 +12825,6 @@ def fn(q, k, v): ) @xfail_if_mps_unimplemented - @expectedFailureXPU @unittest.skipIf( not PLATFORM_SUPPORTS_MEM_EFF_ATTENTION, "Some archs don't support mem eff SDPA" ) @@ -15658,7 +15653,7 @@ def f(x): self.assertFalse("ReductionHint.INNER" in code) @skip_if_halide - @requires_cuda_and_triton + @requires_gpu_and_triton def test_triton_argmin_argmax_transpose_logical_index(self): def fn(x): x.tan_() @@ -15695,7 +15690,7 @@ def fn(x): self.common(fn, (torch.randn(6, 4, device=GPU_TYPE).t().contiguous().t(),)) @skip_if_halide - @requires_cuda_and_triton + @requires_gpu_and_triton def test_unbacked_float_item(self): def fn(x, max_val): return torch.clamp(x, 0, max_val.item()) @@ -15735,7 +15730,7 @@ def fn(a, b, c, d): if self.device.lower() == "cuda": self.assertEqual(torch._inductor.metrics.generated_kernel_count, 2) - @requires_cuda_and_triton + @requires_gpu_and_triton @config.patch(combo_kernels=True) @torch._dynamo.config.patch(assume_static_by_default=False) def test_combo_kernel_store_mask(self): diff --git a/test/test_fake_tensor.py b/test/test_fake_tensor.py index effdacc58852b..bb67c0fbe07ec 100644 --- a/test/test_fake_tensor.py +++ b/test/test_fake_tensor.py @@ -64,6 +64,7 @@ skipIfCrossRef, skipIfTorchDynamo, skipIfWindows, + skipIfXpu, TemporaryFileName, TEST_WITH_TORCHDYNAMO, TestCase, @@ -82,6 +83,8 @@ torch._dynamo.config.fake_tensor_cache_enabled = True torch._dynamo.config.fake_tensor_cache_crosscheck_enabled = True +device_type = acc.type if (acc := torch.accelerator.current_accelerator()) else "cpu" + def expectedFailurePropagateRealTensors(fn): fn._expected_failure_propagate_real_tensors = True @@ -1655,6 +1658,7 @@ def test_cross_entropy_loss(self): self.assertEqual(ref.size(), meta_out.size()) + @skipIfXpu(msg="MetadataMismatchError, torch-xpu-ops: 2802") @unittest.skipIf( not PLATFORM_SUPPORTS_FLASH_ATTENTION, "Does not support SDPA or pre-SM80 hardware", @@ -1671,14 +1675,14 @@ def forward(self, arg1, arg2, arg3): args_new = [ [ - ((1, 48, 64, 64), (0, 4096, 64, 1), torch.float16, "cuda"), - ((1, 48, 64, 64), (0, 4096, 64, 1), torch.float16, "cuda"), - ((1, 48, 64, 64), (0, 4096, 64, 1), torch.float16, "cuda"), + ((1, 48, 64, 64), (0, 4096, 64, 1), torch.float16, device_type), + ((1, 48, 64, 64), (0, 4096, 64, 1), torch.float16, device_type), + ((1, 48, 64, 64), (0, 4096, 64, 1), torch.float16, device_type), ], [ - ((4, 2, 16, 32), (1024, 512, 32, 1), torch.float16, "cuda"), - ((4, 2, 16, 32), (1024, 512, 32, 1), torch.float16, "cuda"), - ((4, 2, 16, 32), (1024, 512, 32, 1), torch.float16, "cuda"), + ((4, 2, 16, 32), (1024, 512, 32, 1), torch.float16, device_type), + ((4, 2, 16, 32), (1024, 512, 32, 1), torch.float16, device_type), + ((4, 2, 16, 32), (1024, 512, 32, 1), torch.float16, device_type), ], ] for args_list in args_new: diff --git a/test/test_scaled_matmul_cuda.py b/test/test_scaled_matmul_cuda.py index 47048f8a4091a..4fdc16fb4f0b6 100644 --- a/test/test_scaled_matmul_cuda.py +++ b/test/test_scaled_matmul_cuda.py @@ -51,6 +51,7 @@ skipIfRocm, TEST_CUDA, TestCase, + skipIfXpu, ) from torch.testing._internal.common_quantized import ( _bfloat16_to_float4_e2m1fn_x2, @@ -922,14 +923,13 @@ def _2d_to_blocked_scaled(X, K, G, offs, format): # Assert outputs are close. torch.testing.assert_close(y_lp, y_bf16, atol=8.0e-2, rtol=8.0e-2) - @unittest.skipIf(not PLATFORM_SUPPORTS_FP8, f8_msg) @parametrize("base_dtype", [torch.float16, torch.bfloat16, torch.float32]) @parametrize("x_cm", [True, False]) @parametrize("y_cm", [True, False]) def test_scaled_mm_vs_emulated(self, base_dtype, x_cm, y_cm, device="cuda"): # Blackwell (SM_10) supports all possible layout permutations, while Hopper only TN - if (x_cm, y_cm) != (True, False) and torch.cuda.get_device_properties(0).major != 10: + if (x_cm, y_cm) != (True, False) and torch.cuda.is_available() and torch.cuda.get_device_properties(0).major != 10: raise unittest.SkipTest("Unsupported layout on the architecture") torch.manual_seed(42) input_dtype = e4m3_type @@ -1105,6 +1105,7 @@ def test_error_message_fp8_pre_sm89(self, device) -> None: lambda: scaled_mm_wrap(x, y, scale_a, scale_b, out_dtype=torch.float32), ) + @skipIfXpu(msg="AssertionError, torch-xpu-ops: 2862") @unittest.skipIf(not PLATFORM_SUPPORTS_FP8, f8_msg) @unittest.skipIf(SM100OrLater, "fast_accum is SM90-only") def test_float8_scale_fast_accum(self, device) -> None: diff --git a/torch/testing/_internal/common_cuda.py b/torch/testing/_internal/common_cuda.py index 0dc3d0b4f10cb..ac8737b5ec301 100644 --- a/torch/testing/_internal/common_cuda.py +++ b/torch/testing/_internal/common_cuda.py @@ -5,7 +5,7 @@ import functools import torch import torch.cuda -from torch.testing._internal.common_utils import LazyVal, TEST_NUMBA, TEST_WITH_ROCM, TEST_CUDA, IS_WINDOWS, IS_MACOS +from torch.testing._internal.common_utils import LazyVal, TEST_NUMBA, TEST_WITH_ROCM, TEST_CUDA, IS_WINDOWS, IS_MACOS, TEST_XPU import inspect import contextlib import os @@ -76,6 +76,8 @@ def evaluate_platform_supports_flash_attention(): return evaluate_gfx_arch_within(arch_list) if TEST_CUDA: return not IS_WINDOWS and SM80OrLater + if TEST_XPU: + return True return False def evaluate_platform_supports_efficient_attention(): @@ -86,6 +88,8 @@ def evaluate_platform_supports_efficient_attention(): return evaluate_gfx_arch_within(arch_list) if TEST_CUDA: return True + if TEST_XPU: + return True return False def evaluate_platform_supports_cudnn_attention(): @@ -117,6 +121,8 @@ def evaluate_platform_supports_bf16(): return SM80OrLater elif torch.version.hip: return True + elif TEST_XPU: + return True return False @@ -153,6 +159,8 @@ def evaluate_platform_supports_fp8(): return True else: return SM90OrLater or torch.cuda.get_device_capability() == (8, 9) + if torch.xpu.is_available(): + return True return False def evaluate_platform_supports_fp8_grouped_gemm(): From 188411752cd8dd6a4d71ea382698bcf8c468a363 Mon Sep 17 00:00:00 2001 From: "Deng, Daisy" Date: Mon, 23 Mar 2026 02:49:07 +0000 Subject: [PATCH 0008/1172] [XPU] Enable skipped inductor test on Intel GPU - generalize 12 inductor test files and enable xpu (#174057) Refactor 12 inductor test files 1. generalize code with requires_gpu_and_triton and torch.accelerator, HAS_GPU_AND_TRITON, GPU_TYPE 2. Remove xpu skips according to latest status 3. Updated a dynamo test file because it is referenced by inductor test, test_higher_order_ops.FuncTorchHigherOrderOpTests is called by test/inductor/test_compiled_autograd.py. - cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @kadeng @muchulee8 @amjames @chauhang @aakhundov @coconutruben @jataylo @Lucaskabela Pull Request resolved: https://github.com/pytorch/pytorch/pull/174057 Approved by: https://github.com/etaf, https://github.com/chuanqi129, https://github.com/mlazos ghstack dependencies: #174053, #174054, #174055, #174056 --- test/dynamo/test_higher_order_ops.py | 13 +++++---- test/inductor/test_aot_inductor_custom_ops.py | 2 +- test/inductor/test_combo_kernels.py | 29 +++++++++++++++++-- test/inductor/test_compiled_autograd.py | 11 ++++--- test/inductor/test_compiled_optimizers.py | 6 ++-- test/inductor/test_control_deps.py | 4 +-- test/inductor/test_control_flow.py | 2 -- test/inductor/test_custom_op_autotune.py | 13 +++++---- test/inductor/test_deterministic.py | 2 -- test/inductor/test_fp8.py | 13 +++++---- test/inductor/test_fused_attention.py | 5 ---- test/inductor/test_gpu_cpp_wrapper.py | 14 --------- 12 files changed, 59 insertions(+), 55 deletions(-) diff --git a/test/dynamo/test_higher_order_ops.py b/test/dynamo/test_higher_order_ops.py index 84c540aefe90d..0c6089d1eee80 100644 --- a/test/dynamo/test_higher_order_ops.py +++ b/test/dynamo/test_higher_order_ops.py @@ -38,6 +38,7 @@ xfailIfTorchDynamo, ) from torch.testing._internal.hop_db import hop_db +from torch.testing._internal.inductor_utils import GPU_TYPE from torch.testing._internal.logging_utils import LoggingTestCase, make_logging_test from torch.testing._internal.triton_utils import ( requires_cuda_and_triton, @@ -3384,7 +3385,7 @@ def outer_body_fn(x): with self.assertRaisesRegex(RuntimeError, msg): fn_with_hints(x, y) - @requires_cuda_and_triton + @requires_gpu_and_triton def test_wrap_inductor_compiled_regions_option(self): """ Test that wrap_inductor_compiled_regions option wraps compiled regions @@ -3406,8 +3407,8 @@ def fn_wrapped(x, y): def fn_not_wrapped(x, y): return torch.matmul(x, y) - x = torch.randn(4, 4, device="cuda") - y = torch.randn(4, 4, device="cuda") + x = torch.randn(4, 4, device=GPU_TYPE) + y = torch.randn(4, 4, device=GPU_TYPE) # Test wrapped version - HOP should be visible in DebugMode with DebugMode() as debug_mode_wrapped: @@ -3428,7 +3429,7 @@ def fn_not_wrapped(x, y): self.assertEqual(result_wrapped, expected) self.assertEqual(result_not_wrapped, expected) - @requires_cuda_and_triton + @requires_gpu_and_triton def test_wrap_inductor_compiled_regions_with_backward(self): """ Test that wrap_inductor_compiled_regions works correctly with autograd. @@ -3443,8 +3444,8 @@ def test_wrap_inductor_compiled_regions_with_backward(self): def fn(x, y): return torch.matmul(x, y) - x = torch.randn(4, 4, device="cuda", requires_grad=True) - y = torch.randn(4, 4, device="cuda", requires_grad=True) + x = torch.randn(4, 4, device=GPU_TYPE, requires_grad=True) + y = torch.randn(4, 4, device=GPU_TYPE, requires_grad=True) # Clone for eager comparison x_eager = x.detach().clone().requires_grad_(True) diff --git a/test/inductor/test_aot_inductor_custom_ops.py b/test/inductor/test_aot_inductor_custom_ops.py index 1b608dc761702..e83d11fe0af38 100644 --- a/test/inductor/test_aot_inductor_custom_ops.py +++ b/test/inductor/test_aot_inductor_custom_ops.py @@ -505,7 +505,7 @@ def forward(self, x): args = (torch.randn(4, 4, device=self.device),) self.check_model(m, args) - @skipIfXpu + @skipIfXpu(msg="compile error - torch-xpu-ops: 2609") @unittest.skipIf(IS_FBCODE, "unable to find library -laoti_custom_ops") def test_custom_op_square(self) -> None: class Model(torch.nn.Module): diff --git a/test/inductor/test_combo_kernels.py b/test/inductor/test_combo_kernels.py index f18ff7c8a8269..36f2af14ca04f 100644 --- a/test/inductor/test_combo_kernels.py +++ b/test/inductor/test_combo_kernels.py @@ -19,7 +19,11 @@ TestCase, ) from torch.testing._internal.inductor_utils import GPU_TYPE, HAS_CPU, HAS_GPU_AND_TRITON -from torch.testing._internal.triton_utils import requires_gpu_and_triton +from torch.testing._internal.triton_utils import ( + requires_cuda_and_triton, + requires_gpu_and_triton, + requires_xpu_and_triton, +) aten = torch.ops.aten @@ -995,7 +999,7 @@ def fn(a0, a1, a2, b0, b1, b2): self.assertEqual(out_eager, out_compiled) self.assertTrue(5 <= torch._inductor.metrics.generated_kernel_count <= 6) - @requires_gpu_and_triton + @requires_cuda_and_triton @torch._dynamo.config.patch("automatic_dynamic_shapes", True) @torch._dynamo.config.patch("assume_static_by_default", True) @torch._inductor.config.patch("triton.autotune_at_compile_time", True) @@ -1016,6 +1020,27 @@ def fn(x, y, z): self.assertEqual(out_eager, out_compiled) + @requires_xpu_and_triton + @torch._dynamo.config.patch("automatic_dynamic_shapes", True) + @torch._dynamo.config.patch("assume_static_by_default", True) + @torch._inductor.config.patch("triton.autotune_at_compile_time", True) + def test_dynamic_shapes_persistent_reduction_mixed_x_dim_xpu(self): + def fn(x, y, z): + return x.sum(1), y.mean(1), z.max(1) + + inps = ( + torch.rand(16, 128, device=GPU_TYPE), + torch.rand(32, 128, device=GPU_TYPE), + torch.rand(32, 256, device=GPU_TYPE), + ) + torch._dynamo.mark_dynamic(inps[0], 0, min=1, max=256) + torch._dynamo.mark_dynamic(inps[1], 0, min=1, max=256) + torch._dynamo.mark_dynamic(inps[2], 0, min=1, max=256) + out_eager = fn(*inps) + out_compiled = torch.compile(fn)(*inps) + + self.assertEqual(out_eager, out_compiled) + @requires_gpu_and_triton def test_helper_fn_defined(self): def fn(x, y, z): diff --git a/test/inductor/test_compiled_autograd.py b/test/inductor/test_compiled_autograd.py index 0e6cbbf3b2d98..921f4937cb0e7 100644 --- a/test/inductor/test_compiled_autograd.py +++ b/test/inductor/test_compiled_autograd.py @@ -53,7 +53,7 @@ HAS_CPU, HAS_CUDA_AND_TRITON, HAS_GPU, - HAS_XPU_AND_TRITON, + HAS_GPU_AND_TRITON, ) from torch.testing._internal.logging_utils import logs_to_string from torch.testing._internal.triton_utils import ( @@ -5419,13 +5419,12 @@ def tearDown(self): test_autograd = load_test_module("test_autograd") test_custom_ops = load_test_module("test_custom_ops") test_higher_order_ops = load_test_module("dynamo/test_higher_order_ops") -if not HAS_XPU_AND_TRITON: - TestAutogradWithCompiledAutograd = wrap_test_class(test_autograd.TestAutograd) + +TestAutogradWithCompiledAutograd = wrap_test_class(test_autograd.TestAutograd) TestNestedCheckpointWithCompiledAutograd = wrap_test_class( test_autograd.TestNestedCheckpoint ) -if not HAS_XPU_AND_TRITON: - TestCustomOpWithCompiledAutograd = wrap_test_class(test_custom_ops.TestCustomOp) +TestCustomOpWithCompiledAutograd = wrap_test_class(test_custom_ops.TestCustomOp) HigherOrderOpTestsWithCompiledAutograd = wrap_test_class( test_higher_order_ops.HigherOrderOpTests ) @@ -5436,7 +5435,7 @@ def tearDown(self): test_higher_order_ops.ActivationCheckpointingTests ) -if torch.distributed.is_available() and HAS_CUDA_AND_TRITON: +if torch.distributed.is_available() and HAS_GPU_AND_TRITON: test_dtensor = load_test_module("distributed/tensor/test_dtensor_compile") TestDTensorCompileWithCompiledAutograd = wrap_test_class( test_dtensor.TestDTensorCompile diff --git a/test/inductor/test_compiled_optimizers.py b/test/inductor/test_compiled_optimizers.py index 225aae2ee23f5..802bea79bcc66 100644 --- a/test/inductor/test_compiled_optimizers.py +++ b/test/inductor/test_compiled_optimizers.py @@ -591,7 +591,7 @@ class CompiledOptimizerParityTests(TestCase): @optims(optim_db, dtypes=[torch.float32]) @parametrize("use_closure", [True, False]) def test_correctness(self, device, dtype, optim_info, use_closure): - torch.cuda.manual_seed_all(0) + torch.get_device_module(device).manual_seed_all(0) torch.manual_seed(0) random.seed(0) optim_cls = optim_info.optim_cls @@ -931,7 +931,7 @@ def fn(xs, ys): self.assertLess(end - start, 90) - @requires_cuda_and_triton + @requires_gpu_and_triton def test_S429861(self): # Just verify we can compile this function without error try: @@ -947,7 +947,7 @@ def test_S429861(self): from torch._inductor.utils import fresh_cache with fresh_cache(): - kwargs = aot_graph_input_parser(forward) + kwargs = aot_graph_input_parser(forward, device=GPU_TYPE) torch.compile(forward)(**kwargs) @requires_gpu_and_triton diff --git a/test/inductor/test_control_deps.py b/test/inductor/test_control_deps.py index 36b365e4f7d82..9adc3458fce75 100644 --- a/test/inductor/test_control_deps.py +++ b/test/inductor/test_control_deps.py @@ -8,7 +8,7 @@ from torch.testing._internal.common_utils import IS_LINUX from torch.testing._internal.inductor_utils import ( GPU_TYPE, - HAS_CUDA_AND_TRITON, + HAS_GPU_AND_TRITON, requires_gpu, ) @@ -259,5 +259,5 @@ def add_control_deps(graph): if __name__ == "__main__": - if IS_LINUX and HAS_CUDA_AND_TRITON: + if IS_LINUX and HAS_GPU_AND_TRITON: run_tests(needs="filelock") diff --git a/test/inductor/test_control_flow.py b/test/inductor/test_control_flow.py index 45761543b3186..680b503e3f815 100644 --- a/test/inductor/test_control_flow.py +++ b/test/inductor/test_control_flow.py @@ -14,7 +14,6 @@ decorateIf, instantiate_parametrized_tests, parametrize, - skipIfXpu, ) from torch.testing._internal.inductor_utils import GPU_TYPE, HAS_CPU, HAS_GPU from torch.testing._internal.triton_utils import requires_gpu @@ -385,7 +384,6 @@ def test_cond_unbacked_symint_closure(self, device, dynamic): dynamic=dynamic, ) - @skipIfXpu(msg="Remove this skip after issue #154949 resolved.") @requires_gpu def test_cond_control_flow_with_precomputed_size(self): class TestModel(torch.nn.Module): diff --git a/test/inductor/test_custom_op_autotune.py b/test/inductor/test_custom_op_autotune.py index b05515406dcbe..d5b9572fa54f3 100644 --- a/test/inductor/test_custom_op_autotune.py +++ b/test/inductor/test_custom_op_autotune.py @@ -26,6 +26,7 @@ skipIfXpu, ) from torch.testing._internal.inductor_utils import ( + GPU_TYPE, HAS_CPU, HAS_GPU, HAS_TRITON, @@ -45,8 +46,12 @@ def setUp(self) -> None: """Set up test environment with appropriate device and dtype.""" super().setUp() torch._dynamo.reset() - self.device = "cuda" if HAS_GPU else "cpu" - self.dtype = torch.float16 if self.device == "cuda" else torch.float32 + self.device = GPU_TYPE if HAS_GPU else "cpu" + self.dtype = ( + torch.float16 + if self.device == "cuda" or self.device == "xpu" + else torch.float32 + ) # Clear any previous lowering registrations to ensure test isolation from torch._inductor.lowering import user_lowerings @@ -164,7 +169,6 @@ def _create_mlp_inputs( ) return input_tensor, gate_weight, up_weight, down_weight - @skipIfXpu def test_rmsnorm_custom_op_autotune_with_dynamic_shape(self): """Test RMSNorm autotuning with multiple decomposition variants and dynamic shapes. @@ -255,7 +259,6 @@ def _create_decompose_k_inputs(self, m=256, k=65536, n=1024): ) return a, b, bias - @skipIfXpu def test_decompose_k_custom_op_autotune_dynamic_config_for_input_shape(self): """Test decompose_k autotuning with with epilogue fusion(matmul+bias+relu+scale) and dynamic config generation based on matmul input shapes. @@ -376,7 +379,6 @@ def reference_model(a, b, bias): msg=f"Failed for shape ({m}, {k}, {n})", ) - @skipIfXpu def test_multi_parameter_tuning(self): """Test autotuning with multiple parameters for combinatorial parameter exploration. @@ -476,7 +478,6 @@ def _( multi_param_op, (test_x, test_factor), expected_result, "MultiParam" ) - @skipIfXpu def test_range_based_static_shape_no_cond_dispatch(self): """Test dispatch code generation for static vs dynamic shapes. diff --git a/test/inductor/test_deterministic.py b/test/inductor/test_deterministic.py index c75e3f9961790..4a5b5a2829f50 100644 --- a/test/inductor/test_deterministic.py +++ b/test/inductor/test_deterministic.py @@ -16,7 +16,6 @@ instantiate_parametrized_tests, IS_FBCODE, parametrize, - skipIfXpu, ) from torch.testing._internal.inductor_utils import ( GPU_TYPE, @@ -48,7 +47,6 @@ def test_use_deterministic_algorithsm(self): finally: torch.use_deterministic_algorithms(old_val, warn_only=True) - @skipIfXpu(msg="pad_mm is not enabled for XPU.") @parametrize("deterministic", [False, True]) def test_mm_padding(self, deterministic): with inductor_config.patch(deterministic=deterministic): diff --git a/test/inductor/test_fp8.py b/test/inductor/test_fp8.py index c91c8b186e7c4..18705ae39835e 100644 --- a/test/inductor/test_fp8.py +++ b/test/inductor/test_fp8.py @@ -33,6 +33,7 @@ _quantize_rowwise, _quantize_tensorwise, _to_fp8_saturated, + GPU_TYPE, HAS_CPU, HAS_CUDA_AND_TRITON, is_big_gpu, @@ -331,7 +332,7 @@ def ln_fp8(x: Tensor, scale: Tensor, amax_buffer: Tensor): amax_buffer_compiled, amax_buffer, rtol=1e-2, atol=1e-2 ) - @onlyCUDA + @onlyOn(["cuda", "xpu"]) @unittest.skipIf(not PLATFORM_SUPPORTS_FP8, f8_msg) @parametrize("float8_dtype", (torch.float8_e4m3fn, torch.float8_e5m2)) @parametrize("shape", ("4,2048,4096",)) @@ -342,7 +343,7 @@ def test_layernorm_fp8_quant_benchmark( shape: str, keepdim: bool, ): - float8_dtype = _fix_fp8_dtype_for_rocm(float8_dtype, device="cuda") + float8_dtype = _fix_fp8_dtype_for_rocm(float8_dtype, device=GPU_TYPE) shape = [int(dim) for dim in shape.split(",")] batch_size, sequence_length, hidden_size = shape @@ -373,11 +374,11 @@ def ln_fp8(x: Tensor, scale: Tensor, amax_buffer: Tensor): compiled_ln_fp8_quant = torch.compile(ln_fp8, backend="inductor") x_shape = (batch_size, sequence_length, hidden_size) - x = torch.rand(*x_shape, device="cuda", dtype=torch.half) - scale = torch.tensor(0.2, device="cuda", dtype=torch.float) + x = torch.rand(*x_shape, device=GPU_TYPE, dtype=torch.half) + scale = torch.tensor(0.2, device=GPU_TYPE, dtype=torch.float) - amax_buffer_compiled = torch.zeros((1), device="cuda", dtype=torch.half) - amax_buffer = torch.zeros((1), device="cuda", dtype=torch.half) + amax_buffer_compiled = torch.zeros((1), device=GPU_TYPE, dtype=torch.half) + amax_buffer = torch.zeros((1), device=GPU_TYPE, dtype=torch.half) _ = compiled_ln_fp8_quant(x, scale, amax_buffer_compiled) compiled_latency = utils.do_bench_using_profiling( functools.partial(compiled_ln_fp8_quant, x, scale, amax_buffer_compiled) diff --git a/test/inductor/test_fused_attention.py b/test/inductor/test_fused_attention.py index 33cf454747c26..c7189b4614cef 100644 --- a/test/inductor/test_fused_attention.py +++ b/test/inductor/test_fused_attention.py @@ -189,11 +189,6 @@ def dot_prod_attention( ) def _test_insignificant_strides(self): - if self.device == "xpu": - self.skipTest( - "The operator 'aten::_scaled_dot_product_efficient_attention'" - " is not currently implemented for the XPU device. " - ) f32 = torch.float32 # repro taken from https://github.com/pytorch/pytorch/issues/124289 diff --git a/test/inductor/test_gpu_cpp_wrapper.py b/test/inductor/test_gpu_cpp_wrapper.py index 714a0e65d2ce1..35461b54e79e1 100644 --- a/test/inductor/test_gpu_cpp_wrapper.py +++ b/test/inductor/test_gpu_cpp_wrapper.py @@ -108,13 +108,6 @@ def test_fn(): "test_mm_plus_mm2_dynamic_shapes": test_torchinductor.TestFailure( ("gpu_wrapper",), is_skip=True ), - # ATen ops: scaled_dot_product_efficient_attention not implemented on XPU. - "test_scaled_dot_product_efficient_attention_xpu": test_torchinductor.TestFailure( - ("gpu_wrapper",), is_skip=False - ), - "test_scaled_dot_product_efficient_attention_xpu_dynamic_shapes": test_torchinductor.TestFailure( - ("gpu_wrapper",), is_skip=False - ), } # Skip only on CUDA as wrapper dynamic shapes passes on ROCm. @@ -189,11 +182,6 @@ class BaseTest(NamedTuple): tests: InductorTestCase = test_torchinductor.GPUTests() check_code: bool = True - # XPU Not implemented yet - XPU_BASE_TEST_SKIP = [ - "test_dynamic_shapes_persistent_reduction_mixed_x_dim", - ] - # Maintain two separate test lists for cuda and cpp for now for item in [ BaseTest("test_add_complex"), @@ -306,8 +294,6 @@ class BaseTest(NamedTuple): tests=test_select_algorithm.TestSelectAlgorithm(), ), ]: - if item.device == "xpu" and item.name in XPU_BASE_TEST_SKIP: - continue make_test_case(item.name, item.device, item.tests, check_code=item.check_code) test_torchinductor.copy_tests( From 6ff79233bae2e3e1d12ddcb56ef00dc050cfb0c5 Mon Sep 17 00:00:00 2001 From: "Deng, Daisy" Date: Mon, 23 Mar 2026 02:49:08 +0000 Subject: [PATCH 0009/1172] [XPU] Enable skipped inductor test on Intel GPU - generalize 21 inductor test files and enable xpu (#174058) Generalize 21 inductor tests and enable XPU: 1. Generailze code with requires_gpu_and_triton, torch.accelerator, instantiate_device_type_tests, HAS_GPU_AND_TRITON, GPU_TYPE 2. Remove skips for XPU according to latest test status 3. Added skipIfXpu for xpu limitations 4. Updated test_dtensor.TestDTensorCompile because it is referenced by test/inductor/test_compiled_autograd.py Pull Request resolved: https://github.com/pytorch/pytorch/pull/174058 Approved by: https://github.com/etaf, https://github.com/chuanqi129, https://github.com/mlazos ghstack dependencies: #174053, #174054, #174055, #174056, #174057 --- .../tensor/test_dtensor_compile.py | 10 +++--- test/inductor/test_aot_inductor.py | 4 ++- test/inductor/test_compiled_optimizers.py | 18 +++++----- test/inductor/test_cuda_repro.py | 22 +++++++----- test/inductor/test_fp8.py | 5 ++- test/inductor/test_inductor_freezing.py | 7 +--- test/inductor/test_inductor_utils.py | 12 +++++-- test/inductor/test_inplace_padding.py | 3 +- test/inductor/test_max_autotune.py | 1 - test/inductor/test_minifier_isolate.py | 7 ++-- test/inductor/test_mmdecomp.py | 14 ++++++-- test/inductor/test_multi_kernel.py | 4 +-- test/inductor/test_needs_exact_strides.py | 8 +++-- test/inductor/test_pattern_matcher.py | 1 - test/inductor/test_profiler.py | 9 +---- test/inductor/test_provenance_tracing.py | 34 +++++++++++-------- test/inductor/test_subgraph_choice.py | 3 -- test/inductor/test_torchinductor_opinfo.py | 5 ++- test/inductor/test_triton_heuristics.py | 4 ++- test/inductor/test_triton_kernels.py | 5 ++- test/inductor/test_unbacked_symints.py | 4 +-- 21 files changed, 104 insertions(+), 76 deletions(-) diff --git a/test/distributed/tensor/test_dtensor_compile.py b/test/distributed/tensor/test_dtensor_compile.py index 6d165783f0e1d..fa35693613c10 100644 --- a/test/distributed/tensor/test_dtensor_compile.py +++ b/test/distributed/tensor/test_dtensor_compile.py @@ -55,6 +55,7 @@ run_tests, skipIfHpu, skipIfTorchDynamo, + skipIfXpu, ) from torch.testing._internal.distributed._tensor.common_dtensor import ( DTensorTestBase, @@ -283,15 +284,16 @@ def forward(self, L_self_buffers_buffer_ : torch.distributed.tensor.DTensor, L_x ) self.assertExpectedInline( str(backend.fw_graphs[0].code).strip(), - """\ + f"""\ def forward(self, arg0_1, arg1_1, arg2_1): - _to_copy = torch.ops.aten._to_copy.default(arg1_1, dtype = torch.float64, layout = torch.strided, device = device(type='cuda', index=0)); arg1_1 = None + _to_copy = torch.ops.aten._to_copy.default(arg1_1, dtype = torch.float64, layout = torch.strided, device = device(type='{self.device_type}', index=0)); arg1_1 = None view = torch.ops.aten.view.default(_to_copy, [4, 4]); _to_copy = None add = torch.ops.aten.add.Tensor(arg0_1, view); arg0_1 = view = None view_1 = torch.ops.aten.view.default(add, [4, 4]); add = None return (view_1,)""", # noqa: B950 ) + @skipIfXpu(msg="AssertionError: torch-xpu-ops: 2958") @unittest.skipIf(not torch.accelerator.is_available(), "accelerator not available") def test_dtensor_basic_export(self): mesh = DeviceMesh(self.device_type, torch.arange(self.world_size)) @@ -344,9 +346,9 @@ def forward(self, args_0): # add is performed in _propagate_tensor_meta_non_cached, hence add_1 instead of add self.assertExpectedInline( str(joint_gm.code).strip(), - """\ + f"""\ def forward(self, arg0_1, arg1_1): - _to_copy = torch.ops.aten._to_copy.default(arg1_1, dtype = torch.float64, layout = torch.strided, device = device(type='cuda', index=0)); arg1_1 = None + _to_copy = torch.ops.aten._to_copy.default(arg1_1, dtype = torch.float64, layout = torch.strided, device = device(type='{self.device_type}', index=0)); arg1_1 = None view = torch.ops.aten.view.default(_to_copy, [4, 4]); _to_copy = None add = torch.ops.aten.add.Tensor(arg0_1, view); arg0_1 = view = None view_1 = torch.ops.aten.view.default(add, [4, 4]); add = None diff --git a/test/inductor/test_aot_inductor.py b/test/inductor/test_aot_inductor.py index 1b2bf752bcb5f..aa9c68a26e634 100644 --- a/test/inductor/test_aot_inductor.py +++ b/test/inductor/test_aot_inductor.py @@ -2080,7 +2080,9 @@ def forward(self, values, repeats, mask, embeddings, x, y, z, lst): } self.check_model(Repro(), example_inputs, dynamic_shapes=spec) - @skipIfXpu(msg="FlashAttentionForward headdim limitation on xpu - xpu-ops: 2698") + @skipIfXpu( + msg="FlashAttentionForward headdim limitation on xpu - torch-xpu-ops: 2698" + ) @unittest.skipIf( not PLATFORM_SUPPORTS_FLASH_ATTENTION, "Some archs don't support flash SDPA" ) diff --git a/test/inductor/test_compiled_optimizers.py b/test/inductor/test_compiled_optimizers.py index 802bea79bcc66..197e3148dc419 100644 --- a/test/inductor/test_compiled_optimizers.py +++ b/test/inductor/test_compiled_optimizers.py @@ -58,18 +58,19 @@ optim_db, optims, ) -from torch.testing._internal.common_utils import parametrize, skipIfRocm, skipIfWindows +from torch.testing._internal.common_utils import ( + parametrize, + skipIfRocm, + skipIfWindows, + skipIfXpu, +) from torch.testing._internal.inductor_utils import ( GPU_TYPE, HAS_CPU, HAS_GPU, has_triton, ) -from torch.testing._internal.triton_utils import ( - requires_cuda_and_triton, - requires_gpu, - requires_gpu_and_triton, -) +from torch.testing._internal.triton_utils import requires_gpu, requires_gpu_and_triton def get_inputs(optim): @@ -992,7 +993,7 @@ def loop(): @skipIfRocm(msg="ROCm may have different numerical behavior") -@requires_cuda_and_triton +@requires_gpu_and_triton class CompiledOptimizerBitwiseTests(TestCase): """ Tests that compiled optimizers produce bitwise identical results to eager @@ -1085,7 +1086,8 @@ def _test_optimizer_bitwise( def _make_bitwise_test(optim_cls, kernel_count=None, **optim_kwargs): @skipIfRocm(msg="ROCm may have different numerical behavior") - @requires_cuda_and_triton + @skipIfXpu(msg="AttributeError, torch-xpu-ops: #2999") + @requires_gpu_and_triton @config.patch( { "score_fusion_memory_threshold": 1, diff --git a/test/inductor/test_cuda_repro.py b/test/inductor/test_cuda_repro.py index e728cbb80c85e..c804d3e365b5d 100644 --- a/test/inductor/test_cuda_repro.py +++ b/test/inductor/test_cuda_repro.py @@ -193,6 +193,7 @@ def forward(self, x: torch.Tensor): self.assertEqual(compiled_out["ten0"], eager_out["ten0"]) self.assertEqual(compiled_out["ten1"], eager_out["ten1"]) + @skipIfXpu(msg="RuntimeError, torch-xpu-ops: 2891") def test_effn_attn_bias_padding(self): batch_size, num_heads, seq_len, head_dim = 2, 32, 512, 128 @@ -245,6 +246,7 @@ def fn( # Greatest absolute difference: 0.07861328125 at index (14, 13, 1008, 36) (up to 1e-05 allowed) # Greatest relative difference: 2.90625 at index (14, 13, 1008, 36) (up to 0.016 allowed) @skipIfRocmArch(MI350_ARCH) + @skipIfXpu(msg="RuntimeError, torch-xpu-ops: 2697") def test_effn_attn_bias_padding_misaligned(self): seqlen_start = 1008 @@ -880,6 +882,7 @@ def foo(x, y, z): check_lowp=False, ) + @skipIfXpu(msg="TypeError, torch-xpu-ops: 3004") def test_memory_history_inductor(self): def called_inside_compile(x, w, b): a = x @ w + b @@ -896,7 +899,7 @@ def fn(x, w, b): def record_memory_history(value: bool): if torch.xpu.is_available(): - torch.xpu._record_memory_history(value) + torch.xpu.memory._record_memory_history(value) else: torch.cuda.memory._record_memory_history(value) @@ -1452,6 +1455,7 @@ def fn(x, y, z): torch._dynamo.reset() gc.collect() + @skipIfXpu(msg="AssertionError, torch-xpu-ops: #3007") @unittest.skipIf( not PLATFORM_SUPPORTS_FLASH_ATTENTION, "flash attention not supported" ) @@ -1526,7 +1530,7 @@ def test_emulate_precision_casts_convert_element_type(self): torch.manual_seed(0) torch.cuda.manual_seed_all(0) - x = torch.rand(1000, device="cuda", dtype=torch.float32) + x = torch.rand(1000, device=device_type, dtype=torch.float32) def fn(x): x_bf16 = x.to(torch.bfloat16) @@ -1558,6 +1562,7 @@ def fn(inp, scale): self.assertEqual(expected, actual) + @skipIfXpu(msg="AssertionError, torch-xpu-ops: #2554") @torch._inductor.config.patch(emulate_precision_casts=True) def test_emulate_precision_casts_min_pow_chain(self): torch.manual_seed(0) @@ -2641,7 +2646,7 @@ def f(x, y): self.assertEqual(eager_out, compile_out) @skipIfXpu( - msg="Explicit attn_mask should not be set when is_causal=True - xpu-ops: 2802" + msg="Explicit attn_mask should not be set when is_causal=True - torch-xpu-ops: 2802" ) def test_qwen2_7b_sdpa_input_alignment_requires_recompile(self): # SDPA constraints ensures inputs have alignment (8). @@ -2772,7 +2777,7 @@ def compiled_divide(x, y): self.assertEqual(eager_div, compiled_div) - @skipIfXpu(msg="triton dependency - xpu-ops: 2554") + @skipIfXpu(msg="triton dependency - torch-xpu-ops: 2554") @config.patch({"eager_numerics.division_rounding": False}) @xfailIfROCm def test_truediv_base_not_bitwise_equivalent(self): @@ -2799,7 +2804,7 @@ def test_disabling_ftz_yields_subnormals(self): from decimal import Decimal x = -127.0 - x_ten = torch.tensor([x], dtype=torch.float32, device="cuda") + x_ten = torch.tensor([x], dtype=torch.float32, device=device_type) def fn(x): return 2.0**x @@ -2809,7 +2814,7 @@ def fn(x): self.assertTrue(compile_decimal > Decimal(0)) - @skipIfXpu(msg="Decimal object comparison failed - xpu-ops: 2810") + @skipIfXpu(msg="Decimal object comparison failed - torch-xpu-ops: 2810") @skipIfRocm(msg="ROCm preserves subnormals by default") @config.patch({"eager_numerics.disable_ftz": False}) def test_not_disabling_ftz_yields_zero(self): @@ -2826,6 +2831,7 @@ def fn(x): self.assertEqual(compile_decimal, Decimal(0)) + @skipIfXpu(msg="AssertionError: torch-xpu-ops: #3006") @config.patch( {"triton.use_block_ptr": True, "triton.codegen_upcast_to_fp32": False} ) @@ -2834,7 +2840,7 @@ def test_float16_reduction_with_int_output(self): def fn(input: torch.Tensor) -> torch.Tensor: return torch.argmax(input, dim=0) - input = torch.randn(20, 20, device="cuda", dtype=torch.float16) + input = torch.randn(20, 20, device=device_type, dtype=torch.float16) _, code = run_and_get_code(fn, input) # There should not be any conversions to float16 in this code, since the input # is already float16 and the output is int64. @@ -2848,7 +2854,7 @@ def test_reciprocal_precision_rounding(self): def fn(x): return torch.reciprocal(x) - x = torch.randn(1000, device="cuda", dtype=torch.float32) + 0.1 + x = torch.randn(1000, device=device_type, dtype=torch.float32) + 0.1 self.common(fn, [x]) diff --git a/test/inductor/test_fp8.py b/test/inductor/test_fp8.py index 18705ae39835e..1da8fba04b44e 100644 --- a/test/inductor/test_fp8.py +++ b/test/inductor/test_fp8.py @@ -27,7 +27,7 @@ skipCUDAIf, ) from torch.testing._internal.common_quantized import ceil_div, to_blocked -from torch.testing._internal.common_utils import parametrize, xfailIf +from torch.testing._internal.common_utils import parametrize, skipIfXpu, xfailIf from torch.testing._internal.inductor_utils import ( _quantize_blockwise, _quantize_rowwise, @@ -169,6 +169,9 @@ def fp8_cast(x): torch.testing.assert_close(y0_fp8, x, rtol=5e-1, atol=5e-1) torch.testing.assert_close(y1_fp8, x, rtol=5e-1, atol=5e-1) + @skipIfXpu( + msg="Conversions between float8_e5m2 and float8_e4m3fn is not supported, torch-xpu-ops: 2888" + ) @unittest.skipIf(not PLATFORM_SUPPORTS_FP8, f8_msg) def test_bad_cast(self, device): def fp8_cast(x, dtype): diff --git a/test/inductor/test_inductor_freezing.py b/test/inductor/test_inductor_freezing.py index f1474964d1865..7cc17f3a65f16 100644 --- a/test/inductor/test_inductor_freezing.py +++ b/test/inductor/test_inductor_freezing.py @@ -17,11 +17,7 @@ from torch._inductor.utils import override_lowering, run_and_get_code from torch.testing import FileCheck from torch.testing._internal.common_cuda import SM80OrLater, tf32_on_and_off -from torch.testing._internal.common_utils import ( - IS_FBCODE, - skipIfXpu, - TEST_WITH_SLOW_GRADCHECK, -) +from torch.testing._internal.common_utils import IS_FBCODE, TEST_WITH_SLOW_GRADCHECK # Make the helper files in test/ importable @@ -768,7 +764,6 @@ def foo(mod, inp): mod_eager = mod(x) self.assertEqual(foo(mod, x), mod_eager) - @skipIfXpu @unittest.skipIf(IS_FBCODE, "Not yet runnable in fbcode") @unittest.skipIf( TEST_WITH_SLOW_GRADCHECK, diff --git a/test/inductor/test_inductor_utils.py b/test/inductor/test_inductor_utils.py index 12468a09103b9..2871a579fe577 100644 --- a/test/inductor/test_inductor_utils.py +++ b/test/inductor/test_inductor_utils.py @@ -11,13 +11,19 @@ log = logging.getLogger(__name__) +device_type = ( + acc.type + if (acc := torch.accelerator.current_accelerator(check_available=True)) + else "cpu" +) + class TestBench(TestCase): @classmethod def setUpClass(cls): super().setUpClass() - x = torch.rand(1024, 10).cuda().half() - w = torch.rand(512, 10).cuda().half() + x = torch.rand(1024, 10).to(device_type).half() + w = torch.rand(512, 10).to(device_type).half() cls._bench_fn = functools.partial(torch.nn.functional.linear, x, w) def test_benchmarker(self): @@ -32,4 +38,4 @@ def test_do_bench_using_profiling(self): if __name__ == "__main__": - run_tests("cuda") + run_tests(device_type) diff --git a/test/inductor/test_inplace_padding.py b/test/inductor/test_inplace_padding.py index 8dc4a3a01a353..58ac6b5f37812 100644 --- a/test/inductor/test_inplace_padding.py +++ b/test/inductor/test_inplace_padding.py @@ -9,7 +9,7 @@ from torch._inductor.test_case import run_tests, TestCase from torch._inductor.utils import run_and_get_code from torch.testing import FileCheck -from torch.testing._internal.common_utils import serialTest +from torch.testing._internal.common_utils import serialTest, skipIfXpu from torch.testing._internal.inductor_utils import ( GPU_TYPE, HAS_GPU, @@ -258,6 +258,7 @@ def f(x, y): @requires_gpu_with_enough_memory(2e10) @inductor_config.patch(max_autotune=True) @serialTest() + @skipIfXpu(msg="AssertionError: torch-xpu-ops: #2997") def test_linear_and_cel_max_autotune(self): self.test_linear_and_cel() diff --git a/test/inductor/test_max_autotune.py b/test/inductor/test_max_autotune.py index 2db856dc67e17..e8c0d3ee12a4a 100644 --- a/test/inductor/test_max_autotune.py +++ b/test/inductor/test_max_autotune.py @@ -2233,7 +2233,6 @@ def misses(): self.assertEqual(misses(), 4) @fresh_cache() - @skipIfXpu @unittest.skipIf( config.cpp_wrapper, "decompose_k not supported for cpp_wrapper yet" ) diff --git a/test/inductor/test_minifier_isolate.py b/test/inductor/test_minifier_isolate.py index 61cf6e3961133..f1862b65f9bce 100644 --- a/test/inductor/test_minifier_isolate.py +++ b/test/inductor/test_minifier_isolate.py @@ -8,7 +8,6 @@ IS_MACOS, skipIfRocm, skipIfWindows, - skipIfXpu, TEST_WITH_ASAN, ) from torch.testing._internal.inductor_utils import GPU_TYPE @@ -40,11 +39,13 @@ def test_after_aot_cpu_runtime_error(self): self._test_after_aot_runtime_error("cpu", "") @skipIfRocm - @skipIfXpu @requires_gpu @inductor_config.patch("triton.inject_relu_bug_TESTING_ONLY", "runtime_error") def test_after_aot_gpu_runtime_error(self): - self._test_after_aot_runtime_error(GPU_TYPE, "device-side assert") + expected_error = ( + "injected assert fail" if GPU_TYPE == "xpu" else "device-side assert" + ) + self._test_after_aot_runtime_error(GPU_TYPE, expected_error) if __name__ == "__main__": diff --git a/test/inductor/test_mmdecomp.py b/test/inductor/test_mmdecomp.py index 6d5d012e733f1..1b25587b5398b 100644 --- a/test/inductor/test_mmdecomp.py +++ b/test/inductor/test_mmdecomp.py @@ -15,7 +15,12 @@ from torch.testing._internal.common_cuda import SM80OrLater from torch.testing._internal.common_device_type import instantiate_device_type_tests from torch.testing._internal.common_nn import NNTestCase -from torch.testing._internal.common_utils import IS_WINDOWS, parametrize, run_tests +from torch.testing._internal.common_utils import ( + IS_WINDOWS, + parametrize, + run_tests, + TEST_XPU, +) from torch.testing._internal.inductor_utils import GPU_TYPE, HAS_GPU @@ -132,7 +137,8 @@ def test_simple_mm(self, device, dtype): @unittest.skipIf(not HAS_GPU, "GPU tests require triton") @parametrize( - "dtype", [torch.float, torch.bfloat16] if SM80OrLater else [torch.float] + "dtype", + [torch.float, torch.bfloat16] if SM80OrLater or TEST_XPU else [torch.float], ) @parametrize("bs", [1, 2, 4, 10]) def test_batched_mm(self, device, dtype, bs): @@ -347,7 +353,9 @@ def test_dynamic_shape_mm(self, device, dtype): device_types = ("cpu", GPU_TYPE) -instantiate_device_type_tests(TestDecomp, globals(), only_for=device_types) +instantiate_device_type_tests( + TestDecomp, globals(), only_for=device_types, allow_xpu=True +) if __name__ == "__main__": # We don't support torch.compile() on Windows diff --git a/test/inductor/test_multi_kernel.py b/test/inductor/test_multi_kernel.py index 55f54756913db..bb5702540487a 100644 --- a/test/inductor/test_multi_kernel.py +++ b/test/inductor/test_multi_kernel.py @@ -111,7 +111,7 @@ def test_softmax(self, expect_multi_kernel=True): # TODO: bobrenjc93 to fix multi-kernel for ROCM @skipIfRocm @unittest.skipIf(not IS_BIG_GPU, "templates require big gpu") - @skipIfXpu(msg="https://github.com/intel/torch-xpu-ops/issues/2295") + @skipIfXpu(msg="driver issue, torch-xpu-ops: 2295") def test_triton_gemm(self): def fn(x, y): return x @ y @@ -136,7 +136,7 @@ def fn(x, y): self.assertEqual(ref, act) self.assertTrue(_contains_size_hint_multi_kernel_code(wrapper_code)) - @skipIfXpu(msg="https://github.com/intel/torch-xpu-ops/issues/2295") + @skipIfXpu(msg="driver issue, torch-xpu-ops: 2295") @requires_triton() # TODO: bobrenjc93 to fix multi-kernel for ROCM @skipIfRocm diff --git a/test/inductor/test_needs_exact_strides.py b/test/inductor/test_needs_exact_strides.py index ee3d4779881f2..dc7c3d55f967e 100644 --- a/test/inductor/test_needs_exact_strides.py +++ b/test/inductor/test_needs_exact_strides.py @@ -13,13 +13,15 @@ IS_LINUX, parametrize, ) -from torch.testing._internal.inductor_utils import HAS_CUDA_AND_TRITON +from torch.testing._internal.inductor_utils import HAS_GPU_AND_TRITON class TestNeedsExactStrides(InductorTestCase): @parametrize("dtype", [torch.float, torch.float8_e8m0fnu]) def test_custom_op(self, dtype): - device = "cuda" # float8_e8m0fnu errors on "cpu" + device = ( + torch.accelerator.current_accelerator() + ) # float8_e8m0fnu errors on "cpu" x = torch.ones(4, 4, 2, 2, device=device, dtype=torch.float8_e8m0fnu) other = torch.ones(4, 4, 2, 2, device=device, dtype=torch.float8_e8m0fnu) @@ -101,5 +103,5 @@ def f(x, other): instantiate_parametrized_tests(TestNeedsExactStrides) if __name__ == "__main__": - if IS_LINUX and HAS_CUDA_AND_TRITON: + if IS_LINUX and HAS_GPU_AND_TRITON: run_tests(needs="filelock") diff --git a/test/inductor/test_pattern_matcher.py b/test/inductor/test_pattern_matcher.py index dcd3ca6689cdf..f2ee6b9a2403c 100644 --- a/test/inductor/test_pattern_matcher.py +++ b/test/inductor/test_pattern_matcher.py @@ -136,7 +136,6 @@ def _test_fused_int_mm_mul_impl(self, fn, args, fused_int_mm_mul_expected=True): ref[indices], test[indices] ) # also checks that dtype is correct - # @skipIfXpu @skipCUDAIf(not SM80OrLater, "need sm_80") @inductor_config.patch( { diff --git a/test/inductor/test_profiler.py b/test/inductor/test_profiler.py index 99f06d5647dcf..2a90c55285f69 100644 --- a/test/inductor/test_profiler.py +++ b/test/inductor/test_profiler.py @@ -11,7 +11,7 @@ from torch import _dynamo as torchdynamo from torch._inductor import config from torch.profiler import ProfilerActivity, record_function -from torch.testing._internal.common_utils import skipIfXpu, TemporaryFileName +from torch.testing._internal.common_utils import TemporaryFileName from torch.testing._internal.inductor_utils import ( GPU_TYPE, HAS_GPU_AND_TRITON, @@ -26,10 +26,6 @@ class DynamoProfilerTests(torch._inductor.test_case.TestCase): - @skipIfXpu( - msg="AssertionError: False is not true, " - "https://github.com/intel/torch-xpu-ops/issues/2335" - ) @unittest.skipIf(not HAS_TRITON, "requires cuda & triton") def test_inductor_profiling_triton_launch(self): # Verify that we get some sort of CPU-side indication of triton kernel launches @@ -224,9 +220,6 @@ def fn(x, y): self.assertTrue(hooks_called["enter"]) self.assertTrue(hooks_called["exit"]) - @skipIfXpu( - msg="TypeError: list indices must be integers or slices, not str, https://github.com/intel/torch-xpu-ops/issues/2335" - ) @unittest.skipIf(not HAS_TRITON, "requires cuda & triton") def test_pt2_triton_attributes(self): from torch._inductor.codecache import code_hash diff --git a/test/inductor/test_provenance_tracing.py b/test/inductor/test_provenance_tracing.py index c58951ef8c7a1..a5e3e8dfc4ad0 100644 --- a/test/inductor/test_provenance_tracing.py +++ b/test/inductor/test_provenance_tracing.py @@ -29,10 +29,7 @@ from torch._inductor.virtualized import V from torch.testing._internal.common_utils import IS_MACOS from torch.testing._internal.inductor_utils import GPU_TYPE -from torch.testing._internal.triton_utils import ( - requires_cuda_and_triton, - requires_gpu_and_triton, -) +from torch.testing._internal.triton_utils import requires_gpu_and_triton try: @@ -613,7 +610,7 @@ def test_tlparse_kernel_stack_traces(self): @torch._inductor.config.patch( {"trace.provenance_tracking_level": 2, "max_autotune_gemm_backends": "ATEN"} ) - @requires_cuda_and_triton + @requires_gpu_and_triton def test_deferred_triton_kernels(self): def foo(m, inp): a = m(inp) @@ -621,8 +618,8 @@ def foo(m, inp): foo_c = torch.compile(mode="max-autotune-no-cudagraphs")(foo) - m = torch.nn.Linear(512, 512, bias=True).half().cuda() - inp = torch.rand([1, 512]).half().cuda() + m = torch.nn.Linear(512, 512, bias=True).half().to(GPU_TYPE) + inp = torch.rand([1, 512]).half().to(GPU_TYPE) with self._setup_provenance_capture() as payload_buffer: with torch.no_grad(): @@ -902,16 +899,21 @@ def forward(self, x, a, b, c): code ) - if self.device == "cuda": + if self.device == "cuda" or self.device == "xpu": + device_type = torch.accelerator.current_accelerator().type FileCheck().check( - """KernelContextGuard _ctx("aoti_torch_cuda_mm_out", R"(""" - ).check("AOTI_TORCH_ERROR_CODE_CHECK(aoti_torch_cuda_mm_out(").check( + f"""KernelContextGuard _ctx("aoti_torch_{device_type}_mm_out", R"(""" + ).check( + f"AOTI_TORCH_ERROR_CODE_CHECK(aoti_torch_{device_type}_mm_out(" + ).check( """KernelContextGuard _ctx("triton_poi_fused_addmm_relu_sigmoid_0", R"(""" ).check("call_triton_poi_fused_addmm_relu_sigmoid_0(").check( """KernelContextGuard _ctx("triton_poi_fused_mul_1", R"(""" ).check("call_triton_poi_fused_mul_1(").check( - """KernelContextGuard _ctx("aoti_torch_cuda_mm_out", R"(""" - ).check("AOTI_TORCH_ERROR_CODE_CHECK(aoti_torch_cuda_mm_out(").check( + f"""KernelContextGuard _ctx("aoti_torch_{device_type}_mm_out", R""" + ).check( + f"AOTI_TORCH_ERROR_CODE_CHECK(aoti_torch_{device_type}_mm_out(" + ).check( """ KernelContextGuard _ctx("triton_poi_fused_addmm_gelu_2", R"(""" ).check("call_triton_poi_fused_addmm_gelu_2(").run(code) else: @@ -942,15 +944,17 @@ class TestProvenanceTracingKernelContextCpu(TestCase): @unittest.skipIf(sys.platform == "darwin", "No CUDA on MacOS") -@unittest.skipIf(not torch.cuda.is_available(), "No CUDA") +@unittest.skipIf( + not torch.cuda.is_available() and not torch.xpu.is_available(), "No CUDA and no XPU" +) class TestProvenanceTracingKernelContextGpu(TestCase): - device = "cuda" + device = GPU_TYPE copy_tests( ProvenanceTracingKernelContextTemplate, TestProvenanceTracingKernelContextGpu, - "cuda", + GPU_TYPE, ) diff --git a/test/inductor/test_subgraph_choice.py b/test/inductor/test_subgraph_choice.py index 098b70b591787..27100f7b66b8f 100644 --- a/test/inductor/test_subgraph_choice.py +++ b/test/inductor/test_subgraph_choice.py @@ -7,7 +7,6 @@ from torch._inductor.lowering import register_lowering from torch._inductor.select_algorithm import autotune_select_algorithm from torch._inductor.test_case import run_tests, TestCase -from torch.testing._internal.common_utils import skipIfXpu from torch.testing._internal.inductor_utils import GPU_TYPE, HAS_CPU, HAS_GPU @@ -35,7 +34,6 @@ def _create_buffer(self, name, shape, dtype): layout=FixedLayout(torch.device(f"{GPU_TYPE}:0"), dtype=dtype, size=shape), ) - @skipIfXpu def test_subgraph_decompose_k(self): from torch._inductor.kernel.mm import aten_mm from torch._inductor.kernel.mm_common import mm_args @@ -96,7 +94,6 @@ def func(mat1, mat2): # Check same results of compiled result and regular torch.mm torch.testing.assert_close(res, a_in @ b_in, atol=1e-1, rtol=1e-1) - @skipIfXpu def test_subgraph_freeze_layout(self): from torch._inductor.kernel.mm_common import mm_args diff --git a/test/inductor/test_torchinductor_opinfo.py b/test/inductor/test_torchinductor_opinfo.py index e729d9575c8a2..d244a0742e6e6 100644 --- a/test/inductor/test_torchinductor_opinfo.py +++ b/test/inductor/test_torchinductor_opinfo.py @@ -233,6 +233,10 @@ def format_op(op): inductor_skips["xpu"] = {} +# torch-xpu-ops: #2956 +inductor_skips["xpu"]["lu"] = {f32} +inductor_skips["xpu"]["nn.functional.linear"] = {f16} + inductor_expected_failures_single_sample = defaultdict(dict) inductor_expected_failures_single_sample["cpu"] = { @@ -296,7 +300,6 @@ def format_op(op): i32, i64, }, # align with cuda. - ("linalg.pinv", "singular"): {f64}, # could not create a primitive "fft.fft": {f16}, "fft.fft2": {f16}, diff --git a/test/inductor/test_triton_heuristics.py b/test/inductor/test_triton_heuristics.py index 96bed8bf23711..413cdd88db5c5 100644 --- a/test/inductor/test_triton_heuristics.py +++ b/test/inductor/test_triton_heuristics.py @@ -270,7 +270,9 @@ def fn(x): res = torch.compile(fn)(x) self.assertEqual(ref, res) - @skipIfXpu(msg="lack _get_exceeding_shared_memory_checker support - xpu-ops: 2331") + @skipIfXpu( + msg="lack _get_exceeding_shared_memory_checker support - torch-xpu-ops: 2331" + ) @skipUnless(HAS_GPU_AND_TRITON, "requires gpu and triton") @parametrize("do_pruning", [False, True]) def test_prune_configs_over_shared_memory_limit(self, do_pruning): diff --git a/test/inductor/test_triton_kernels.py b/test/inductor/test_triton_kernels.py index 5e9c51e61ebf8..df71a8171ed1e 100644 --- a/test/inductor/test_triton_kernels.py +++ b/test/inductor/test_triton_kernels.py @@ -2819,7 +2819,10 @@ def fn(sz): self.assertEqual(actual, expected) @requires_gpu - @skipIfXpu(msg="`tl.inline_asm_elementwise` is not yet supported on Intel GPUs") + @skipIfXpu( + msg="`tl.inline_asm_elementwise` is not yet supported on Intel GPUs, " + "https://github.com/pytorch/pytorch/pull/167786" + ) @inductor_config.patch({"triton.autotune_at_compile_time": True}) @parametrize("quotes", ["single", "double"]) def test_kernel_inline_asm(self, quotes): diff --git a/test/inductor/test_unbacked_symints.py b/test/inductor/test_unbacked_symints.py index bd9576ac6b466..74fd0c5369a93 100644 --- a/test/inductor/test_unbacked_symints.py +++ b/test/inductor/test_unbacked_symints.py @@ -516,9 +516,9 @@ def fn(x): expected = fn(*example_inputs) torch.testing.assert_close(actual, expected) + @skipIfXpu(msg="FlashAttentionForward headdim limitation on xpu") @skipGPUIf(not HAS_GPU, "requires gpu and triton") @skipCUDAIf(not SM80OrLater, "Requires sm80 or later.") - @skipIfXpu(msg="_scaled_dot_product_flash_attention is not supported on XPU yet") @dynamo_config.patch({"capture_dynamic_output_shape_ops": True}) def test_sdpfa(self, device): if device == "cpu": @@ -543,9 +543,9 @@ def fn(x): x = torch.tensor([1.0, 0.0, 1.0, 0.0], device=device) torch.compile(fn, fullgraph=True)(x) + @skipIfXpu(msg="FlashAttentionForward headdim limitation on xpu") @skipGPUIf(not HAS_GPU, "requires gpu and triton") @skipCUDAIf(not SM80OrLater, "Requires sm80 or later.") - @skipIfXpu(msg="scaled_dot_product_attention is not supported on XPU yet") @dynamo_config.patch({"capture_dynamic_output_shape_ops": True}) def test_sdfpa_unbacked_strides(self, device): if device == "cpu": From 7b45702237852c500f256730439dddb726ad59c6 Mon Sep 17 00:00:00 2001 From: Paul Zhang Date: Mon, 23 Mar 2026 13:54:14 +0000 Subject: [PATCH 0010/1172] [Inductor] repeat_interleave fx graph runnable fix (#177909) (#177909) Summary: Pull Request resolved: https://github.com/pytorch/pytorch/pull/177909 Test Plan: `test_repeat_interleave_with_output_size` Reviewed By: nmacchioni Differential Revision: D97335376 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177909 Approved by: https://github.com/Lucaskabela --- test/dynamo/test_fx_graph_runnable.py | 38 +++++++++++++++++ torch/_dynamo/repro/after_aot.py | 59 +++++++++++++++++++++++++++ 2 files changed, 97 insertions(+) diff --git a/test/dynamo/test_fx_graph_runnable.py b/test/dynamo/test_fx_graph_runnable.py index 9ff32cbc2e60a..748e2c682d4bd 100644 --- a/test/dynamo/test_fx_graph_runnable.py +++ b/test/dynamo/test_fx_graph_runnable.py @@ -549,6 +549,44 @@ def f(view, weights): torch.compile(f)(view, weights) self._exec_and_verify_payload() + @torch._dynamo.config.patch(assume_static_by_default=False) + def test_repeat_interleave_with_output_size(self): + def f(data, repeats, output_size): + indices = torch.repeat_interleave(repeats, output_size=output_size.item()) + return data[indices] + + num_segments = 128 + data = torch.randn(1000, 16) + repeats = torch.randint(5, 15, (num_segments,), dtype=torch.int64) + output_size = repeats.sum() + + torch.compile(f, dynamic=True)(data, repeats, output_size) + + self._exec_and_verify_payload() + + # Verify the payload contains the repeat_interleave fixup + payload = self.buffer.getvalue().strip() + self.assertIn("def forward", payload) + self.assertIn("repeat_interleave", payload) + # Verify the fixup code is present + self.assertIn("# Fixup: ensure sum(repeats) == output_size", payload) + self.assertIn("_repeats.fill_", payload) + + def test_repeat_interleave_with_constant_output_size(self): + def f(data, repeats): + # output_size is a constant, not a dynamic input + indices = torch.repeat_interleave(repeats, output_size=1280) + return data[indices] + + num_segments = 128 + data = torch.randn(1000, 16) + repeats = torch.full((num_segments,), 10, dtype=torch.int64) + + torch.compile(f)(data, repeats) + self._exec_and_verify_payload() + payload = self.buffer.getvalue().strip() + self.assertNotIn("# Fixup: ensure sum(repeats) == output_size", payload) + @unittest.skipIf(IS_FBCODE or IS_SANDCASTLE, "Skip in fbcode/sandcastle") class TestFxGraphRunnableMultiProcessGroup(TestCase): diff --git a/torch/_dynamo/repro/after_aot.py b/torch/_dynamo/repro/after_aot.py index 37b2344d959b0..7ab24e0800d60 100644 --- a/torch/_dynamo/repro/after_aot.py +++ b/torch/_dynamo/repro/after_aot.py @@ -98,6 +98,39 @@ class TritonConstexpr: # type: ignore[no-redef] from .. import config +def _find_repeat_interleave_constraints( + gm: torch.fx.GraphModule, +) -> list[tuple[str, str]]: + """ + Find repeat_interleave operations with output_size constraints. + + Returns list of (repeats_placeholder_name, output_size_placeholder_name) pairs. + These represent constraints where sum(repeats) must equal output_size. + """ + constraints = [] + for node in gm.graph.nodes: + if ( + node.op != "call_function" + or "repeat_interleave" not in str(node.target) + or not node.args + ): + continue + + output_size_node = node.kwargs.get("output_size") + repeats_node = node.args[0] + + # Both must be FX nodes (not constants) and direct placeholders + if ( + isinstance(repeats_node, torch.fx.Node) + and isinstance(output_size_node, torch.fx.Node) + and repeats_node.op == "placeholder" + and output_size_node.op == "placeholder" + ): + constraints.append((str(repeats_node.target), str(output_size_node.target))) + + return constraints + + if TYPE_CHECKING: from collections.abc import Callable, Sequence @@ -701,6 +734,32 @@ def write_kernel_dependencies( ) model_str = f"{hint_lines}\n\n{model_str}" + # Add fixup code for repeat_interleave constraints + # When inputs are regenerated randomly, sum(repeats) != output_size + # This fixup adjusts the repeats tensor to satisfy the constraint + constraints = _find_repeat_interleave_constraints(gm) + if constraints: + placeholder_to_idx = {name: idx for idx, name in enumerate(placeholder_targets)} + for repeats_name, output_size_name in constraints: + repeats_idx = placeholder_to_idx.get(repeats_name) + output_size_idx = placeholder_to_idx.get(output_size_name) + if repeats_idx is not None and output_size_idx is not None: + # Guard with hasattr since NopInputReader doesn't have args + writer._lines.append( + "# Fixup: ensure sum(repeats) == output_size for repeat_interleave" + ) + writer._lines.append("if hasattr(reader, 'args'):") + writer._lines.append(f" _repeats = reader.args[{repeats_idx}]") + writer._lines.append( + f" _output_size = reader.args[{output_size_idx}]" + ) + writer._lines.append( + " if isinstance(_repeats, torch.Tensor) and _repeats.dtype == torch.int64:" + ) + writer._lines.append(" _n = _repeats.numel()") + writer._lines.append(" _repeats.fill_(_output_size // _n)") + writer._lines.append(" _repeats[:_output_size % _n] += 1") + load_args_lines = writer.lines() load_args_code = "\n".join(load_args_lines) model_str += load_args_code + "\n" From 11d35a3ad40d50ae7cdae5635b33359a4c60f208 Mon Sep 17 00:00:00 2001 From: Witold Dziurdz Date: Mon, 23 Mar 2026 15:43:54 +0000 Subject: [PATCH 0011/1172] Add GEMM configs to XPU autotuning heuristic (#177647) Add two XPU-specific mm configs to improve autotuning for tall-skinny GEMM shapes (e.g. M=10000, N=64, K=64, fp16). Benchmarked on BMG. The new configs cover BLOCK_N=64 which matches the N dimension exactly, reducing workgroup count and improving GPU occupancy. Fixes: https://github.com/intel/intel-xpu-backend-for-triton/issues/6012 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177647 Approved by: https://github.com/EikanWang, https://github.com/jansel --- torch/_inductor/template_heuristics/triton.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/torch/_inductor/template_heuristics/triton.py b/torch/_inductor/template_heuristics/triton.py index 793a6ae49a3aa..9d51077d668c3 100644 --- a/torch/_inductor/template_heuristics/triton.py +++ b/torch/_inductor/template_heuristics/triton.py @@ -1801,6 +1801,10 @@ class XPUConfigHeuristic(BaseConfigHeuristic): def __init__(self) -> None: super().__init__() + self.mm_configs = self.mm_configs + [ + GemmConfig(32, 64, 128, 2, 2), + GemmConfig(64, 64, 32, 2, 8), + ] self.xpu_default_flex_config = { (torch.float32, 64): FlexConfig(128, 32, 1, 16), (torch.float32, 128): FlexConfig(128, 32, 1, 16), From 1d3abfe96b4d1ab0232e515e310fed45cb6c99f8 Mon Sep 17 00:00:00 2001 From: PyTorch MergeBot Date: Mon, 23 Mar 2026 15:45:53 +0000 Subject: [PATCH 0012/1172] Revert "[dynamo] Adjust more tests to prepare for inline_inbuilt_nn_module deprecation (#178067)" This reverts commit 27c9a3533a6717ea1c8439ba9ff95bebe84d0204. Reverted https://github.com/pytorch/pytorch/pull/178067 on behalf of https://github.com/yangw-dev due to it seems this maybe the cause of FAILED CONSISTENTLY: test/export/test_hop.py::TestHOPCUDA::test_retrace_export_inline_asm_elementwise_simple_cuda_float32, please take a look ([comment](https://github.com/pytorch/pytorch/pull/178067#issuecomment-4111617299)) --- test/dynamo/test_activation_checkpointing.py | 1 + test/dynamo/test_decorators.py | 3 +- test/dynamo/test_dynamic_shapes.py | 8 +++ test/dynamo/test_higher_order_ops.py | 59 ++++++++++++++--- test/dynamo/test_misc.py | 6 +- test/dynamo/test_modules.py | 20 ++++-- test/dynamo/test_repros.py | 6 +- test/dynamo/test_structured_trace.py | 66 ++++++++++++++++++-- test/dynamo/test_utils.py | 22 ++++--- test/functorch/test_control_flow.py | 9 +-- test/inductor/test_cudagraph_trees.py | 18 +++++- 11 files changed, 180 insertions(+), 38 deletions(-) diff --git a/test/dynamo/test_activation_checkpointing.py b/test/dynamo/test_activation_checkpointing.py index 23fa43efb4864..7c7717441217c 100644 --- a/test/dynamo/test_activation_checkpointing.py +++ b/test/dynamo/test_activation_checkpointing.py @@ -1871,6 +1871,7 @@ def forward(self, x): @requires_distributed() @requires_cuda_and_triton def test_dynamo_does_not_trace_getattr_as_top_frame(self): + # inline_inbuilt_nn_modules is a proxy to emulate what FSDP tests do. from torch.distributed.algorithms._checkpoint.checkpoint_wrapper import ( CheckpointWrapper, ) diff --git a/test/dynamo/test_decorators.py b/test/dynamo/test_decorators.py index 273b915c46a57..96d25d9022ebc 100644 --- a/test/dynamo/test_decorators.py +++ b/test/dynamo/test_decorators.py @@ -1323,7 +1323,8 @@ def forward(self, a, *args): def _test_mark_static_address(self, guarded): # This test verifies that dynamo properly marks inputs as static # when using the mark_static_address API. - # We expect the tensor to be present in the buffers attribute of the graph. + # For both inline_inbuilt_nn_modules True and False, we expect the + # tensor to be present in the buffers attribute of the graph. compiles_with_buffers = 0 compiles = 0 diff --git a/test/dynamo/test_dynamic_shapes.py b/test/dynamo/test_dynamic_shapes.py index a585c65c82157..3bebcfc345b70 100644 --- a/test/dynamo/test_dynamic_shapes.py +++ b/test/dynamo/test_dynamic_shapes.py @@ -1,4 +1,5 @@ # Owner(s): ["module: dynamo"] +import unittest import warnings from torch._dynamo import config @@ -84,6 +85,13 @@ def make_dynamic_cls(cls): make_dynamic_cls(test) del test +if TEST_Z3: + if not config.inline_inbuilt_nn_modules: + # TODO model is somehow not being freed when z3 is available + unittest.expectedFailure( + DynamicShapesMiscTests.test_parameter_free_dynamic_shapes # noqa: F821 + ) + # Test takes too long ~700s as of 414a1fd29f04d06e41b7f895368dd1f83a4be29d DynamicShapesExportTests.test_retracibility_dynamic_shapes = slowTest( # noqa: F821 DynamicShapesExportTests.test_retracibility_dynamic_shapes # noqa: F821 diff --git a/test/dynamo/test_higher_order_ops.py b/test/dynamo/test_higher_order_ops.py index 0c6089d1eee80..e3ca2aae54a2a 100644 --- a/test/dynamo/test_higher_order_ops.py +++ b/test/dynamo/test_higher_order_ops.py @@ -2583,6 +2583,10 @@ def f(x): # 3 args - 1 for input, and other 2 for the weight and bias self.assertTrue(len(wrap_node.args), 3) + # Check that the linear bias and weight are getattr in the outer graph + if not torch._dynamo.config.inline_inbuilt_nn_modules: + self.assertTrue(len(dict(backend.graphs[0].named_parameters())) == 2) + # Check that the inner function has one op and its a linear op body_function = getattr(backend.graphs[0], wrap_node.args[0].name) self.assertEqual(op_count(body_function), 1) @@ -2711,6 +2715,10 @@ def f(x): wrap_node = find_first_node(backend.graphs[0], wrap) self.assertTrue(len(wrap_node.args), 3) + # Check that the linear bias and weight are getattr in the outer graph + if not torch._dynamo.config.inline_inbuilt_nn_modules: + self.assertTrue(len(dict(backend.graphs[0].named_parameters())) == 2) + # Check that the inner function has one op and its a linear op body_function = getattr(backend.graphs[0], wrap_node.args[0].name) self.assertEqual(op_count(body_function), 1) @@ -4496,9 +4504,10 @@ def wrapper_fn(model, params, inputs, targets): return actual = normalize_gm(wrapped_gm.print_readable(print_output=False)) - self.assertExpectedInline( - actual, - """\ + if torch._dynamo.config.inline_inbuilt_nn_modules: + self.assertExpectedInline( + actual, + """\ class GraphModule(torch.nn.Module): def forward(self, L_model_parameters_weight_: "f32[3, 3]", L_model_parameters_bias_: "f32[3]", L_inputs_: "f32[64, 3]", L_targets_: "f32[64, 3]"): l_model_parameters_weight_ = L_model_parameters_weight_ @@ -4511,7 +4520,22 @@ def forward(self, L_model_parameters_weight_: "f32[3, 3]", L_model_parameters_bi mse_loss: "f32[]" = torch.nn.functional.mse_loss(prediction, l_targets_); prediction = l_targets_ = None return (mse_loss,) """, - ) + ) + else: + self.assertExpectedInline( + actual, + """\ +class GraphModule(torch.nn.Module): + def forward(self, L_inputs_: "f32[64, 3]", L_targets_: "f32[64, 3]"): + l_inputs_ = L_inputs_ + l_targets_ = L_targets_ + + prediction: "f32[64, 3]" = self.model(l_inputs_); l_inputs_ = None + + mse_loss: "f32[]" = torch.nn.functional.mse_loss(prediction, l_targets_); prediction = l_targets_ = None + return (mse_loss,) +""", + ) def test_functional_call_sequential_params_and_buffers(self): # copied from test/test_stateless.py @@ -4542,7 +4566,8 @@ def wrapper_fn(model, params, buffers, inputs): return actual = normalize_gm(wrapped_gm.print_readable(print_output=False)) - expected = """\ + if torch._dynamo.config.inline_inbuilt_nn_modules: + expected = """\ class GraphModule(torch.nn.Module): def forward(self, L_inputs_: "f32[1, 1]", L_model_modules_l1_parameters_weight_: "f32[1, 1]", L_model_modules_l1_parameters_bias_: "f32[1]", L_model_buffers_buffer_: "f32[1]"): l_inputs_ = L_inputs_ @@ -4553,11 +4578,25 @@ def forward(self, L_inputs_: "f32[1, 1]", L_model_modules_l1_parameters_weight_: add: "f32[1, 1]" = linear + l_model_buffers_buffer_; linear = l_model_buffers_buffer_ = None return (add,) """ - # We found Windows/Linux have some empty line difference, empty_line_normalizer will help fix it. - self.assertExpectedInline( - empty_line_normalizer(actual), - empty_line_normalizer(normalize_gm(expected)), - ) + # We found Windows/Linux have some empty line difference, empty_line_normalizer will help fix it. + self.assertExpectedInline( + empty_line_normalizer(actual), + empty_line_normalizer(normalize_gm(expected)), + ) + else: + self.assertExpectedInline( + actual, + """\ +class GraphModule(torch.nn.Module): + def forward(self, L_x_: "f32[1, 1]"): + l_x_ = L_x_ + + l__self___l1: "f32[1, 1]" = self.L__self___l1(l_x_); l_x_ = None + l__self___buffer: "f32[1]" = self.L__self___buffer + add: "f32[1, 1]" = l__self___l1 + l__self___buffer; l__self___l1 = l__self___buffer = None + return (add,) +""", + ) def test_grad(self): counters.clear() diff --git a/test/dynamo/test_misc.py b/test/dynamo/test_misc.py index 6b2397192e9f6..a94c1d27b7471 100644 --- a/test/dynamo/test_misc.py +++ b/test/dynamo/test_misc.py @@ -7248,7 +7248,11 @@ def body(x): mod = Module() - error_message = r"Higher Order Operator: torch\.ops\.higher_order\.map_impl" + error_message = "" + if torch._dynamo.config.inline_inbuilt_nn_modules: + error_message = r"Higher Order Operator: torch\.ops\.higher_order\.map_impl" + else: + error_message = "Can't inplace modify module params/buffers" with self.assertRaisesRegex( torch._dynamo.exc.UncapturedHigherOrderOpError, error_message diff --git a/test/dynamo/test_modules.py b/test/dynamo/test_modules.py index 28ec48a82b45d..054ac7dcb7c8e 100644 --- a/test/dynamo/test_modules.py +++ b/test/dynamo/test_modules.py @@ -2140,7 +2140,10 @@ def forward(self, x): ): x = torch.randn(*size, requires_grad=True) mod(x) - self.assertEqual(cnts.frame_count, 1) + if torch._dynamo.config.inline_inbuilt_nn_modules: + self.assertEqual(cnts.frame_count, 1) + else: + self.assertEqual(cnts.frame_count, num_submodules) def test_inline_inbuilt_nn_modules(self): size = (10, 10) @@ -2222,7 +2225,10 @@ def forward(self, x): ]: x = torch.randn(size) mod(x) - self.assertEqual(cnts.frame_count, 2) + if torch._dynamo.config.inline_inbuilt_nn_modules: + self.assertEqual(cnts.frame_count, 2) + else: + self.assertEqual(cnts.frame_count, 2 * num_submodules) def test_recursion(self): mod = MockModule() @@ -2794,10 +2800,16 @@ def foo(mod, x): mod = Mod() foo(mod, torch.rand([4])) - self.assertEqual(compiles_without_buffers, 1) + if torch._dynamo.config.inline_inbuilt_nn_modules: + self.assertEqual(compiles_without_buffers, 1) + else: + self.assertEqual(compiles_without_buffers, 0) foo(mod, torch.rand([4], dtype=torch.half)) - self.assertEqual(compiles_without_buffers, 2) + if torch._dynamo.config.inline_inbuilt_nn_modules: + self.assertEqual(compiles_without_buffers, 2) + else: + self.assertEqual(compiles_without_buffers, 1) class Mod2(Mod): def __setattr__(self, name, value): diff --git a/test/dynamo/test_repros.py b/test/dynamo/test_repros.py index 66541b9cacfbf..7c5bf178fc44d 100644 --- a/test/dynamo/test_repros.py +++ b/test/dynamo/test_repros.py @@ -1399,8 +1399,12 @@ def test_reformer_eval(self): def test_reformer_train(self): with torch.enable_grad(): cnt = self._reformer(nopython=False) + expected_op_count = ( + """10""" if torch._dynamo.config.inline_inbuilt_nn_modules else """4""" + ) + self.assertExpectedInline(cnt.frame_count, """1""") - self.assertExpectedInline(cnt.op_count, """10""") + self.assertExpectedInline(cnt.op_count, expected_op_count) def test_longformer_chunk(self): input1 = torch.randn([1, 4096, 1]) diff --git a/test/dynamo/test_structured_trace.py b/test/dynamo/test_structured_trace.py index a0230ca562b8b..526eb88cbda8f 100644 --- a/test/dynamo/test_structured_trace.py +++ b/test/dynamo/test_structured_trace.py @@ -655,9 +655,67 @@ def forward(self, x): dist.destroy_process_group() - self.assertExpectedInline( - self.buffer.getvalue(), - """\ + if not torch._dynamo.config.inline_inbuilt_nn_modules: + self.assertExpectedInline( + self.buffer.getvalue(), + """\ +{"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} +{"artifact": {"name": "dynamo_graph_break_reason", "encoding": "string"}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 1, "has_payload": "HASH"} +{"compilation_metrics": "METRICS", "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 1} +{"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} +{"artifact": {"name": "dynamo_graph_break_reason", "encoding": "string"}, "rank": 0, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 1, "frame_compile_id": 0, "attempt": 1, "has_payload": "HASH"} +{"compilation_metrics": "METRICS", "rank": 0, "frame_id": 1, "frame_compile_id": 0, "attempt": 1} +{"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 2, "frame_compile_id": 0, "attempt": 0} +{"artifact": {"name": "dynamo_graph_break_reason", "encoding": "string"}, "rank": 0, "frame_id": 2, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 2, "frame_compile_id": 0, "attempt": 1, "has_payload": "HASH"} +{"compilation_metrics": "METRICS", "rank": 0, "frame_id": 2, "frame_compile_id": 0, "attempt": 1} +{"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} +{"describe_storage": {"id": 0, "describer_id": "ID", "size": 4194304}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} +{"describe_tensor": {"id": 0, "ndim": 2, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1024, 1024], "is_leaf": true, "stride": [1024, 1], "storage": 0, "view_func": "VIEW_FUNC", "describer_id": "ID"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} +{"describe_source": {"describer_id": "ID", "id": 0, "source": "L['x']"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} +{"dynamo_output_graph": {"sizes": {"l_x_": [1024, 1024], "l__self___layers_0": [1024, 1024], "l__self___layers_1": [1024, 1024]}}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"optimize_ddp_split_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"optimize_ddp_split_child": {"name": "submod_0"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"optimize_ddp_split_child": {"name": "submod_1"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"describe_storage": {"id": 0, "describer_id": "ID", "size": 4194304}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} +{"describe_tensor": {"id": 0, "ndim": 2, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1024, 1024], "is_leaf": true, "stride": [1024, 1], "storage": 0, "view_func": "VIEW_FUNC", "describer_id": "ID"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} +{"describe_source": {"describer_id": "ID", "id": 0, "source": "L['x']"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} +{"artifact": {"name": "before_pre_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "after_pre_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"aot_joint_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "torch._functorch.config", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "aot_forward_graph_fw_metadata", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"aot_forward_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"aot_backward_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "fx_graph_runnable", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "before_post_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "inductor_post_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"inductor_output_code": {"filename": "FILENAME", "file_path": "FILENAME"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "fx_graph_cache_miss", "encoding": "json"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "aotautograd_cache_bypass", "encoding": "json"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "before_pre_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "after_pre_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"aot_joint_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "torch._functorch.config", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "aot_forward_graph_fw_metadata", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"aot_forward_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"aot_backward_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "fx_graph_runnable", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "before_post_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "inductor_post_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"inductor_output_code": {"filename": "FILENAME", "file_path": "FILENAME"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "fx_graph_cache_miss", "encoding": "json"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"artifact": {"name": "aotautograd_cache_bypass", "encoding": "json"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} +{"compilation_metrics": "METRICS", "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} +""", # noqa: B950 + ) + else: + self.assertExpectedInline( + self.buffer.getvalue(), + """\ {"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} {"artifact": {"name": "dynamo_graph_break_reason", "encoding": "string"}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} {"describe_storage": {"id": 0, "describer_id": "ID", "size": 4194304}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 1} @@ -740,7 +798,7 @@ def forward(self, x): {"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 4, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} {"compilation_metrics": "METRICS", "rank": 0, "frame_id": 4, "frame_compile_id": 0, "attempt": 0} """, # noqa: B950 - ) + ) self.assertParses() diff --git a/test/dynamo/test_utils.py b/test/dynamo/test_utils.py index 721f99bc317d4..2cd19fb3334d6 100644 --- a/test/dynamo/test_utils.py +++ b/test/dynamo/test_utils.py @@ -84,6 +84,7 @@ def test_larger_multiplier_for_even_smaller_tensor(self): @dynamo_config.patch( { "log_compilation_metrics": True, + "inline_inbuilt_nn_modules": False, } ) def test_graph_break_counting(self): @@ -444,6 +445,7 @@ def backward(grad_output): @dynamo_config.patch( { "log_compilation_metrics": True, + "inline_inbuilt_nn_modules": False, } ) @inductor_config.patch( @@ -686,7 +688,7 @@ def filter_expected(s: str) -> str: 'compile_time_autotune_time_us': None, 'compiler_config': None, 'compliant_custom_ops': set(), - 'config_inline_inbuilt_nn_modules': True, + 'config_inline_inbuilt_nn_modules': False, 'config_suppress_errors': False, 'cuda_version': None, 'cudagraph_skip_reason': None, @@ -705,11 +707,11 @@ def filter_expected(s: str) -> str: 'fail_user_frame_lineno': None, 'frame_key': '1', 'gc_time_us': 0, - 'graph_input_count': 3, - 'graph_node_count': 5, + 'graph_input_count': 1, + 'graph_node_count': 3, 'graph_node_shapes': None, 'graph_op_count': 1, - 'guard_count': 31, + 'guard_count': 10, 'has_guarded_code': True, 'inductor_code_gen_cumulative_compile_time_us': 0, 'inductor_compile_time_s': 0.0, @@ -779,7 +781,7 @@ def filter_expected(s: str) -> str: 'compile_time_autotune_time_us': None, 'compiler_config': None, 'compliant_custom_ops': set(), - 'config_inline_inbuilt_nn_modules': True, + 'config_inline_inbuilt_nn_modules': False, 'config_suppress_errors': False, 'cuda_version': None, 'cudagraph_skip_reason': None, @@ -798,11 +800,11 @@ def filter_expected(s: str) -> str: 'fail_user_frame_lineno': None, 'frame_key': '1', 'gc_time_us': 0, - 'graph_input_count': 3, - 'graph_node_count': 5, + 'graph_input_count': 1, + 'graph_node_count': 3, 'graph_node_shapes': None, 'graph_op_count': 1, - 'guard_count': 31, + 'guard_count': 10, 'has_guarded_code': True, 'inductor_code_gen_cumulative_compile_time_us': 0, 'inductor_compile_time_s': 0.0, @@ -886,7 +888,7 @@ def filter_expected(s: str) -> str: 'compile_time_autotune_time_us': None, 'compiler_config': None, 'compliant_custom_ops': None, - 'config_inline_inbuilt_nn_modules': True, + 'config_inline_inbuilt_nn_modules': False, 'config_suppress_errors': False, 'cuda_version': None, 'cudagraph_skip_reason': None, @@ -979,7 +981,7 @@ def filter_expected(s: str) -> str: 'compile_time_autotune_time_us': None, 'compiler_config': None, 'compliant_custom_ops': None, - 'config_inline_inbuilt_nn_modules': True, + 'config_inline_inbuilt_nn_modules': False, 'config_suppress_errors': False, 'cuda_version': None, 'cudagraph_skip_reason': None, diff --git a/test/functorch/test_control_flow.py b/test/functorch/test_control_flow.py index 4c5a6b0cf0993..72df12ad1822f 100644 --- a/test/functorch/test_control_flow.py +++ b/test/functorch/test_control_flow.py @@ -5910,9 +5910,10 @@ def test_while_loop_simple_with_linear_compile_check_graph(self): torch.compile(fn, backend=backend)(*inp) self.assertEqual(len(backend.graphs), 1) gm = backend.graphs[0] - self.assertExpectedInline( - normalize_gm(gm.print_readable(print_output=False)), - """\ + if torch._dynamo.config.inline_inbuilt_nn_modules: + self.assertExpectedInline( + normalize_gm(gm.print_readable(print_output=False)), + """\ class GraphModule(torch.nn.Module): def forward(self, L_iter_: "i64[]", L_x_: "f32[2, 2]", L_self_buffers_dec_: "i64[]", L_self_modules_linear_parameters_weight_: "f32[2, 2]", L_self_modules_linear_parameters_bias_: "f32[2]"): l_iter_ = L_iter_ @@ -5940,7 +5941,7 @@ def forward(self, child_2: "i64[]", child_3: "f32[2, 2]", l_self_buffers_dec__co child_4: "f32[2, 2]" = torch._C._nn.linear(child_3, l_self_modules_linear_parameters_weight__body_fn, l_self_modules_linear_parameters_bias__body_fn); child_3 = l_self_modules_linear_parameters_weight__body_fn = l_self_modules_linear_parameters_bias__body_fn = None return (child, child_4) """, # noqa: B950 - ) + ) def test_while_loop_nested2_traced(self): fn, inp = WHILE_LOOP_TESTS["nested2"] diff --git a/test/inductor/test_cudagraph_trees.py b/test/inductor/test_cudagraph_trees.py index 40c5f2e8247df..3fa2ec8ffc127 100644 --- a/test/inductor/test_cudagraph_trees.py +++ b/test/inductor/test_cudagraph_trees.py @@ -22,7 +22,7 @@ from torch._inductor.codecache import FxGraphCache from torch._inductor.compile_fx import compile_fx_inner from torch._inductor.cudagraph_trees import cudagraphify_impl as tree_cudagraphify_impl -from torch._inductor.cudagraph_utils import PlaceholderInfo +from torch._inductor.cudagraph_utils import FunctionID, PlaceholderInfo from torch._inductor.test_case import TestCase as InductorTestCase from torch._inductor.utils import run_and_get_code from torch._ops import OpOverload @@ -3303,8 +3303,20 @@ def forward(self, x) -> torch.Tensor: foo.static_tensor = torch.ones((2, 2), device="cuda") foo.goo.linear.bias = torch.nn.Parameter(torch.ones((2,), device="cuda")) - for _ in range(3): - foo(inp) + if torch._dynamo.config.inline_inbuilt_nn_modules: + for _ in range(3): + foo(inp) + else: + # Run with specific function id to avoid dynamo recompiling + self.get_manager().run( + [ + foo.goo.linear.weight, + foo.goo.linear.bias, + foo.static_tensor, + inp, + ], + FunctionID(0), + ) self.assertEqual(self.get_manager().new_graph_id().id, 2) From 57927012e4360a14acf8f48801a1f4f2c49a32ad Mon Sep 17 00:00:00 2001 From: Haoyu Zhang Date: Mon, 23 Mar 2026 16:04:43 +0000 Subject: [PATCH 0013/1172] [ROCm] Reland: Enable expandable segments (#173330) (#177974) Summary: Original pull request: https://github.com/pytorch/pytorch/pull/173330 Fixes https://github.com/pytorch/pytorch/issues/168737. Fixes https://github.com/pytorch/pytorch/issues/168736. The original diff enabled expandable segments for ROCm by adding `#ifdef USE_ROCM` guards throughout CUDACachingAllocator.cpp to use HIP APIs (hipMemAddressReserve, hipMemCreate, hipMemMap, etc.) instead of CUDA driver APIs when building for ROCm. Root cause: In HIP/ROCm 6.2.1, the field name for memory allocation properties is `requestedHandleType` (singular), not `requestedHandleTypes` (plural) as in CUDA. Additionally, `hipMemHandleTypeFabric` does not exist in HIP, so the `CU_MEM_HANDLE_TYPE_FABRIC` assignment must be skipped on ROCm. Fix applied on top of the original diff (from D96652342): - Use `prop.requestedHandleType = hipMemHandleTypePosixFileDescriptor` under `#ifdef USE_ROCM` (singular field name, HIP constant) - Use `prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR` for CUDA (plural field name, CUDA constant) - Skip the `CU_MEM_HANDLE_TYPE_FABRIC` assignment entirely on ROCm under `#ifndef USE_ROCM`, as `hipMemHandleTypeFabric` does not exist in HIP Co-authored-by: Prachi Gupta prachi.gupta@amd.com Co-authored-by: Jeff Daily jeff.daily@amd.com Co-authored-by: moonshadow-25 moonshadow-25@users.noreply.github.com Co-authored-by: Vighanesh Sharma vighaneshsharma@gmail.com Test Plan: ``` fbpkg build //aps_models/ads/ecosystem/eval/cogwheel_tests/amd:cogwheel_aps_ads_icvr_kd_eval_amd_test_harness --build-remote ``` https://www.internalfb.com/sandcastle/workflow/1049338713192153464 Differential Revision: D97211385 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177974 Approved by: https://github.com/jeffdaily, https://github.com/echen4096 --- c10/cuda/CUDAAllocatorConfig.h | 2 +- c10/cuda/CUDACachingAllocator.cpp | 122 +++++++++++++++++++++++- test/distributed/test_cupy_as_tensor.py | 6 +- test/test_cuda.py | 9 +- test/test_cuda_expandable_segments.py | 9 +- torch/_C/__init__.pyi.in | 1 + torch/_dynamo/trace_rules.py | 1 + torch/csrc/DeviceAccelerator.cpp | 4 + 8 files changed, 142 insertions(+), 12 deletions(-) diff --git a/c10/cuda/CUDAAllocatorConfig.h b/c10/cuda/CUDAAllocatorConfig.h index 4e6097a406bc2..cd9c9b86285c4 100644 --- a/c10/cuda/CUDAAllocatorConfig.h +++ b/c10/cuda/CUDAAllocatorConfig.h @@ -34,7 +34,7 @@ class C10_CUDA_API CUDAAllocatorConfig { static bool expandable_segments() { bool enabled = c10::CachingAllocator::AcceleratorAllocatorConfig:: use_expandable_segments(); -#ifndef PYTORCH_C10_DRIVER_API_SUPPORTED +#if !defined(PYTORCH_C10_DRIVER_API_SUPPORTED) && !defined(USE_ROCM) if (enabled) { TORCH_WARN_ONCE("expandable_segments not supported on this platform") } diff --git a/c10/cuda/CUDACachingAllocator.cpp b/c10/cuda/CUDACachingAllocator.cpp index eb628b51968a4..c043a2c58ce1a 100644 --- a/c10/cuda/CUDACachingAllocator.cpp +++ b/c10/cuda/CUDACachingAllocator.cpp @@ -17,11 +17,17 @@ #include #include -#if !defined(USE_ROCM) && defined(PYTORCH_C10_DRIVER_API_SUPPORTED) +#if defined(PYTORCH_C10_DRIVER_API_SUPPORTED) || defined(USE_ROCM) +#if defined(PYTORCH_C10_DRIVER_API_SUPPORTED) #include +#endif +#ifndef _WIN32 #include #include #include +#else +#include +#endif #endif #include @@ -269,7 +275,8 @@ struct SegmentRange { SegmentRange(void* p, size_t s) : ptr(static_cast(p)), size(s) {} }; -#if !defined(USE_ROCM) && defined(PYTORCH_C10_DRIVER_API_SUPPORTED) +#if !defined(USE_ROCM) && defined(PYTORCH_C10_DRIVER_API_SUPPORTED) || \ + defined(USE_ROCM) /* Note [Expandable Segments] @@ -383,8 +390,13 @@ struct ExpandableSegment { // This allows for some cases where we have to unmap pages earlier in the // segment to put them at the end. max_handles_ = numSegments(prop.totalGlobalMem + prop.totalGlobalMem / 8); +#ifdef USE_ROCM + C10_CUDA_CHECK(hipMemAddressReserve( + &ptr_, segment_size_ * max_handles_, 0ULL, 0, 0ULL)); +#else C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemAddressReserve_( &ptr_, segment_size_ * max_handles_, 0ULL, 0, 0ULL)); +#endif } ExpandableSegment(const ExpandableSegment&) = delete; ExpandableSegment(ExpandableSegment&&) = delete; @@ -408,12 +420,14 @@ struct ExpandableSegment { // if it fails, use posix file handle if (CUDAAllocatorConfig::expandable_segments_handle_type() == Expandable_Segments_Handle_Type::UNSPECIFIED) { +#ifndef USE_ROCM CUDAAllocatorConfig::set_expandable_segments_handle_type( Expandable_Segments_Handle_Type::FABRIC_HANDLE); auto output = map(range); if (output.ptr != nullptr) { return output; } +#endif // if fabric handle is not supported, use posix file handle. CUDAAllocatorConfig::set_expandable_segments_handle_type( Expandable_Segments_Handle_Type::POSIX_FD); @@ -445,33 +459,60 @@ struct ExpandableSegment { if (enable_ipc_handles) { if (CUDAAllocatorConfig::expandable_segments_handle_type() != Expandable_Segments_Handle_Type::FABRIC_HANDLE) { +#ifdef USE_ROCM + prop.requestedHandleType = hipMemHandleTypePosixFileDescriptor; +#else prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR; +#endif } else { +#ifndef USE_ROCM prop.requestedHandleTypes = CU_MEM_HANDLE_TYPE_FABRIC; +#endif } } int flag = 0; +#ifndef USE_ROCM C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuDeviceGetAttribute_( &flag, CU_DEVICE_ATTRIBUTE_GPU_DIRECT_RDMA_WITH_CUDA_VMM_SUPPORTED, device_)); +#endif if (flag) prop.allocFlags.gpuDirectRDMACapable = 1; prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE; // NOLINTNEXTLINE(bugprone-signed-char-misuse) prop.location.id = static_cast(device_); +#ifdef USE_ROCM + auto status = hipMemCreate(&handle, segment_size_, &prop, 0); +#else auto status = DriverAPI::get()->cuMemCreate_(&handle, segment_size_, &prop, 0); +#endif if (status != CUDA_SUCCESS) { if (status == CUDA_ERROR_OUT_OF_MEMORY) { +#ifdef USE_ROCM + // hipMemCreate above returned hipErrorOutOfMemory and treated it + // like a sticky runtime error. Which means we need to clear it. + // Unlike the corresponding CUDA Driver API. + (void)hipGetLastError(); +#endif for (auto j : c10::irange(begin, i)) { // NOLINTNEXTLINE(bugprone-unchecked-optional-access) auto h = handles_.at(j).value(); handles_.at(j) = std::nullopt; +#ifdef USE_ROCM + C10_CUDA_CHECK(hipMemRelease(h.handle)); +#else C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemRelease_(h.handle)); +#endif } trimHandles(); return rangeFromHandles(begin, begin); +#ifdef USE_ROCM + } else { + C10_CUDA_CHECK(status); + } +#else } else if ( CUDAAllocatorConfig::expandable_segments_handle_type() == Expandable_Segments_Handle_Type::FABRIC_HANDLE) { @@ -487,6 +528,7 @@ struct ExpandableSegment { } else { C10_CUDA_DRIVER_CHECK(status); } +#endif } handles_.at(i) = Handle{handle, std::nullopt}; } @@ -522,7 +564,11 @@ struct ExpandableSegment { // thereby ensuring that the handle can be correctly matched in // ipcMemHandle_to_devptr. ShareHeader header{}; +#ifdef _WIN32 + header.pid = _getpid(); +#else header.pid = getpid(); +#endif header.segment_size = segment_size_; header.num_handles = end - begin; @@ -534,8 +580,13 @@ struct ExpandableSegment { Expandable_Segments_Handle_Type::FABRIC_HANDLE) { if (!handle.shareable_handle) { int fd = 0; +#ifdef USE_ROCM + C10_CUDA_CHECK(hipMemExportToShareableHandle( + &fd, handle.handle, hipMemHandleTypePosixFileDescriptor, 0)); +#else C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemExportToShareableHandle_( &fd, handle.handle, CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR, 0)); +#endif handle.shareable_handle = fd; LOG(INFO) << "use posix fd to share expandable segments."; } @@ -546,6 +597,10 @@ struct ExpandableSegment { reinterpret_cast(&*handle.shareable_handle), sizeof(int)); } else { +#ifdef USE_ROCM + TORCH_INTERNAL_ASSERT( + false, "expandable segment with fabric handle not supported"); +#else if (!handle.shareable_handle) { CUmemFabricHandle fabric_handle; C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemExportToShareableHandle_( @@ -559,6 +614,7 @@ struct ExpandableSegment { buf.write( reinterpret_cast(&*handle.shareable_handle), sizeof(CUmemFabricHandle)); +#endif } } return rangeFromHandles(begin, end); @@ -574,14 +630,20 @@ struct ExpandableSegment { device, std::nullopt, header.segment_size, std::move(peers)); // older build setups (e.g. multiwheels) do not have this syscall, added 2020 // but the kernel on the system might still support it. +#ifndef _WIN32 #ifndef SYS_pidfd_open #define SYS_pidfd_open 434 #endif #ifndef SYS_pidfd_getfd #define SYS_pidfd_getfd 438 #endif +#endif // !_WIN32 if (CUDAAllocatorConfig::expandable_segments_handle_type() != Expandable_Segments_Handle_Type::FABRIC_HANDLE) { +#ifdef _WIN32 + TORCH_CHECK( + false, "IPC expandable segments are not supported on Windows"); +#else auto pidfd = syscall(SYS_pidfd_open, header.pid, 0); TORCH_CHECK( pidfd != -1 || errno != ENOSYS, @@ -597,9 +659,13 @@ struct ExpandableSegment { auto err = errno; close(static_cast(pidfd)); for (auto& h : segment->handles_) { +#ifdef USE_ROCM + C10_CUDA_CHECK(hipMemRelease(h.value().handle)); +#else C10_CUDA_DRIVER_CHECK( // NOLINTNEXTLINE(bugprone-unchecked-optional-access) DriverAPI::get()->cuMemRelease_(h.value().handle)); +#endif h = std::nullopt; } TORCH_CHECK( @@ -609,17 +675,33 @@ struct ExpandableSegment { TORCH_CHECK(false, "pidfd_getfd: ", c10::utils::str_error(err)); } CUmemGenericAllocationHandle handle = 0; +#ifdef USE_ROCM +#if ROCM_VERSION >= 70100 + void* myfd_handle = + reinterpret_cast(static_cast(myfd)); +#else + void* myfd_handle = (void*)(uintptr_t)&myfd; +#endif + C10_CUDA_CHECK(hipMemImportFromShareableHandle( + &handle, myfd_handle, hipMemHandleTypePosixFileDescriptor)); +#else C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemImportFromShareableHandle_( &handle, // NOLINTNEXTLINE(performance-no-int-to-ptr) (void*)(uintptr_t)myfd, CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR)); +#endif LOG(INFO) << "use posix fd to import expandable segments."; close(static_cast(myfd)); segment->handles_.emplace_back(Handle{handle, std::nullopt}); } close(static_cast(pidfd)); +#endif // !_WIN32 } else { +#ifdef USE_ROCM + TORCH_INTERNAL_ASSERT( + false, "expandable segment with fabric handle not supported"); +#else for (auto i : c10::irange(header.num_handles)) { (void)i; CUmemFabricHandle fabric_handle; @@ -634,6 +716,7 @@ struct ExpandableSegment { LOG(INFO) << "use fabric handle to import expandable segments."; segment->handles_.emplace_back(Handle{handle, std::nullopt}); } +#endif } segment->mapAndSetAccess(0, header.num_handles); return segment; @@ -669,8 +752,12 @@ struct ExpandableSegment { ~ExpandableSegment() { forEachAllocatedRange( [&](size_t begin, size_t end) { unmapHandles(begin, end); }); +#ifdef USE_ROCM + C10_CUDA_CHECK(hipMemAddressFree(ptr_, segment_size_ * max_handles_)); +#else C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemAddressFree_( ptr_, segment_size_ * max_handles_)); +#endif } private: @@ -680,12 +767,28 @@ struct ExpandableSegment { // NOLINTNEXTLINE(bugprone-signed-char-misuse) desc.location.id = static_cast(device); desc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; +#ifdef USE_ROCM + C10_CUDA_CHECK(hipMemSetAccess( + ptr() + begin * segment_size_, + (end - begin) * segment_size_, + &desc, + 1)); +#else C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemSetAccess_( ptr_ + begin * segment_size_, (end - begin) * segment_size_, &desc, 1)); +#endif } void mapAndSetAccess(size_t begin, size_t end) { for (auto i : c10::irange(begin, end)) { +#ifdef USE_ROCM + C10_CUDA_CHECK(hipMemMap( + ptr() + i * segment_size_, + segment_size_, + 0, + handles_.at(i).value().handle, + 0ULL)); +#else C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemMap_( ptr_ + i * segment_size_, segment_size_, @@ -693,6 +796,7 @@ struct ExpandableSegment { // NOLINTNEXTLINE(bugprone-unchecked-optional-access) handles_.at(i).value().handle, 0ULL)); +#endif } mapped_size_ += (end - begin) * segment_size_; setAccess(device_, begin, end); @@ -719,12 +823,22 @@ struct ExpandableSegment { // NOLINTNEXTLINE(bugprone-unchecked-optional-access) Handle h = handles_.at(i).value(); handles_.at(i) = std::nullopt; +#ifdef USE_ROCM + C10_CUDA_CHECK(hipMemUnmap(ptr() + segment_size_ * i, segment_size_)); +#else C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemUnmap_( ptr_ + segment_size_ * i, segment_size_)); +#endif if (h.shareable_handle) { +#ifndef _WIN32 close(std::get(*h.shareable_handle)); +#endif } +#ifdef USE_ROCM + C10_CUDA_CHECK(hipMemRelease(h.handle)); +#else C10_CUDA_DRIVER_CHECK(DriverAPI::get()->cuMemRelease_(h.handle)); +#endif } trimHandles(); } @@ -772,7 +886,11 @@ struct ExpandableSegment { std::optional> shareable_handle; }; struct ShareHeader { +#ifdef _WIN32 + int pid; +#else pid_t pid; +#endif size_t segment_size; size_t num_handles; }; diff --git a/test/distributed/test_cupy_as_tensor.py b/test/distributed/test_cupy_as_tensor.py index 06b978ca91236..63b290e2e8e66 100644 --- a/test/distributed/test_cupy_as_tensor.py +++ b/test/distributed/test_cupy_as_tensor.py @@ -8,7 +8,10 @@ import torch from torch.multiprocessing.reductions import reduce_tensor from torch.testing._internal.common_cuda import SM100OrLater -from torch.testing._internal.common_distributed import MultiProcContinuousTest +from torch.testing._internal.common_distributed import ( + MultiProcContinuousTest, + skip_if_rocm_multiprocess, +) from torch.testing._internal.common_utils import ( requires_cuda_p2p_access, run_tests, @@ -67,6 +70,7 @@ def _init_device(self) -> None: def device(self) -> torch.device: return torch.device(device_type, self.rank) + @skip_if_rocm_multiprocess # RuntimeError: pidfd_getfd Operation not permitted" @skip_but_pass_in_sandcastle_if( SM100OrLater, "Fails if ran in docker environment without privileged access (https://github.com/pytorch/pytorch/issues/165170)", diff --git a/test/test_cuda.py b/test/test_cuda.py index 24db075c8bf2c..b69e3ba0867d1 100644 --- a/test/test_cuda.py +++ b/test/test_cuda.py @@ -5071,6 +5071,14 @@ def cb(device, alloc, device_alloc, device_free): def test_allocator_fuzz(self): # fuzz + if ( + torch.version.hip + and "expandable_segments:True" + in torch._C._accelerator_getAllocatorSettings() + ): + raise unittest.SkipTest( + "ROCm needs https://github.com/ROCm/rocm-systems/pull/3023" + ) state = random.getstate() random.seed(123) N = 10000 @@ -6710,7 +6718,6 @@ def test_graph_capture_pre_capture_stream_use(self): "graph_capture_record_stream_reuse:False" ) - @skipIfRocm(msg="expandable_segments mode is not supported on ROCm") @unittest.skipIf(IS_FBCODE or IS_SANDCASTLE, "Load_inline doesn't work in fbcode") def test_mempool_expandable(self): torch.cuda.empty_cache() diff --git a/test/test_cuda_expandable_segments.py b/test/test_cuda_expandable_segments.py index 262c53ab23ca4..25c2e9eaff2c5 100644 --- a/test/test_cuda_expandable_segments.py +++ b/test/test_cuda_expandable_segments.py @@ -12,7 +12,7 @@ import torch from torch.testing._internal.common_cuda import IS_JETSON, IS_WINDOWS -from torch.testing._internal.common_utils import run_tests, TEST_WITH_ROCM +from torch.testing._internal.common_utils import run_tests REPO_ROOT = pathlib.Path(__file__).resolve().parent.parent @@ -25,12 +25,7 @@ sys.path.remove(str(REPO_ROOT)) if __name__ == "__main__": - if ( - torch.cuda.is_available() - and not IS_JETSON - and not IS_WINDOWS - and not TEST_WITH_ROCM - ): + if torch.cuda.is_available() and not IS_JETSON and not IS_WINDOWS: get_disabled_tests(".") torch.cuda.memory._set_allocator_settings("expandable_segments:True") diff --git a/torch/_C/__init__.pyi.in b/torch/_C/__init__.pyi.in index d925d7e616d2f..0bce936f884d3 100644 --- a/torch/_C/__init__.pyi.in +++ b/torch/_C/__init__.pyi.in @@ -2610,6 +2610,7 @@ def _accelerator_getDeviceStats(device_index: _int) -> dict[str, Any]: ... def _accelerator_resetAccumulatedStats(device_index: _int) -> None: ... def _accelerator_resetPeakStats(device_index: _int) -> None: ... def _accelerator_getMemoryInfo(device_index: _int) -> tuple[_int, _int]: ... +def _accelerator_getAllocatorSettings() -> str: ... def _accelerator_setAllocatorSettings(env: str) -> None: ... class _acceleratorGraph: diff --git a/torch/_dynamo/trace_rules.py b/torch/_dynamo/trace_rules.py index c269bf26c41e5..c4a472de35945 100644 --- a/torch/_dynamo/trace_rules.py +++ b/torch/_dynamo/trace_rules.py @@ -464,6 +464,7 @@ "torch._C._accelerator_getAccelerator", "torch._C._accelerator_getDeviceIndex", "torch._C._accelerator_getStream", + "torch._C._accelerator_getAllocatorSettings", "torch._C._accelerator_setAllocatorSettings", "torch._C._accelerator_setStream", "torch._C._accelerator_synchronizeDevice", diff --git a/torch/csrc/DeviceAccelerator.cpp b/torch/csrc/DeviceAccelerator.cpp index 9281bf2608d75..9e5aa1e5eaa69 100644 --- a/torch/csrc/DeviceAccelerator.cpp +++ b/torch/csrc/DeviceAccelerator.cpp @@ -165,6 +165,10 @@ void initModule(PyObject* module) { return at::accelerator::getMemoryInfo(device_index); }); + m.def("_accelerator_getAllocatorSettings", []() { + return c10::CachingAllocator::getAllocatorSettings(); + }); + m.def("_accelerator_setAllocatorSettings", [](std::string env) { c10::CachingAllocator::setAllocatorSettings(env); }); From 93d4f56b5ed1f72af855d0197134cd4b0ced0363 Mon Sep 17 00:00:00 2001 From: Ankita George Date: Mon, 23 Mar 2026 16:18:10 +0000 Subject: [PATCH 0014/1172] [dcp][oss] Fix Metadata.storage_meta regression from dataclasses.replace() (#178001) (#178001) Summary: `_FileSystemWriter.finish()` used `dataclasses.replace(metadata, version=CURRENT_DCP_VERSION)` which creates a new `Metadata` object and rebinds the local variable. Since `finish()` returns `None`, the caller (`_save_state_dict`) still holds the original object, which never gets `storage_meta`, `storage_data`, or `version` updated. The on-disk metadata is correct (the new copy is pickled), but the in-memory `Metadata` returned by `dcp.save()` has `None` for these fields. Fix by mutating the passed-in metadata in-place (`metadata.version = CURRENT_DCP_VERSION`) instead of creating a new copy. `Metadata` is not a frozen dataclass, so in-place mutation is safe and consistent with the subsequent `storage_data` and `storage_meta` assignments. Fixes https://github.com/pytorch/pytorch/issues/177887 Test Plan: ``` import tempfile import torch import torch.distributed.checkpoint as dcp from torch.distributed.checkpoint._fsspec_filesystem import FsspecWriter state_dict = {"weight": torch.randn(4, 4)} with tempfile.TemporaryDirectory() as tmpdir: writer = FsspecWriter(tmpdir) meta = dcp.save(state_dict, storage_writer=writer, no_dist=True) assert meta.storage_meta is not None assert meta.storage_data is not None assert meta.version is not None print(f"storage_meta: {meta.storage_meta}") print(f"version: {meta.version}") ``` Before the fix, `meta.storage_meta` is `None`. After the fix, it contains the expected `StorageMeta` with `checkpoint_id` and `save_id`. Differential Revision: D97535692 Pull Request resolved: https://github.com/pytorch/pytorch/pull/178001 Approved by: https://github.com/LucasLLC, https://github.com/Skylion007 --- torch/distributed/checkpoint/filesystem.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/torch/distributed/checkpoint/filesystem.py b/torch/distributed/checkpoint/filesystem.py index d693790356a9d..831fa16394563 100644 --- a/torch/distributed/checkpoint/filesystem.py +++ b/torch/distributed/checkpoint/filesystem.py @@ -756,7 +756,7 @@ def _write_data( return fut def finish(self, metadata: Metadata, results: list[list[WriteResult]]) -> None: - metadata = dataclasses.replace(metadata, version=CURRENT_DCP_VERSION) + metadata.version = CURRENT_DCP_VERSION storage_md = {} for wr_list in results: From 8ae229dce08e03318c75ded9e9564723c43df9d5 Mon Sep 17 00:00:00 2001 From: zpcore Date: Mon, 23 Mar 2026 16:56:48 +0000 Subject: [PATCH 0015/1172] [DTensor] Fix `None IValue == DTensorSpec` and cache key collision for ops with optional Tensor (#177441) Two bugs in NativeShardingPropagatorCache: 1. IValueOrDTensorSpec::operator== returned true when comparing a None IValue against a DTensorSpec, because both have default iv=None. On hash bucket collisions this caused false-positive equality, leading to wrong OutputSharding results. Fixed by checking !rhs.dtensor_spec before comparing iv fields. 2. handle_non_dtensor_arg dropped None/undefined tensor args from the cache key when idx < static_argnum. For ops with optional Tensor? parameters at those positions, this made op(t1, t2, None) and op(t1, None, t2) produce identical keys. Fixed by always including None/undefined args regardless of static_argnum. Note that Bug 1 was the root cause of the flaky test_dtensor_op_db_clamp_cpu_float32 failure (hash bucket collisions between different samples triggered the broken operator==). Bug 2 is a latent issue for any op registered with high static_argnum and optional Tensor? parameters. Resolves https://github.com/pytorch/pytorch/issues/176974 Authored with Claude. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177441 Approved by: https://github.com/wconstab --- test/distributed/tensor/test_tensor_ops.py | 71 ++++++++++++++++++++++ torch/csrc/autograd/python_variable.cpp | 25 +++----- 2 files changed, 79 insertions(+), 17 deletions(-) diff --git a/test/distributed/tensor/test_tensor_ops.py b/test/distributed/tensor/test_tensor_ops.py index ca589e53c5187..d9a5799182e18 100644 --- a/test/distributed/tensor/test_tensor_ops.py +++ b/test/distributed/tensor/test_tensor_ops.py @@ -1297,6 +1297,18 @@ def build_device_mesh(self): ) +@torch.library.custom_op("testlib::optional_clamp_op", mutates_args=()) +def optional_clamp_op( + x: torch.Tensor, lo: torch.Tensor | None, hi: torch.Tensor | None +) -> torch.Tensor: + return torch.clamp(x, lo, hi) + + +@optional_clamp_op.register_fake +def _optional_clamp_op_fake(x, lo, hi): + return torch.empty_like(x) + + @torch.library.custom_op("testlib::modified_cat_op", mutates_args=()) def modified_cat_op(a: list[torch.Tensor], b: list[torch.Tensor]) -> torch.Tensor: return torch.cat(a, dim=0) @@ -1400,6 +1412,65 @@ def test_two_list_op_cache_collision(self): self.assertEqual(misses, 2) self.assertEqual(result.shape, torch.Size([8, 8])) + def test_optional_tensor_cache_key(self): + """Cache must distinguish op(t1, None, t2) from op(t1, t2, None). + + Uses a custom op with high static_argnum so that None args for + optional Tensor? parameters would be dropped from the cache key + without the is_none_or_undefined fix in handle_non_dtensor_arg. + With same-spec DTensors the collision is deterministic. + """ + from test_op_strategy import op_strategy_context + + from torch.distributed.tensor._op_schema import RuntimeSchemaInfo + from torch.distributed.tensor._ops.utils import replicate_op_strategy + from torch.distributed.tensor.debug import ( + _clear_fast_path_sharding_prop_cache, + _get_fast_path_sharding_prop_cache_stats, + ) + + mesh = self.build_device_mesh() + op = torch.ops.testlib.optional_clamp_op + + with op_strategy_context( + op.default, + replicate_op_strategy, + schema_info=RuntimeSchemaInfo(static_argnum=3), + ): + _clear_fast_path_sharding_prop_cache() + a = distribute_tensor(torch.randn(4, 8), mesh, [Shard(0)]) + + op(a, a, None) # miss 1: key should include None at position 2 + op(a, None, a) # miss 2: key should include None at position 1 + + hits, misses = _get_fast_path_sharding_prop_cache_stats() + self.assertEqual(hits, 0) + self.assertEqual(misses, 2) + + def test_optional_tensor_cache_key_python_slow_path(self): + """Same as above but exercises the Python slow path via OpSchema directly. + + DTensor_OpSchema_recompute_comparison_key_impl must include None + args for optional Tensor? parameters in the comparison key. + """ + from torch.distributed.tensor._dtensor_spec import DTensorSpec + from torch.distributed.tensor._op_schema import OpSchema, RuntimeSchemaInfo + + mesh = self.build_device_mesh() + spec = DTensorSpec(mesh, (Shard(0),)) + op = torch.ops.testlib.optional_clamp_op.default + + schema1 = OpSchema(op, (spec, None, spec), {}) + schema1.schema_info = RuntimeSchemaInfo(static_argnum=3) + schema1._recompute_comparison_key() + + schema2 = OpSchema(op, (spec, spec, None), {}) + schema2.schema_info = RuntimeSchemaInfo(static_argnum=3) + schema2._recompute_comparison_key() + + self.assertNotEqual(hash(schema1), hash(schema2)) + self.assertNotEqual(schema1, schema2) + class TestNewEmptyStridedUneven(DTensorTestBase): @with_comms diff --git a/torch/csrc/autograd/python_variable.cpp b/torch/csrc/autograd/python_variable.cpp index d220c6dec130e..b8c0633437cfd 100644 --- a/torch/csrc/autograd/python_variable.cpp +++ b/torch/csrc/autograd/python_variable.cpp @@ -1087,9 +1087,10 @@ struct IValueOrDTensorSpec { py::object dtensor_spec; bool operator==(const IValueOrDTensorSpec& rhs) const { - return dtensor_spec - ? (rhs.dtensor_spec && dtensor_spec.equal(rhs.dtensor_spec)) - : (iv == rhs.iv); + if (dtensor_spec) { + return rhs.dtensor_spec && dtensor_spec.equal(rhs.dtensor_spec); + } + return !rhs.dtensor_spec && iv == rhs.iv; } }; @@ -1162,18 +1163,6 @@ class NativeOpSchema { // have no guarantees about its lifetime. This class is cheap anyway. c10::OperatorHandle op_; std::size_t hash_; - // Subtle point: consider clamp.Tensor(Tensor self, Tensor? - // min=None, Tensor? max=None). The invocations clamp(t1, None, t2) - // and clamp(t1, t2, None) have the same comparison key (t1, t2) - // because we drop non-static non-tensor args from comparison. The - // only way we happen to be able to tell them apart is that we omit - // trailing defaulted arguments from the args tuple passed to - // __torch_dispatch__ (and hence to DTensor dispatch as well), so - // they have different args_schema_len_. - // - // I am preserving this existing behavior, but I suspect we should - // make an algorithm change to be less brittle, such as including - // None defaults for Tensor arguments in the comparison. std::size_t args_schema_len_; // There is no particular justification for the choice of 8 // here. Feel free to change it. @@ -1842,7 +1831,7 @@ static bool DTensor_OpSchema_recompute_comparison_key_impl( size_t idx = 0; for (const auto& e : args_schema) { if (idx >= native_info.static_argnum || - arg_type_tensor_or_tensor_list_like(e)) { + arg_type_tensor_or_tensor_list_like(e) || e.is_none()) { if (PyList_Check(e.ptr())) { args_to_hash.push_back( py::reinterpret_steal(PyList_AsTuple(e.ptr()))); @@ -2304,7 +2293,9 @@ create_native_op_schema( const auto handle_non_dtensor_arg = [&comparison_key, &comparison_key_hash, &native_info]( size_t idx, c10::IValue arg) { - if (idx >= native_info.static_argnum) { + bool is_none_or_undefined = + arg.isNone() || (arg.isTensor() && !arg.toTensor().defined()); + if (idx >= native_info.static_argnum || is_none_or_undefined) { if (arg.isList()) { const auto& list = arg.toList(); if (list.empty()) { From 0737a27363b31a942ed93fbf3c69bf0fc804bd39 Mon Sep 17 00:00:00 2001 From: Wei Feng Date: Sun, 22 Mar 2026 23:14:50 -0700 Subject: [PATCH 0016/1172] [FSDP2] support non-float parameters (#177948) Co-authored-by: roycho96 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177948 Approved by: https://github.com/xmfan, https://github.com/aditvenk --- .../_composable/fsdp/test_fully_shard_init.py | 108 +++++++++++++++++- .../fsdp/_fully_shard/_fsdp_param.py | 15 ++- 2 files changed, 117 insertions(+), 6 deletions(-) diff --git a/test/distributed/_composable/fsdp/test_fully_shard_init.py b/test/distributed/_composable/fsdp/test_fully_shard_init.py index b73ce10c48fe6..fe15449f3f3df 100644 --- a/test/distributed/_composable/fsdp/test_fully_shard_init.py +++ b/test/distributed/_composable/fsdp/test_fully_shard_init.py @@ -37,7 +37,14 @@ ) from torch.distributed.tensor.placement_types import _StridedShard from torch.testing._internal.common_distributed import skip_if_lt_x_gpu -from torch.testing._internal.common_fsdp import FSDPTestMultiThread, get_devtype, MLP +from torch.testing._internal.common_fsdp import ( + FSDPTest, + FSDPTestMultiThread, + get_devtype, + MLP, + patch_all_gather, + patch_reduce_scatter, +) from torch.testing._internal.common_utils import run_tests from torch.testing._internal.distributed._tensor.common_dtensor import ( ModelArgs, @@ -1393,5 +1400,104 @@ def forward(self, input): model(0) +class TestFullyShardNonFloatParam(FSDPTest): + @property + def world_size(self) -> int: + return 2 + + @skip_if_lt_x_gpu(2) + def test_non_float_param(self): + """Non-float params (e.g. uint8, int64) are supported in all-gather, + excluded from reduce-scatter, and not cast by mp_policy.param_dtype.""" + self.run_subtests( + { + "non_float_dtypes": [ + [torch.uint8], + [torch.int64], + [torch.uint8, torch.int8, torch.int16, torch.int32, torch.int64], + ], + "mp_policy": [ + MixedPrecisionPolicy(), + MixedPrecisionPolicy( + param_dtype=torch.bfloat16, reduce_dtype=torch.float32 + ), + ], + "frozen_float": [False, True], + }, + self._test_non_float_param, + ) + + def _test_non_float_param( + self, + non_float_dtypes: list, + mp_policy: MixedPrecisionPolicy, + frozen_float: bool, + ): + class Model(nn.Module): + def __init__(self): + super().__init__() + self.linear = nn.Linear(16, 16) + if frozen_float: + self.frozen_float = nn.Parameter( + torch.randn(16), requires_grad=False + ) + for i, dtype in enumerate(non_float_dtypes): + self.register_parameter( + f"non_float_{i}", + nn.Parameter( + torch.randint(0, 127, (16,), dtype=dtype), + requires_grad=False, + ), + ) + + def forward(self, x): + return self.linear(x) + + model = Model() + fully_shard(model, mp_policy=mp_policy) + for p in model.parameters(): + self.assertEqual(p.size(0) % self.world_size, 0) + ag_input_dtypes = set() + expected_ag_output_bytes = 0 + for p in model.parameters(): + if p.dtype.is_floating_point and mp_policy.param_dtype is not None: + ag_input_dtypes.add(mp_policy.param_dtype) + expected_ag_output_bytes += p.numel() * mp_policy.param_dtype.itemsize + else: + # Non-float params keep their original dtype; param_dtype + # only applies to floating-point params + ag_input_dtypes.add(p.dtype) + expected_ag_output_bytes += p.numel() * p.element_size() + expected_ag_dtype = ( + next(iter(ag_input_dtypes)) if len(ag_input_dtypes) == 1 else torch.uint8 + ) + orig_ag = dist.all_gather_into_tensor + + def assert_all_gather(*args, **kw): + output = kw.get("output", args[0] if len(args) > 0 else None) + input = kw.get("input_tensor", args[1] if len(args) > 1 else None) + self.assertEqual(input.dtype, expected_ag_dtype) + self.assertEqual(output.nbytes, expected_ag_output_bytes) + return orig_ag(*args, **kw) + + expected_rs_input_numel = sum( + p.numel() for p in model.parameters() if p.requires_grad + ) + orig_rs = dist.reduce_scatter_tensor + + def assert_reduce_scatter(*args, **kw): + input = kw.get("input", args[1] if len(args) > 1 else None) + self.assertEqual(input.numel(), expected_rs_input_numel) + return orig_rs(*args, **kw) + + x = torch.randn(4, 16, device=device_type) + with ( + patch_all_gather(assert_all_gather), + patch_reduce_scatter(assert_reduce_scatter), + ): + loss = model(x).sum() + loss.backward() + + if __name__ == "__main__": run_tests() diff --git a/torch/distributed/fsdp/_fully_shard/_fsdp_param.py b/torch/distributed/fsdp/_fully_shard/_fsdp_param.py index 8b1227ee6d92e..82f1c85b4c19c 100644 --- a/torch/distributed/fsdp/_fully_shard/_fsdp_param.py +++ b/torch/distributed/fsdp/_fully_shard/_fsdp_param.py @@ -317,8 +317,10 @@ def _init_sharded_param( raise AssertionError( f"Expected contiguous tensor with {self.fsdp_placement=}" ) - self.sharded_param = nn.Parameter(self.to_sharded_dtensor(sharded_param)) - self.sharded_param.requires_grad_(param.requires_grad) + self.sharded_param = nn.Parameter( + self.to_sharded_dtensor(sharded_param), + requires_grad=param.requires_grad, + ) # Let `param_data` be freed normally when its ref count reaches 0 when # the `fully_shard` call returns to allow provided parameters to alias self._setattr_on_modules(self.sharded_param) @@ -529,8 +531,10 @@ def init_dtype_attrs(self, mp_policy: MixedPrecisionPolicy): # then we do not need extra casting if reduce_dtype == param_dtype: reduce_dtype = None - # Clamp `param_dtype` to `None` if no casting is required - if param_dtype == self.orig_dtype: + # Clamp `param_dtype` to `None` if no casting is required or if the + # parameter is non-floating-point (mixed precision is only meaningful + # for floating-point parameters) + if param_dtype == self.orig_dtype or not self.orig_dtype.is_floating_point: param_dtype = None self.param_dtype = param_dtype self.reduce_dtype = reduce_dtype @@ -659,7 +663,8 @@ def to_sharded_post_forward(self) -> None: storage_offset=0, ) self._sharded_post_forward_param = nn.Parameter( - self.to_sharded_post_forward_dtensor(sharded_post_forward_tensor) + self.to_sharded_post_forward_dtensor(sharded_post_forward_tensor), + requires_grad=self.sharded_param.requires_grad, ) self._setattr_on_modules(self._sharded_post_forward_param) self.free_unsharded_param() From 21408e59e8a45365e2ab87ed1f0d527f67e83e06 Mon Sep 17 00:00:00 2001 From: Frank Lin Date: Thu, 19 Mar 2026 21:24:20 +0000 Subject: [PATCH 0017/1172] [CUDA Graph] Add stream capture ID query utilities (#176751) Add helper APIs for querying CUDA stream capture IDs, in preparation for the upcoming per-capture RNG state refactor. This consolidates direct `cudaStreamGetCaptureInfo` / `cudaStreamIsCapturing` usage behind a small, cleaner interface. * Add `c10::cuda` helpers: * `currentStreamCaptureIdMayInitCtx()` * `isStreamCapturingMayInitCtx()` * extend `currentStreamCaptureStatusMayInitCtx()` with an optional stream argument * Add context-safe `at::cuda` wrappers: * `currentStreamCaptureId()` * `isStreamCapturing()` * Migrate existing callsites in: * `CUDAGraph` (`capture_begin`, `get_currently_capturing_graph`, `create_allocate_filter`, `create_child_allocate_filter`, `begin_capture_to_if_node`) * `CachingHostAllocator` * cuDNN RNN `DropoutState` Pull Request resolved: https://github.com/pytorch/pytorch/pull/176751 Approved by: https://github.com/ngimel --- aten/src/ATen/cuda/CUDAGraph.cpp | 78 ++++++++++----------- aten/src/ATen/cuda/CUDAGraph.h | 9 ++- aten/src/ATen/cuda/CUDAGraphsUtils.cuh | 11 ++- aten/src/ATen/cuda/CachingHostAllocator.cpp | 6 +- aten/src/ATen/native/cudnn/RNN.cpp | 14 +--- c10/cuda/CUDAGraphsC10Utils.h | 41 +++++++++-- 6 files changed, 96 insertions(+), 63 deletions(-) diff --git a/aten/src/ATen/cuda/CUDAGraph.cpp b/aten/src/ATen/cuda/CUDAGraph.cpp index f1e25cfe011a6..0f4e16b9fd5b3 100644 --- a/aten/src/ATen/cuda/CUDAGraph.cpp +++ b/aten/src/ATen/cuda/CUDAGraph.cpp @@ -2,6 +2,7 @@ #include #include #include +#include #include #include #include @@ -9,6 +10,7 @@ #include #include +#include namespace at::cuda { @@ -78,6 +80,23 @@ void CUDAGraph::register_generator_state(const at::Generator& generator) { cuda_gen->register_graph(this); } +template <> +std::function CUDAGraph::create_allocate_filter() const { + return [this](cudaStream_t stream) { + auto capture_id_opt = c10::cuda::captureIdMayInitCtx(stream); + return capture_id_opt.has_value() && capture_id_opt.value() == capture_id_; + }; +} + +template <> +std::function CUDAGraph::create_allocate_filter() const { + return [this](c10::Stream stream) { + cudaStream_t cuda_stream = CUDAStream(CUDAStream::UNCHECKED, stream); + auto capture_id_opt = c10::cuda::captureIdMayInitCtx(cuda_stream); + return capture_id_opt.has_value() && capture_id_opt.value() == capture_id_; + }; +} + void CUDAGraph::capture_begin(MempoolId_t pool/*={0,0}*/, cudaStreamCaptureMode capture_mode) { TORCH_CHECK(!has_graph_exec_, "This CUDAGraph instance already owns a captured graph. " @@ -122,25 +141,22 @@ void CUDAGraph::capture_begin(MempoolId_t pool/*={0,0}*/, cudaStreamCaptureMode // Addendum: beginAllocateStreamToPool is now called before cudaStreamBeginCapture to prevent an // autograd thread's free() call triggering an invalid cudaEventRecord in the caching allocator // due to the capture status being updated _after_ a capture had already started. - c10::cuda::CUDACachingAllocator::beginAllocateToPool(capture_dev_, mempool_id_, create_allocate_filter()); + c10::cuda::CUDACachingAllocator::beginAllocateToPool(capture_dev_, mempool_id_, create_allocate_filter()); - auto filter = create_allocate_filter(); - - at::getHostAllocator(at::kCUDA)->begin_allocate_to_pool(mempool_id_, [filter](c10::Stream stream) { - return filter(CUDAStream(CUDAStream::UNCHECKED, stream)); - }); + at::getHostAllocator(at::kCUDA)->begin_allocate_to_pool(mempool_id_, create_allocate_filter()); // cudaStreamCaptureModeGlobal is the most conservative option to // prevent potentially unsafe CUDA API calls during capture. See // https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__STREAM.html#group__CUDART__STREAM_1g9d0535d93a214cbf126835257b16ba85 AT_CUDA_CHECK(cudaStreamBeginCapture(capture_stream_, capture_mode)); - cudaStreamCaptureStatus status{}; - AT_CUDA_CHECK(cudaStreamGetCaptureInfo(stream, &status, &capture_id_)); - TORCH_INTERNAL_ASSERT(status == cudaStreamCaptureStatus::cudaStreamCaptureStatusActive); + auto capture_id_opt = c10::cuda::captureIdMayInitCtx(stream); + TORCH_INTERNAL_ASSERT(capture_id_opt.has_value(), + "Stream should be actively capturing after cudaStreamBeginCapture"); + capture_id_ = capture_id_opt.value(); { - std::unique_lock lock(_currently_capturing_graphs_mutex); + std::lock_guard lock(_currently_capturing_graphs_mutex); _currently_capturing_graphs.emplace(capture_id_, this); } } @@ -343,17 +359,14 @@ CUDAGraph::~CUDAGraph() { CUDAGraph* CUDAGraph::get_currently_capturing_graph() { std::unique_lock lock(_currently_capturing_graphs_mutex); - cudaStreamCaptureStatus status{}; - CaptureId_t current_capture_id = 0; - auto stream = at::cuda::getCurrentCUDAStream(); - AT_CUDA_CHECK(cudaStreamGetCaptureInfo(stream, &status, ¤t_capture_id)); + auto capture_id_opt = c10::cuda::currentStreamCaptureIdMayInitCtx(); TORCH_CHECK( - status == cudaStreamCaptureStatus::cudaStreamCaptureStatusActive, + capture_id_opt.has_value(), "The current stream is not currently capturing."); TORCH_CHECK( - _currently_capturing_graphs.count(current_capture_id), + _currently_capturing_graphs.count(capture_id_opt.value()), "get_currently_capturing_graph() can be used only between capture_begin() and capture_end(). Did you use a stream without making it depend upon the original stream used for capture?"); - return _currently_capturing_graphs.at(current_capture_id); + return _currently_capturing_graphs.at(capture_id_opt.value()); } void CUDAGraph::begin_capture_to_if_node( @@ -459,10 +472,10 @@ getCurrentCUDAStream(), &cond_node, nullptr, 1, cudaStreamSetCaptureDependencies AT_CUDA_CHECK(cudaStreamBeginCaptureToGraph( child_stream, if_node_child_graph, nullptr, nullptr, 0, capture_mode_)); - AT_CUDA_CHECK(cudaStreamGetCaptureInfo( - child_stream, &status, &conditional_graph_capture_ids_.top())); - TORCH_INTERNAL_ASSERT( - status == cudaStreamCaptureStatus::cudaStreamCaptureStatusActive); + auto child_capture_id_opt = c10::cuda::captureIdMayInitCtx(child_stream); + TORCH_INTERNAL_ASSERT(child_capture_id_opt.has_value(), + "Child stream should be actively capturing after cudaStreamBeginCaptureToGraph"); + conditional_graph_capture_ids_.top() = child_capture_id_opt.value(); conditional_node_streams_.emplace(child_stream); @@ -522,11 +535,8 @@ void CUDAGraph::end_capture_to_conditional_node() { at::getHostAllocator(at::kCUDA)->end_allocate_to_pool(mempool_id_); if (conditional_graph_capture_ids_.empty()) { c10::cuda::CUDACachingAllocator::beginAllocateToPool( - capture_dev_, mempool_id_, create_allocate_filter()); - auto filter = create_allocate_filter(); - at::getHostAllocator(at::kCUDA)->begin_allocate_to_pool(mempool_id_, [filter](c10::Stream stream) { - return filter(CUDAStream(CUDAStream::UNCHECKED, stream)); - }); + capture_dev_, mempool_id_, create_allocate_filter()); + at::getHostAllocator(at::kCUDA)->begin_allocate_to_pool(mempool_id_, create_allocate_filter()); } else { c10::cuda::CUDACachingAllocator::beginAllocateToPool( capture_dev_, mempool_id_, create_child_allocate_filter()); @@ -535,7 +545,6 @@ void CUDAGraph::end_capture_to_conditional_node() { return filter(CUDAStream(CUDAStream::UNCHECKED, stream)); }); } - constexpr const char* rng_with_conditional_nodes_error = "RNG within data-dependent conditional nodes is not supported yet."; TORCH_CHECK(!rng_or_generators_changed, rng_with_conditional_nodes_error); @@ -547,22 +556,11 @@ void CUDAGraph::end_capture_to_conditional_node() { #endif } -std::function CUDAGraph::create_allocate_filter() { - return [this](cudaStream_t stream) { - cudaStreamCaptureStatus status{}; - CaptureId_t stream_capture_id = 0; - AT_CUDA_CHECK(cudaStreamGetCaptureInfo(stream, &status, &stream_capture_id)); - return status == cudaStreamCaptureStatus::cudaStreamCaptureStatusActive && stream_capture_id == capture_id_; - }; -} - std::function CUDAGraph::create_child_allocate_filter() { #if !defined(USE_ROCM) && (defined(CUDA_VERSION) && CUDA_VERSION >= 12040) return [¤t_capture_id = conditional_graph_capture_ids_.top()](cudaStream_t stream) { - cudaStreamCaptureStatus status{}; - CaptureId_t stream_capture_id{}; - AT_CUDA_CHECK(cudaStreamGetCaptureInfo(stream, &status, &stream_capture_id)); - return status == cudaStreamCaptureStatus::cudaStreamCaptureStatusActive && stream_capture_id == current_capture_id; + auto capture_id_opt = c10::cuda::captureIdMayInitCtx(stream); + return capture_id_opt.has_value() && capture_id_opt.value() == current_capture_id; }; #else // !defined(USE_ROCM) && (defined(CUDA_VERSION) && CUDA_VERSION >= 12040) AT_ERROR( diff --git a/aten/src/ATen/cuda/CUDAGraph.h b/aten/src/ATen/cuda/CUDAGraph.h index edaad31b42bb5..d1dc6e8740f1c 100644 --- a/aten/src/ATen/cuda/CUDAGraph.h +++ b/aten/src/ATen/cuda/CUDAGraph.h @@ -9,6 +9,7 @@ #include #include +#include #include #if defined(USE_ROCM) || !(defined(CUDA_VERSION) && CUDA_VERSION >= 12040) @@ -88,7 +89,8 @@ struct TORCH_CUDA_CPP_API CUDAGraph { const Tensor& scalar_cuda_pred_tensor); private: - std::function create_allocate_filter(); + template + std::function create_allocate_filter() const; std::function create_child_allocate_filter(); protected: @@ -152,5 +154,10 @@ struct TORCH_CUDA_CPP_API CUDAGraph { #endif // !defined(USE_ROCM) && defined(CUDA_VERSION) && CUDA_VERSION >= 12040 }; +template <> +std::function CUDAGraph::create_allocate_filter() const; +template <> +std::function CUDAGraph::create_allocate_filter() const; + } // namespace cuda } // namespace at diff --git a/aten/src/ATen/cuda/CUDAGraphsUtils.cuh b/aten/src/ATen/cuda/CUDAGraphsUtils.cuh index d3a5b306eeea4..64ab0140793bb 100644 --- a/aten/src/ATen/cuda/CUDAGraphsUtils.cuh +++ b/aten/src/ATen/cuda/CUDAGraphsUtils.cuh @@ -19,12 +19,17 @@ using CaptureStatus = c10::cuda::CaptureStatus; // Use this version where you don't want to create a CUDA context if none exists. inline CaptureStatus currentStreamCaptureStatus() { - // don't create a context if we don't have to if (c10::cuda::hasPrimaryContext(c10::cuda::current_device())) { return c10::cuda::currentStreamCaptureStatusMayInitCtx(); - } else { - return CaptureStatus::None; } + return CaptureStatus::None; +} + +inline std::optional currentStreamCaptureId() { + if (c10::cuda::hasPrimaryContext(c10::cuda::current_device())) { + return c10::cuda::currentStreamCaptureIdMayInitCtx(); + } + return std::nullopt; } inline void assertNotCapturing(const std::string& attempt) { diff --git a/aten/src/ATen/cuda/CachingHostAllocator.cpp b/aten/src/ATen/cuda/CachingHostAllocator.cpp index dcff81333cb2f..075cdda3ba8ae 100644 --- a/aten/src/ATen/cuda/CachingHostAllocator.cpp +++ b/aten/src/ATen/cuda/CachingHostAllocator.cpp @@ -3,6 +3,7 @@ #include #include #include +#include #include #include @@ -245,9 +246,8 @@ struct CUDACachingHostAllocatorImpl } bool stream_is_capturing(CUDAStream s) const override { - cudaStreamCaptureStatus status{cudaStreamCaptureStatusNone}; - C10_CUDA_CHECK(cudaStreamIsCapturing(s, &status)); - return status != cudaStreamCaptureStatusNone; + return c10::cuda::isStreamCapturingMayInitCtx( + static_cast(s)); } }; diff --git a/aten/src/ATen/native/cudnn/RNN.cpp b/aten/src/ATen/native/cudnn/RNN.cpp index 9c12ad5f96b3e..a45fda2b22db4 100644 --- a/aten/src/ATen/native/cudnn/RNN.cpp +++ b/aten/src/ATen/native/cudnn/RNN.cpp @@ -2376,12 +2376,7 @@ struct DropoutState { if (event) { #if !defined(USE_ROCM) // See Note [DropoutState and CUDA graph capture] - cudaStreamCaptureStatus status; - AT_CUDA_CHECK(cudaStreamGetCaptureInfo( - cuda::getCurrentCUDAStream(), &status, &capture_id_last_lock)); - if (status == cudaStreamCaptureStatus::cudaStreamCaptureStatusNone) { - capture_id_last_lock = 0; - } + capture_id_last_lock = at::cuda::currentStreamCaptureId().value_or(0); if (capture_id_last_lock == capture_id_last_unlock) { event->block(cuda::getCurrentCUDAStream()); } @@ -2396,12 +2391,7 @@ struct DropoutState { event->record(); #if !defined(USE_ROCM) // See Note [DropoutState and CUDA graph capture] - cudaStreamCaptureStatus status; - AT_CUDA_CHECK(cudaStreamGetCaptureInfo( - cuda::getCurrentCUDAStream(), &status, &capture_id_last_unlock)); - if (status == cudaStreamCaptureStatus::cudaStreamCaptureStatusNone) { - capture_id_last_unlock = 0; - } + capture_id_last_unlock = at::cuda::currentStreamCaptureId().value_or(0); TORCH_INTERNAL_ASSERT(capture_id_last_unlock == capture_id_last_lock); #endif } diff --git a/c10/cuda/CUDAGraphsC10Utils.h b/c10/cuda/CUDAGraphsC10Utils.h index 936875fd71d5c..3709aa64c7845 100644 --- a/c10/cuda/CUDAGraphsC10Utils.h +++ b/c10/cuda/CUDAGraphsC10Utils.h @@ -1,8 +1,10 @@ #pragma once +#include #include + #include -#include +#include // CUDA Graphs utils used by c10 and aten. // aten/cuda/CUDAGraphsUtils.cuh adds utils used by aten only. @@ -67,10 +69,41 @@ inline std::ostream& operator<<(std::ostream& os, CaptureStatus status) { // Use this version where you're sure a CUDA context exists already. inline CaptureStatus currentStreamCaptureStatusMayInitCtx() { - cudaStreamCaptureStatus is_capturing{cudaStreamCaptureStatusNone}; + cudaStreamCaptureStatus status{cudaStreamCaptureStatusNone}; C10_CUDA_CHECK( - cudaStreamIsCapturing(c10::cuda::getCurrentCUDAStream(), &is_capturing)); - return CaptureStatus(is_capturing); + cudaStreamIsCapturing(c10::cuda::getCurrentCUDAStream(), &status)); + return CaptureStatus(status); +} + +inline CaptureStatus captureStatusMayInitCtx(cudaStream_t stream) { + cudaStreamCaptureStatus status{cudaStreamCaptureStatusNone}; + C10_CUDA_CHECK(cudaStreamIsCapturing(stream, &status)); + return CaptureStatus(status); +} + +inline bool isStreamCapturingMayInitCtx(cudaStream_t stream) { + return captureStatusMayInitCtx(stream) == CaptureStatus::Active; +} + +inline std::optional currentStreamCaptureIdMayInitCtx() { + cudaStreamCaptureStatus status{}; + CaptureId_t capture_id = 0; + C10_CUDA_CHECK(cudaStreamGetCaptureInfo( + c10::cuda::getCurrentCUDAStream(), &status, &capture_id)); + if (status == cudaStreamCaptureStatus::cudaStreamCaptureStatusActive) { + return capture_id; + } + return std::nullopt; +} + +inline std::optional captureIdMayInitCtx(cudaStream_t stream) { + cudaStreamCaptureStatus status{}; + CaptureId_t capture_id = 0; + C10_CUDA_CHECK(cudaStreamGetCaptureInfo(stream, &status, &capture_id)); + if (status == cudaStreamCaptureStatus::cudaStreamCaptureStatusActive) { + return capture_id; + } + return std::nullopt; } } // namespace c10::cuda From 37664e83b8ec2e5578c076940ed2c839d60587ee Mon Sep 17 00:00:00 2001 From: Benedikt Johannes Date: Mon, 23 Mar 2026 17:21:31 +0000 Subject: [PATCH 0018/1172] [C10] Fix incorrect channels-last contiguity flag handling in `c10/core/SymbolicShapeMeta.h` (#178103) # Summary This pull request fixes #178102 by correcting a bug in `c10/core/SymbolicShapeMeta.h` where `assume_channels_last_contiguous` that consistet of an incorrect change `is_contiguous_` instead of `is_channels_last_contiguous_`. This caused inconsistent state and may have leaded to wrong contiguity information for channels-last tensors. # Change(s) - Replaced assignment `is_contiguous_ = std::move(val);` with `is_channels_last_contiguous_ = std::move(val);` # Impact - Corrects the internal state of `SymbolicShapeMeta` when a tensor is assumed to be channels-last contiguous. - Prevents potential incorrect optimizations and ensures proper propagation of contiguity flags in symbolic shape handling. Contributed by **Benedikt Johannes** Pull Request resolved: https://github.com/pytorch/pytorch/pull/178103 Approved by: https://github.com/Skylion007 --- c10/core/SymbolicShapeMeta.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/c10/core/SymbolicShapeMeta.h b/c10/core/SymbolicShapeMeta.h index 0820038968a8e..59c3147c878f7 100644 --- a/c10/core/SymbolicShapeMeta.h +++ b/c10/core/SymbolicShapeMeta.h @@ -141,7 +141,7 @@ class C10_API SymbolicShapeMeta { available_.fetch_or(is_contiguous_avail); } void assume_channels_last_contiguous(SymBool val = true) { - is_contiguous_ = std::move(val); + is_channels_last_contiguous_ = std::move(val); available_.fetch_or(is_channels_last_contiguous_avail); } void assume_channels_last_3d_contiguous(SymBool val = true) { From f3ae965591b96301c638139b4206b280453f08d5 Mon Sep 17 00:00:00 2001 From: Aaron Orenstein Date: Sun, 22 Mar 2026 19:04:44 -0700 Subject: [PATCH 0019/1172] [aotautograd] Clean up AOT autograd subclass handling (#178107) Tighten type annotations across AOT autograd's subclass utilities to eliminate pyrefly suppressions and type: ignore comments. The main structural change is in flatten_subclass, where the SubclassCreationMeta assertion is moved before the attrs access that already required it (previously suppressed with a pyrefly ignore). The rest is mechanical: replacing OpaqueType with OpaqueBase, using the public DTensor import path, and casting args to typed locals instead of using type: ignore. Authored with Claude. Pull Request resolved: https://github.com/pytorch/pytorch/pull/178107 Approved by: https://github.com/bobrenjc93, https://github.com/Lucaskabela --- torch/_dynamo/variables/torch.py | 4 +- .../collect_metadata_analysis.py | 16 ++++--- .../_aot_autograd/graph_capture_wrappers.py | 46 +++++++++++++------ .../_aot_autograd/runtime_wrappers.py | 4 +- .../_aot_autograd/subclass_utils.py | 40 ++++++++-------- torch/_subclasses/fake_tensor.py | 5 +- 6 files changed, 67 insertions(+), 48 deletions(-) diff --git a/torch/_dynamo/variables/torch.py b/torch/_dynamo/variables/torch.py index 047ef88d7e4f5..83193c9717c97 100644 --- a/torch/_dynamo/variables/torch.py +++ b/torch/_dynamo/variables/torch.py @@ -114,7 +114,7 @@ if TYPE_CHECKING: from torch._dynamo.symbolic_convert import InstructionTranslator - from torch._library.opaque_object import OpaqueType + from torch._opaque_base import OpaqueBase from torch.utils._pytree import TreeSpec @@ -285,7 +285,7 @@ def _collect_all_grad_fns(tensor: torch.Tensor) -> set[torch.autograd.graph.Node grad_fns: set[torch.autograd.graph.Node] = set() - plain_tensors: list[torch.SymInt | torch.Tensor | int | OpaqueType] = [] + plain_tensors: list[torch.SymInt | torch.Tensor | int | OpaqueBase] = [] # Get all plain tensors (handles nested subclasses) if is_traceable_wrapper_subclass(tensor): get_plain_tensors(tensor, out=plain_tensors) diff --git a/torch/_functorch/_aot_autograd/collect_metadata_analysis.py b/torch/_functorch/_aot_autograd/collect_metadata_analysis.py index f6edf2f4036a4..d52f55d5b8b33 100644 --- a/torch/_functorch/_aot_autograd/collect_metadata_analysis.py +++ b/torch/_functorch/_aot_autograd/collect_metadata_analysis.py @@ -870,6 +870,12 @@ def _is_subclass_mutated_input_tangent_always_subclass(inp: object) -> bool: grad_enabled_mutation, ) + subclass_inp_meta = create_subclass_meta(flat_args) + subclass_fw_graph_out_meta = create_subclass_meta(fw_graph_outs) + subclass_tangent_meta = create_subclass_meta( + traced_tangents, count_symints=False, with_memory_format=True + ) + metadata = ViewAndMutationMeta( input_info=input_info, output_info=output_info, @@ -877,13 +883,9 @@ def _is_subclass_mutated_input_tangent_always_subclass(inp: object) -> bool: keep_input_mutations=keep_input_mutations, traced_tangents=traced_tangents, traced_tangents_descs=traced_tangents_descs, - subclass_inp_meta=create_subclass_meta(flat_args), - subclass_fw_graph_out_meta=create_subclass_meta(fw_graph_outs), - subclass_tangent_meta=create_subclass_meta( - traced_tangents, - count_symints=False, - with_memory_format=True, - ), + subclass_inp_meta=subclass_inp_meta, + subclass_fw_graph_out_meta=subclass_fw_graph_out_meta, + subclass_tangent_meta=subclass_tangent_meta, is_train=is_train, grad_enabled_mutation=grad_enabled_mutation, static_input_indices=static_input_indices, diff --git a/torch/_functorch/_aot_autograd/graph_capture_wrappers.py b/torch/_functorch/_aot_autograd/graph_capture_wrappers.py index 71a92a7f0d647..f596aafdd3ea1 100644 --- a/torch/_functorch/_aot_autograd/graph_capture_wrappers.py +++ b/torch/_functorch/_aot_autograd/graph_capture_wrappers.py @@ -10,6 +10,7 @@ 4. dispatching subclasses """ +import typing import warnings from collections.abc import Callable, Generator from contextlib import AbstractContextManager, contextmanager, ExitStack, nullcontext @@ -1356,11 +1357,15 @@ def inner_fn(fn: Callable[..., Any], args: Any, *, use_trace_joint: bool) -> Any # Add extra symints as outputs to the forward/backward graphs # ignore nested ints here forward_outs, forward_outs_descs = unwrap_tensor_subclasses( - wrapped_outs[0], wrapped_outs_descs[0], append_symints=True + wrapped_outs[0], + wrapped_outs_descs[0], + append_symints=True, ) # ignore nested ints here backward_outs, backward_outs_descs = unwrap_tensor_subclasses( - wrapped_outs[1], wrapped_outs_descs[1], append_symints=True + wrapped_outs[1], + wrapped_outs_descs[1], + append_symints=True, ) return ( (forward_outs, backward_outs), @@ -1394,42 +1399,53 @@ def inner_fw_only(*args: Any) -> Any: return inner_fn(inner_fw_only, primals, use_trace_joint=False) if is_joint_structure: + primals_wrapped: list[FxValue] = typing.cast(list[FxValue], args[0]) + primals_wrapped_descs: list[AOTInput] = typing.cast( + list[AOTInput], args_descs[0] + ) + tangents_wrapped: list[FxValue] = typing.cast(list[FxValue], args[1]) + tangents_wrapped_descs: list[AOTInput] = typing.cast( + list[AOTInput], args_descs[1] + ) + # Add extra symints (size/strides) as input to the forward graph primals_unwrapped_pair = unwrap_tensor_subclasses( - args[0], # type: ignore[arg-type] - args_descs[0], # type: ignore[arg-type] + primals_wrapped, + primals_wrapped_descs, append_symints=True, ) # We pass append_symints=False here because the partitioner will # capture and add any extra argument. tangents_unwrapped_pair = unwrap_tensor_subclasses( - args[1], # type: ignore[arg-type] - args_descs[1], # type: ignore[arg-type] + tangents_wrapped, + tangents_wrapped_descs, append_symints=False, ) args_unwrapped = (primals_unwrapped_pair[0], tangents_unwrapped_pair[0]) args_descs_unwrapped = (primals_unwrapped_pair[1], tangents_unwrapped_pair[1]) remapped_static_indices = remap_unwrapped_subclass_arg_indices( - args[0], # type: ignore[arg-type] + primals_wrapped, meta.static_input_indices, # type: ignore[arg-type] ) + + primals_unwrapped = args_unwrapped[0] # type: ignore[assignment] + primals_unwrapped_descs = args_descs_unwrapped[0] # type: ignore[assignment] + fn_to_trace = joint_fn # type: ignore[assignment] else: + primals_wrapped: list[FxValue] = typing.cast(list[FxValue], args) + primals_wrapped_descs: list[AOTInput] = typing.cast(list[AOTInput], args_descs) + args_unwrapped, args_descs_unwrapped = unwrap_tensor_subclasses( # type: ignore[assignment] - args, # type: ignore[arg-type] - args_descs, # type: ignore[arg-type] + primals_wrapped, + primals_wrapped_descs, append_symints=True, ) remapped_static_indices = remap_unwrapped_subclass_arg_indices( - args, # type: ignore[arg-type] + primals_wrapped, meta.static_input_indices, # type: ignore[arg-type] ) - if is_joint_structure: - primals_unwrapped = args_unwrapped[0] # type: ignore[assignment] - primals_unwrapped_descs = args_descs_unwrapped[0] # type: ignore[assignment] - fn_to_trace = joint_fn # type: ignore[assignment] - else: primals_unwrapped = args_unwrapped # type: ignore[assignment] primals_unwrapped_descs = args_descs_unwrapped # type: ignore[assignment] fn_to_trace = fw_fn # type: ignore[assignment] diff --git a/torch/_functorch/_aot_autograd/runtime_wrappers.py b/torch/_functorch/_aot_autograd/runtime_wrappers.py index 41eb18c504b6a..ee158d2965e2b 100644 --- a/torch/_functorch/_aot_autograd/runtime_wrappers.py +++ b/torch/_functorch/_aot_autograd/runtime_wrappers.py @@ -114,12 +114,12 @@ def _describe_arg_for_logging(arg: object) -> str: from torch._library import opaque_object try: - is_dtensor = isinstance(arg, torch.distributed._tensor.DTensor) + is_dtensor = isinstance(arg, torch.distributed.tensor.DTensor) except AttributeError: is_dtensor = False if is_dtensor: - arg = typing.cast(torch.distributed._tensor.DTensor, arg) + arg = typing.cast(torch.distributed.tensor.DTensor, arg) mesh = arg.device_mesh return ( f"DTensor(shape={arg.shape}, dtype={arg.dtype}, " diff --git a/torch/_functorch/_aot_autograd/subclass_utils.py b/torch/_functorch/_aot_autograd/subclass_utils.py index b9548107eb6fe..c866f554550d9 100644 --- a/torch/_functorch/_aot_autograd/subclass_utils.py +++ b/torch/_functorch/_aot_autograd/subclass_utils.py @@ -7,7 +7,7 @@ import collections import typing from collections.abc import Callable, Iterable, Sequence -from typing import Any, TYPE_CHECKING, TypeGuard, TypeVar +from typing import Any, TypeGuard, TypeVar import torch import torch.utils._pytree as pytree @@ -42,10 +42,6 @@ from .utils import strict_zip -if TYPE_CHECKING: - from torch._library.opaque_object import OpaqueType - - zip = strict_zip T = TypeVar("T", bound=torch.Tensor) @@ -328,16 +324,21 @@ def runtime_unwrap_tensor_subclasses( *, append_symints: bool, subclass_metas: list[PlainTensorMeta | SubclassCreationMeta] | None = None, -) -> list[Any]: +) -> list[int | Tensor | SymInt | OpaqueBase]: def flatten_subclass( - x: Tensor, meta: SubclassCreationMeta | None, *, out: list[Any] - ) -> list[Any]: + x: Tensor, + subclass_meta: PlainTensorMeta | SubclassCreationMeta | OpaqueMeta | None, + *, + out: list[OpaqueBase | SymInt | Tensor | int], + ) -> list[OpaqueBase | SymInt | Tensor | int]: if not is_traceable_wrapper_subclass(x): out.append(x) return out if not isinstance(x, Tensor): raise AssertionError(f"expected Tensor, got {type(x)}") + if not isinstance(subclass_meta, SubclassCreationMeta): + raise AssertionError("subclass_meta should be a SubclassCreationMeta") attrs, _ = x.__tensor_flatten__() @@ -347,8 +348,7 @@ def flatten_subclass( case OpaqueBase(): out.append(inner_value) case Tensor(): - # pyrefly: ignore [missing-attribute] - inner_meta = meta.attrs.get(attr) + inner_meta = subclass_meta.attrs.get(attr) flatten_subclass(inner_value, inner_meta, out=out) case _: raise AssertionError( @@ -356,11 +356,9 @@ def flatten_subclass( ) if append_symints: - if not isinstance(meta, SubclassCreationMeta): - raise AssertionError(f"expected SubclassCreationMeta, got {type(meta)}") # outer_size size = x.size() - symint_placeholders = compute_symint_placeholders(meta.outer_size) + symint_placeholders = compute_symint_placeholders(subclass_meta.outer_size) if len(size) != len(symint_placeholders): raise AssertionError( f"size length mismatch: {len(size)} != {len(symint_placeholders)}" @@ -371,7 +369,9 @@ def flatten_subclass( # outer_stride stride = x.stride() - symint_placeholders = compute_symint_placeholders(meta.outer_stride) + symint_placeholders = compute_symint_placeholders( + subclass_meta.outer_stride + ) if len(stride) != len(symint_placeholders): raise AssertionError( f"stride length mismatch: {len(stride)} != {len(symint_placeholders)}" @@ -381,7 +381,7 @@ def flatten_subclass( ) return out - xs_inner: list[int | Tensor | SymInt | OpaqueType] = [] + xs_inner: list[int | Tensor | SymInt | OpaqueBase] = [] if append_symints: if subclass_metas is None: @@ -397,10 +397,12 @@ def flatten_subclass( if subclass_metas is None: get_plain_tensors(typing.cast(Tensor, x), out=xs_inner) else: - meta = subclass_metas[idx] - if not isinstance(meta, SubclassCreationMeta): - raise AssertionError(f"expected SubclassCreationMeta, got {type(meta)}") - flatten_subclass(typing.cast(Tensor, x), meta, out=xs_inner) + subclass_meta = subclass_metas[idx] + if not isinstance(subclass_meta, SubclassCreationMeta): + raise AssertionError( + f"expected SubclassCreationMeta, got {type(subclass_meta)}" + ) + flatten_subclass(typing.cast(Tensor, x), subclass_meta, out=xs_inner) return xs_inner diff --git a/torch/_subclasses/fake_tensor.py b/torch/_subclasses/fake_tensor.py index 3146f0662518d..660a5ad3789ee 100644 --- a/torch/_subclasses/fake_tensor.py +++ b/torch/_subclasses/fake_tensor.py @@ -57,7 +57,6 @@ from types import TracebackType from torch._guards import Source - from torch._library.opaque_object import OpaqueType from torch._ops import OpOverload from torch.fx.experimental.symbolic_shapes import ShapeEnv, SymbolicContext @@ -183,8 +182,8 @@ def disable_fake_tensor_cache(fake_mode: FakeTensorMode) -> Generator[None, None def get_plain_tensors( - subclass: Tensor, *, out: list[Tensor | int | SymInt | OpaqueType] -) -> list[Tensor | int | SymInt | OpaqueType]: + subclass: Tensor, *, out: list[Tensor | int | SymInt | OpaqueBase] +) -> list[Tensor | int | SymInt | OpaqueBase]: # This function is used in Runtime, do not add redundant asserts todo = [subclass] while todo: From 56647622406f49346c436aceb0fc081e0dc717a7 Mon Sep 17 00:00:00 2001 From: angelayi Date: Fri, 20 Mar 2026 22:47:39 -0700 Subject: [PATCH 0020/1172] [opaque obj] Filter out opaque objs from inductor handling (#177991) Fixes https://github.com/pytorch/pytorch/issues/175973 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177991 Approved by: https://github.com/zou3519 --- test/test_opaque_obj_v2.py | 104 +++++++++++++++++++++++++++++++++++- torch/_inductor/lowering.py | 22 +++++--- 2 files changed, 119 insertions(+), 7 deletions(-) diff --git a/test/test_opaque_obj_v2.py b/test/test_opaque_obj_v2.py index 33bb6c33a15e9..4457913b0090d 100644 --- a/test/test_opaque_obj_v2.py +++ b/test/test_opaque_obj_v2.py @@ -12,7 +12,6 @@ import torch.distributed as dist import torch.utils._pytree as pytree from torch._dynamo.functional_export import _dynamo_graph_capture_for_export -from torch._dynamo.test_case import run_tests, TestCase from torch._dynamo.testing import ( AotEagerAndRecordGraphs, CompileCounter, @@ -30,6 +29,7 @@ ) from torch._inductor import config as inductor_config from torch._inductor.compile_fx import compile_fx +from torch._inductor.test_case import run_tests, TestCase from torch._inductor.utils import fresh_inductor_cache from torch._library.effects import EffectType from torch._library.fake_class_registry import FakeScriptObject, maybe_to_fake_obj @@ -1444,6 +1444,108 @@ def foo(counter, x): NestedCounters(Counter(1, 5)), torch.ones(2, 3) ) + def test_compile_fixed_stride_order(self): + hs_name = get_opaque_type_name(HoistedString) + + torch.library.define( + "_TestOpaqueObject::stride_op", + f"(Tensor x, {hs_name} s) -> Tensor", + tags=(torch.Tag.needs_fixed_stride_order,), + lib=self.lib, + ) + + @torch.library.impl( + "_TestOpaqueObject::stride_op", + "CompositeExplicitAutograd", + lib=self.lib, + ) + def stride_op_impl(x: torch.Tensor, s: HoistedString) -> torch.Tensor: + return x * 2.0 + + @torch.library.register_fake("_TestOpaqueObject::stride_op", lib=self.lib) + def stride_op_fake(x, s): + return torch.empty_like(x) + + def fn(x, s): + return torch.ops._TestOpaqueObject.stride_op(x, s) + + s = HoistedString("double") + x = torch.randn(4, 4) + + compiled_fn = torch.compile(fn, backend="inductor") + result = compiled_fn(x, s) + + expected = x * 2.0 + self.assertEqual(result, expected) + + def test_compile_exact_strides(self): + hs_name = get_opaque_type_name(HoistedString) + + torch.library.define( + "_TestOpaqueObject::exact_op", + f"(Tensor x, {hs_name} s) -> Tensor", + tags=(torch.Tag.needs_exact_strides,), + lib=self.lib, + ) + + @torch.library.impl( + "_TestOpaqueObject::exact_op", + "CompositeExplicitAutograd", + lib=self.lib, + ) + def exact_op_impl(x: torch.Tensor, s: HoistedString) -> torch.Tensor: + return x * 3.0 + + @torch.library.register_fake("_TestOpaqueObject::exact_op", lib=self.lib) + def exact_op_fake(x, s): + return torch.empty_like(x) + + def fn(x, s): + return torch.ops._TestOpaqueObject.exact_op(x, s) + + s = HoistedString("double") + x = torch.randn(4, 4) + + compiled_fn = torch.compile(fn, backend="inductor") + result = compiled_fn(x, s) + + expected = x * 3.0 + self.assertEqual(result, expected) + + def test_compile_contiguous_strides(self): + hs_name = get_opaque_type_name(HoistedString) + + torch.library.define( + "_TestOpaqueObject::contig_op", + f"(Tensor x, {hs_name} s) -> Tensor", + tags=(torch.Tag.needs_contiguous_strides,), + lib=self.lib, + ) + + @torch.library.impl( + "_TestOpaqueObject::contig_op", + "CompositeExplicitAutograd", + lib=self.lib, + ) + def contig_op_impl(x: torch.Tensor, s: HoistedString) -> torch.Tensor: + return x * 4.0 + + @torch.library.register_fake("_TestOpaqueObject::contig_op", lib=self.lib) + def contig_op_fake(x, s): + return torch.empty_like(x) + + def fn(x, s): + return torch.ops._TestOpaqueObject.contig_op(x, s) + + s = HoistedString("double") + x = torch.randn(4, 4) + + compiled_fn = torch.compile(fn, backend="inductor") + result = compiled_fn(x, s) + + expected = x * 4.0 + self.assertEqual(result, expected) + def test_export_joint(self): torch.library.define( "_TestOpaqueObject::module_mul", diff --git a/torch/_inductor/lowering.py b/torch/_inductor/lowering.py index a4a36534e2152..a72afb1c4c57f 100644 --- a/torch/_inductor/lowering.py +++ b/torch/_inductor/lowering.py @@ -2882,16 +2882,22 @@ def inner_fn(index): return result +def _is_tensor_irnode(x): + return isinstance(x, ir.IRNode) and not isinstance(x, ir.NonTensorObj) + + def require_dense(_, *args, **kwargs): args, kwargs = pytree.tree_map_only( - ir.IRNode, ir.ExternKernel.require_stride1, (args, kwargs) + _is_tensor_irnode, ir.ExternKernel.require_stride1, (args, kwargs) ) return args, kwargs def require_contiguous(_, *args, **kwargs): args, kwargs = pytree.tree_map_only( - ir.IRNode, ir.ExternKernel.require_contiguous, (args, kwargs) + _is_tensor_irnode, + ir.ExternKernel.require_contiguous, + (args, kwargs), ) return args, kwargs @@ -2900,14 +2906,18 @@ def require_contiguous_strides(_, *args, **kwargs): # TODO: combine this with require_contiguous after # https://github.com/pytorch/pytorch/pull/148235 lands. args, kwargs = pytree.tree_map_only( - ir.IRNode, ir.ExternKernel.require_contiguous_strides, (args, kwargs) + _is_tensor_irnode, + ir.ExternKernel.require_contiguous_strides, + (args, kwargs), ) return args, kwargs def require_channels_last(_, *args, **kwargs): args, kwargs = pytree.tree_map_only( - ir.IRNode, ir.ExternKernel.require_channels_last, (args, kwargs) + _is_tensor_irnode, + ir.ExternKernel.require_channels_last, + (args, kwargs), ) return args, kwargs @@ -2939,7 +2949,7 @@ def constrain_to_fake_tensors(args, kwargs, fake_args, fake_kwargs): def constrain_to_fx_strides(fx_node, *args, **kwargs): def apply_constraint(arg, fx_arg): - if isinstance(arg, ir.IRNode): + if _is_tensor_irnode(arg): stride_order = ir.get_stride_order( fx_arg.meta["val"].stride(), V.graph.sizevars.shape_env ) @@ -2965,7 +2975,7 @@ def sdpa_constraint(fx_node, *args, **kwargs): """Apply stride constraints to SDPA inputs, ensuring dense last dimension.""" def apply_constraint(idx, arg, fx_arg): - if not isinstance(arg, ir.IRNode): + if not _is_tensor_irnode(arg): return arg meta_val = fx_arg.meta["val"] From 60cc45cc126949cf39d877dbc7e9859d27ad1b2d Mon Sep 17 00:00:00 2001 From: Aaron Gokaslan Date: Mon, 23 Mar 2026 18:07:11 +0000 Subject: [PATCH 0021/1172] [BE][Ez]: Make distribution more efficient with aminmax call. (#178018) Make multinomial distribution check more efficient. Pull Request resolved: https://github.com/pytorch/pytorch/pull/178018 Approved by: https://github.com/malfet --- aten/src/ATen/native/Distributions.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/aten/src/ATen/native/Distributions.cpp b/aten/src/ATen/native/Distributions.cpp index 5f34ed9d24c17..6382b1ecaace1 100644 --- a/aten/src/ATen/native/Distributions.cpp +++ b/aten/src/ATen/native/Distributions.cpp @@ -594,7 +594,8 @@ Tensor& multinomial_out(const Tensor& self, // https://github.com/pytorch/pytorch/issues/11931#issuecomment-625882503 if (!with_replacement || n_sample == 1) { // Sanity checks on `self`. - auto is_valid = ((self.max() < INFINITY) & (self.min() >= 0)); + auto [self_min, self_max] = self.aminmax(); + auto is_valid = ((self_max < INFINITY) & (self_min >= 0)); at::_assert_async(is_valid, "probability tensor contains either `inf`, `nan` or element < 0"); at::Tensor zero_prob_condition; if (self.dim() == 1){ From 549e5c77bb09262be27b441ee27eeaebffcc7905 Mon Sep 17 00:00:00 2001 From: Yang Wang Date: Mon, 23 Mar 2026 18:18:52 +0000 Subject: [PATCH 0022/1172] fix binary validation tag and mino verification (#177993) fix binary validation tag and mino verification Pull Request resolved: https://github.com/pytorch/pytorch/pull/177993 Approved by: https://github.com/atalman --- .ci/pytorch/smoke_test/check_wheel_tags.py | 181 ++++++++++++--------- 1 file changed, 104 insertions(+), 77 deletions(-) diff --git a/.ci/pytorch/smoke_test/check_wheel_tags.py b/.ci/pytorch/smoke_test/check_wheel_tags.py index 47853183e0270..212e302f169dd 100644 --- a/.ci/pytorch/smoke_test/check_wheel_tags.py +++ b/.ci/pytorch/smoke_test/check_wheel_tags.py @@ -1,5 +1,4 @@ """Validate wheel platform tags and macOS dylib minos. - Supports two modes: 1. Pre-install: reads .whl files from PYTORCH_FINAL_PACKAGE_DIR 2. Post-install: reads metadata from installed torch package (soft warnings) @@ -11,6 +10,7 @@ import re import subprocess import sys +import tempfile import zipfile from pathlib import Path @@ -66,8 +66,17 @@ def check_wheel_platform_tag() -> None: if target_os == "linux" and platform.machine() == "aarch64": target_os = "linux-aarch64" expected_python = f"cp{sys.version_info.major}{sys.version_info.minor}" + import sysconfig + abiflags = getattr(sys, "abiflags", "") + if not abiflags and ( + os.getenv("MATRIX_PYTHON_VERSION", "").endswith("t") + or bool(sysconfig.get_config_var("Py_GIL_DISABLED")) + or not getattr(sys, "_is_gil_enabled", lambda: True)() + ): + abiflags = "t" expected_abi = f"cp{sys.version_info.major}{sys.version_info.minor}{abiflags}" + print(f"Expected ABI tag: {expected_abi}") platform_pattern = EXPECTED_PLATFORM_TAGS.get(target_os) if not platform_pattern: @@ -113,7 +122,6 @@ def check_wheel_platform_tag() -> None: f"expected format: --" ) raise RuntimeError(msg) - continue python_tag, abi_tag, platform_tag = parts @@ -143,92 +151,111 @@ def check_wheel_platform_tag() -> None: print(f"OK: Wheel tag(s) valid for {source}: {', '.join(tags)}") -def check_mac_wheel_minos() -> None: - """Check that dylib minos matches the wheel platform tag on macOS. +def _check_dylibs_minos(dylibs: list, expected_minos: str, source: str) -> None: + mismatches = [] + for dylib in dylibs: + try: + result = subprocess.run( + ["otool", "-l", str(dylib)], + capture_output=True, + text=True, + timeout=30, + ) + except Exception: + continue - Extracts dylibs from the .whl in PYTORCH_FINAL_PACKAGE_DIR to a temp dir, - then verifies each dylib's minos (from otool -l) matches the platform tag. - """ + minos = None + lines = result.stdout.splitlines() + for i, line in enumerate(lines): + s = line.strip() + if "LC_BUILD_VERSION" in s: + for j in range(i + 1, min(i + 6, len(lines))): + if lines[j].strip().startswith("minos"): + minos = lines[j].strip().split()[1] + break + break + if "LC_VERSION_MIN_MACOSX" in s: + for j in range(i + 1, min(i + 4, len(lines))): + if lines[j].strip().startswith("version"): + minos = lines[j].strip().split()[1] + break + break + + if minos and minos != expected_minos: + mismatches.append(f"{dylib.name}: minos={minos}, expected={expected_minos}") + + if mismatches: + raise RuntimeError( + f"minos/platform tag mismatch in {len(mismatches)} dylib(s):\n" + + "\n".join(f" {m}" for m in mismatches) + ) + print( + f"OK: All {len(dylibs)} dylib(s) have minos matching " + f"platform tag ({expected_minos}) for {source}" + ) + + +def check_mac_wheel_minos() -> None: if sys.platform != "darwin": return wheel_dir = os.getenv("PYTORCH_FINAL_PACKAGE_DIR", "") - if not wheel_dir or not os.path.isdir(wheel_dir): - print("PYTORCH_FINAL_PACKAGE_DIR not set, skipping wheel minos check") - return - whls = list(Path(wheel_dir).glob("*.whl")) - if not whls: - print(f"No .whl files in {wheel_dir}, skipping wheel minos check") - return + if wheel_dir and os.path.isdir(wheel_dir): + # Mode 1: extract dylibs from .whl file + whls = list(Path(wheel_dir).glob("*.whl")) + if not whls: + print(f"No .whl files in {wheel_dir}, skipping wheel minos check") + return + + macos_whl_re = re.compile(r"macosx_(\d+)_(\d+)_(\w+)\.whl$") + for whl in whls: + print(f"Checking wheel tag minos for: {whl.name}") + m = macos_whl_re.search(whl.name) + if not m: + print(f"No macOS platform tag in {whl.name}, skipping") + continue + expected_minos = f"{m.group(1)}.{m.group(2)}" + + with tempfile.TemporaryDirectory() as tmpdir: + with zipfile.ZipFile(whl, "r") as zf: + dylib_names = [n for n in zf.namelist() if n.endswith(".dylib")] + if not dylib_names: + print("No .dylib files in wheel, skipping minos check") + continue + for name in dylib_names: + zf.extract(name, tmpdir) + dylibs = list(Path(tmpdir).rglob("*.dylib")) + _check_dylibs_minos(dylibs, expected_minos, whl.name) + else: + # Mode 2: read from installed torch package + print("PYTORCH_FINAL_PACKAGE_DIR not set, checking installed torch dylibs") + try: + tags = _extract_installed_wheel_tags("torch") + except Exception as e: + print(f"Could not read installed torch metadata: {e}, skipping") + return - import tempfile + expected_minos = None + for tag_str in tags: + m = re.search(r"macosx_(\d+)_(\d+)_\w+", tag_str) + if m: + expected_minos = f"{m.group(1)}.{m.group(2)}" + break - for whl in whls: - print(f"Checking wheel tag minos for: {whl.name}") + if not expected_minos: + print("No macOS platform tag found in installed torch metadata, skipping") + return - m = re.search(r"macosx_(\d+)_(\d+)_(\w+)\.whl$", whl.name) - if not m: - print(f"No macOS platform tag in {whl.name}, skipping") - continue + print(f"Expected minos from installed wheel tag: {expected_minos}") - expected_minos = f"{m.group(1)}.{m.group(2)}" - print(f"Expected minos from platform tag: {expected_minos}") - - # Extract dylibs from wheel to temp dir - with tempfile.TemporaryDirectory() as tmpdir: - with zipfile.ZipFile(whl, "r") as zf: - dylib_names = [n for n in zf.namelist() if n.endswith(".dylib")] - if not dylib_names: - print("No .dylib files in wheel, skipping minos check") - continue - for name in dylib_names: - zf.extract(name, tmpdir) - - dylibs = list(Path(tmpdir).rglob("*.dylib")) - mismatches = [] - for dylib in dylibs: - try: - result = subprocess.run( - ["otool", "-l", str(dylib)], - capture_output=True, - text=True, - timeout=30, - ) - except Exception: - continue - - minos = None - lines = result.stdout.splitlines() - for i, line in enumerate(lines): - s = line.strip() - if "LC_BUILD_VERSION" in s: - for j in range(i + 1, min(i + 6, len(lines))): - if lines[j].strip().startswith("minos"): - minos = lines[j].strip().split()[1] - break - break - if "LC_VERSION_MIN_MACOSX" in s: - for j in range(i + 1, min(i + 4, len(lines))): - if lines[j].strip().startswith("version"): - minos = lines[j].strip().split()[1] - break - break + import torch - if minos and minos != expected_minos: - mismatches.append( - f"{dylib.name}: minos={minos}, expected={expected_minos}" - ) - - if mismatches: - raise RuntimeError( - f"minos/platform tag mismatch in {len(mismatches)} dylib(s):\n" - + "\n".join(f" {m}" for m in mismatches) - ) - print( - f"OK: All {len(dylibs)} dylib(s) have minos matching " - f"platform tag ({expected_minos})" - ) + torch_dir = Path(torch.__file__).parent + dylibs = list(torch_dir.rglob("*.dylib")) + if not dylibs: + raise RuntimeError("No .dylib files found in installed torch") + _check_dylibs_minos(dylibs, expected_minos, "installed torch") if __name__ == "__main__": From dd2bf2cfccdad92fe6f557ac7fc838d155cd7558 Mon Sep 17 00:00:00 2001 From: Lucas Kabela Date: Mon, 23 Mar 2026 18:21:20 +0000 Subject: [PATCH 0023/1172] [Bugfix] Fix meta conv backwards strides (#177175) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Summary This pr makes changes to memory format handling by: 1. fake_impls.py: - Unbacked SymInts now fall back to a strides-based heuristic instead of mem_fmt = None - The backward now applies a separate weight_fmt to grad_weight (using suggest_memory_format(weight) for conv2d/3d, mem_fmt for conv1d) instead of the same mem_fmt for both outputs - convert now explicitly fixes the strides when the meta kernel set a non-contiguous format but we want contiguous, and grad_bias is explicitly passed contiguous_format instead of None 2. _meta_registrations.py: Two explanatory comments only — no logic change. ### Test ```bash python test/test_fake_tensor.py FakeTensorOperatorInvariants.test_convolution_backward_meta_kernel_channels_last ``` Breaks on main, but passes on this PR Authored with Claude Opus 4.6 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177175 Approved by: https://github.com/aorenste --- test/test_meta.py | 53 +++++++++++++++++++++++ torch/_meta_registrations.py | 23 +++++----- torch/_subclasses/fake_impls.py | 76 +++++++++++++-------------------- 3 files changed, 93 insertions(+), 59 deletions(-) diff --git a/test/test_meta.py b/test/test_meta.py index b0c9e66c7a581..7ac6581223d3a 100644 --- a/test/test_meta.py +++ b/test/test_meta.py @@ -20,6 +20,7 @@ from torch.testing._internal.common_utils import ( TestCase, skipIfCrossRef, + skipIfTorchDynamo, suppress_warnings, TEST_WITH_TORCHDYNAMO, run_tests, @@ -1895,8 +1896,60 @@ def fn(input, weight, bias, need_grad_input): else: self.assertEqual(out_dtype, [in_dtype,]) +class TestMetaKernelConv(TestCase): + @skipIfTorchDynamo("tests raw meta kernel, not dynamo") + def test_convolution_backward_meta_kernel_channels_last(self): + """Test the meta kernel directly (device='meta', no FakeTensorMode). + This exercises the @register_meta path used by torch.export, which + does NOT go through the FakeTensor intercept in fake_impls.py. + """ + # channels_last grad_output + contiguous input/weight -> contiguous + grad_out = torch.empty(2, 3, 4, 4, device="meta").to( + memory_format=torch.channels_last + ) + inp = torch.empty(2, 3, 4, 4, device="meta") + w = torch.empty(3, 3, 3, 3, device="meta") + gi, gw, _ = torch.ops.aten.convolution_backward( + grad_out, + inp, + w, + [3], + [1, 1], + [1, 1], + [1, 1], + False, + [0, 0], + 1, + [True, True, True], + ) + self.assertTrue(gi.is_contiguous()) + self.assertTrue(gw.is_contiguous()) + + # contiguous grad_output + channels_last input -> channels_last + grad_out2 = torch.empty(2, 3, 4, 4, device="meta") + inp2 = torch.empty(2, 3, 4, 4, device="meta").to( + memory_format=torch.channels_last + ) + gi2, gw2, _ = torch.ops.aten.convolution_backward( + grad_out2, + inp2, + w, + [3], + [1, 1], + [1, 1], + [1, 1], + False, + [0, 0], + 1, + [True, True, True], + ) + self.assertTrue(gi2.is_contiguous(memory_format=torch.channels_last)) + self.assertTrue(gw2.is_contiguous(memory_format=torch.channels_last)) + + instantiate_device_type_tests(TestMeta, globals()) + def print_op_str_if_not_supported(op_str): op = OperatorName.parse(op_str) packet = getattr(torch.ops.aten, str(op.name)) diff --git a/torch/_meta_registrations.py b/torch/_meta_registrations.py index 9227ff808ced7..ac1b1bad1bac3 100644 --- a/torch/_meta_registrations.py +++ b/torch/_meta_registrations.py @@ -2618,6 +2618,11 @@ def meta_conv( if guard_or_false(input_tensor.size(input_channels_dim) == 0): shape_out[output_channels_dim] = 0 + # Memory format is left as contiguous: meta tensors have no device info, + # so _select_conv_backend returns Overrideable and the correct format + # cannot be determined here. The FakeTensor path (torch.compile, export) + # intercepts via a register_op_impl in fake_impls.py before reaching this + # kernel and uses FakeTensor.fake_device for an accurate answer. out = input_tensor.new_empty(shape_out) return out @@ -3676,18 +3681,13 @@ def meta_convolution_backward( backend_grad_weight = None backend_grad_bias = None - # Backend layout expectation: GPU backends (CUDA via cudnn_conv_suggest_memory_format, - # MPS via mps_conv_use_channels_last) return channels_last outputs when either input - # tensor is channels_last. This must be matched here to avoid stride assertion failures - # in inductor when the predicted strides don't match actual backend output strides. + # All GPU backends compute output memory format via + # determine_backend_memory_format(input, weight, backend) — which calls + # cudnn_conv_suggest_memory_format(input, weight), mps_conv_use_channels_last(input, weight), + # etc. The format depends only on input and weight, NOT on grad_output. + # Both grad_input and grad_weight use this same backend_memory_format. # See: https://github.com/pytorch/pytorch/issues/171622 - # - # Memory format inference rules (matching backend behavior): - # - grad_input format: derived from grad_output and weight - # - grad_weight format: derived from input and grad_output def _conv_memory_format(t1, t2): - # Match the logic in cudnn_conv_suggest_memory_format and mps_conv_use_channels_last: - # Use channels_last if either tensor suggests it fmt1 = suggest_memory_format(t1) fmt2 = suggest_memory_format(t2) if fmt1 == torch.channels_last or fmt2 == torch.channels_last: @@ -3696,13 +3696,12 @@ def _conv_memory_format(t1, t2): return torch.channels_last_3d return torch.contiguous_format + memory_format = _conv_memory_format(input_, weight_) if output_mask[0]: - memory_format = _conv_memory_format(grad_output_, weight_) backend_grad_input = grad_output_.new_empty(input_.size()).to( memory_format=memory_format ) if output_mask[1]: - memory_format = _conv_memory_format(input_, grad_output_) backend_grad_weight = grad_output_.new_empty(weight_.size()).to( memory_format=memory_format ) diff --git a/torch/_subclasses/fake_impls.py b/torch/_subclasses/fake_impls.py index a602e266ba11e..16c0b530b92af 100644 --- a/torch/_subclasses/fake_impls.py +++ b/torch/_subclasses/fake_impls.py @@ -1329,71 +1329,53 @@ def conv( _, new_kwargs = _normalize_function_or_error( func, args=args, kwargs=kwargs, normalize_to_only_use_kwargs=True ) - device = new_kwargs["input"].fake_device + input_ = new_kwargs["input"] + weight = new_kwargs["weight"] + device = input_.fake_device # need to re-enable mode so the tensors report fake device with fake_mode: - # if the input is unsqueezed is done in Convolution.cpp we get segfault - k = new_kwargs["weight"].ndim + # if the input is unsqueezed in Convolution.cpp we get segfault + k = weight.ndim # Avoid importing sympy at a module level from torch.fx.experimental.symbolic_shapes import has_guarding_hint - all_hinted = all( - has_guarding_hint(s) for s in new_kwargs["input"].shape - ) and all(has_guarding_hint(s) for s in new_kwargs["weight"].shape) + all_hinted = all(has_guarding_hint(s) for s in input_.shape) and all( + has_guarding_hint(s) for s in weight.shape + ) if not all_hinted: # TODO: We can make this a little more faithful with best effort # channels last detection (but only if it's statically obvious!) mem_fmt = None else: - if func is aten.convolution.default: - conv_backend = torch._C._select_conv_backend(**new_kwargs) - else: - conv_backend = torch._C._select_conv_backend( - new_kwargs["input"], - new_kwargs["weight"], - bias=None, - stride=new_kwargs["stride"], - padding=new_kwargs["padding"], - dilation=new_kwargs["dilation"], - transposed=new_kwargs["transposed"], - output_padding=new_kwargs["output_padding"], - groups=new_kwargs["groups"], - bias_sizes=new_kwargs["bias_sizes"], - ) + # convolution has "bias" but not "bias_sizes"; convolution_backward + # has "bias_sizes" but not "bias". .get() handles both with one call. + bias = new_kwargs.get("bias") + select_kwargs: dict[str, object] = dict( + stride=new_kwargs["stride"], + padding=new_kwargs["padding"], + dilation=new_kwargs["dilation"], + transposed=new_kwargs["transposed"], + output_padding=new_kwargs["output_padding"], + groups=new_kwargs["groups"], + bias=bias, + ) + if bias is None: + select_kwargs["bias_sizes"] = new_kwargs.get("bias_sizes") + conv_backend = torch._C._select_conv_backend( + input_, weight, **select_kwargs + ) # Expand 1d -> 2d. # Note: Avoid expanding before calling _select_conv_backend, # as the function handles 2D expansion internally. - if ( - k == 3 - and not new_kwargs["input"].is_mkldnn - and not new_kwargs["input"].is_xpu - ): + if k == 3 and not input_.is_mkldnn and not input_.is_xpu: # Note: Using input.to(memory_format=contiguous) does not work. - new_kwargs["input"] = new_kwargs["input"].contiguous().unsqueeze(2) - new_kwargs["weight"] = new_kwargs["weight"].unsqueeze(2) - if len(new_kwargs["stride"]) == 1: - new_kwargs["stride"].insert(0, 1) - new_kwargs["padding"].insert(0, 0) - new_kwargs["dilation"].insert(0, 1) - new_kwargs["output_padding"].insert(0, 0) + input_ = input_.contiguous().unsqueeze(2) + weight = weight.unsqueeze(2) mem_fmt = torch._C._conv_determine_backend_memory_format( - new_kwargs["input"], new_kwargs["weight"], conv_backend + input_, weight, conv_backend ) - # revert 2d -> 1d - if ( - k == 3 - and not new_kwargs["input"].is_mkldnn - and not new_kwargs["input"].is_xpu - ): - new_kwargs["input"] = new_kwargs["input"].squeeze(2) - new_kwargs["weight"] = new_kwargs["weight"].squeeze(2) - if len(new_kwargs["stride"]) == 2: - new_kwargs["stride"].pop(0) - new_kwargs["padding"].pop(0) - new_kwargs["dilation"].pop(0) - new_kwargs["output_padding"].pop(0) def convert( t: torch.Tensor | None, mem_fmt: torch.memory_format | None From 1c5103a1ced8a3753e6c89b1cc2f9019dee1776d Mon Sep 17 00:00:00 2001 From: drisspg Date: Fri, 20 Mar 2026 23:42:48 +0000 Subject: [PATCH 0024/1172] Add scratch space for experiments (#178026) Pull Request resolved: https://github.com/pytorch/pytorch/pull/178026 Approved by: https://github.com/ezyang, https://github.com/izaitsevfb --- .gitignore | 1 + CLAUDE.md | 4 ++++ 2 files changed, 5 insertions(+) diff --git a/.gitignore b/.gitignore index 54e3fa32b2fdb..5670e404441f5 100644 --- a/.gitignore +++ b/.gitignore @@ -8,6 +8,7 @@ ## PyTorch +agent_space/ .coverage coverage.xml .dmypy.json diff --git a/CLAUDE.md b/CLAUDE.md index e83bcdf5e0e12..640a24bbcb410 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -1,3 +1,7 @@ +# Scratch Space + +Use `agent_space/` (git-ignored, at repo root) for temporary scripts, scratch files, and throwaway experiments. Do not commit files from this directory. + # PR Review When asked to review a PR, always use the /pr-review skill. From 488f6f9bd2974b65b80e4d0a9a360f095847a052 Mon Sep 17 00:00:00 2001 From: Sean McGovern Date: Mon, 23 Mar 2026 19:05:53 +0000 Subject: [PATCH 0025/1172] [DTensor] Add sharding strategy for aten.squeeze.dims (#173563) Fixes #173521 Fixes #166124 Extend `dim_squeeze` to handle multiple dimensions by normalizing all dim variants to a target dimension set. This unifies the logic into a single code path. Fix the long-standing FIXME in dim_squeeze where squeeze(dim=None) could incorrectly remove sharded dimensions whose local size happened to be 1 (despite global size > 1). Canonicalizes all squeeze variants to squeeze.dims at the sharding propagation level using global shape to determine which dimensions are truly singleton. Strategy validator: 74 correct, 0 incorrect, 0 missing. This is without the P(max/min) - R rules mentioned below. - Add test_squeeze_variants to test all squeeze variants with DTensor ~~Note: op_db test remains xfail due to pre-existing bug where local squeeze removes sharded dims with local size 1 (see PR #166862).~~ That PR is/will be closed in favor of this approach that avoids a custom handler Pull Request resolved: https://github.com/pytorch/pytorch/pull/173563 Approved by: https://github.com/wconstab --- test/distributed/tensor/test_dtensor_ops.py | 18 ++- test/distributed/tensor/test_view_ops.py | 148 ++++++++++++++++++++ torch/distributed/tensor/_dispatch.py | 32 ++++- torch/distributed/tensor/_ops/_view_ops.py | 19 ++- torch/distributed/tensor/_sharding_prop.py | 62 ++++++++ 5 files changed, 259 insertions(+), 20 deletions(-) diff --git a/test/distributed/tensor/test_dtensor_ops.py b/test/distributed/tensor/test_dtensor_ops.py index 5a228d3e9a275..65145b37c42ff 100644 --- a/test/distributed/tensor/test_dtensor_ops.py +++ b/test/distributed/tensor/test_dtensor_ops.py @@ -207,8 +207,6 @@ def repurpose_ops(op_db, base_test_name, derived_test_name): # DTensorConverter can't convert sparse tensor inputs xfail("sparse.sampled_addmm"), xfail("sparse.mm", "reduce"), - # bug in squeeze.dims strategy: TypeError with empty dims arg - xfail("squeeze", "multiple"), # meta tensor data not allocated yet during tensor_split xfail("tensor_split"), # output_specs count mismatch in unsafe_split strategy @@ -251,7 +249,6 @@ def repurpose_ops(op_db, base_test_name, derived_test_name): skip("_segment_reduce", "lengths"), skip("_segment_reduce", "offsets"), # TODO: fix the following ops - skip("squeeze"), skip("empty"), skip("empty_strided"), skip("empty_like"), @@ -289,6 +286,8 @@ def repurpose_ops(op_db, base_test_name, derived_test_name): xfail("permute"), xfail("select"), xfail("slice"), + xfail("squeeze"), + xfail("squeeze", "multiple"), xfail("t"), xfail("transpose_copy"), xfail("unsqueeze"), @@ -320,6 +319,11 @@ def repurpose_ops(op_db, base_test_name, derived_test_name): xfail("lu_unpack"), xfail("scatter"), xfail("scatter_add"), + # batch_norm variants decompose through squeeze.dims → as_strided under + # compilation, and DTensor has no as_strided strategy. + xfail("_native_batch_norm_legit"), + xfail("native_batch_norm"), + xfail("nn.functional.batch_norm"), # False positives: these have no sharding strategy and their # eager DTensor failure is registered elsewhere. xfail("nn.functional.margin_ranking_loss"), @@ -368,7 +372,6 @@ def repurpose_ops(op_db, base_test_name, derived_test_name): dtensor_fails_no_strategy = { xfail("_batch_norm_with_update"), xfail("_chunk_cat"), - xfail("_native_batch_norm_legit"), xfail("_unsafe_masked_index"), xfail("_unsafe_masked_index_put_accumulate"), xfail("_upsample_bilinear2d_aa"), @@ -406,8 +409,6 @@ def repurpose_ops(op_db, base_test_name, derived_test_name): xfail("max_pool2d_with_indices_backward"), xfail("multinomial"), xfail("nanquantile"), - xfail("native_batch_norm"), - xfail("nn.functional.batch_norm"), xfail("nn.functional.bilinear"), xfail("nn.functional.grid_sample"), xfail("nn.functional.group_norm"), @@ -807,6 +808,7 @@ def assertEqualOnRank(self, x, y, msg=None, *, rank=0): xfail("__rmatmul__"), xfail("_segment_reduce", "lengths"), xfail("_segment_reduce", "offsets"), + xfail("_native_batch_norm_legit"), xfail("_unsafe_masked_index"), xfail("addmm"), xfail("addmm", "decomposed"), @@ -858,7 +860,9 @@ def assertEqualOnRank(self, x, y, msg=None, *, rank=0): xfail("new_empty_strided"), xfail("new_full"), xfail("new_ones"), + xfail("native_batch_norm"), xfail("new_zeros"), + xfail("nn.functional.batch_norm"), xfail("nn.functional.celu"), xfail("nn.functional.conv1d"), xfail("nn.functional.conv2d"), @@ -904,7 +908,7 @@ def assertEqualOnRank(self, x, y, msg=None, *, rank=0): xfail("special.log_ndtr"), xfail("special.ndtri"), xfail("special.spherical_bessel_j0"), - xfail("squeeze", "multiple"), + xfail("squeeze_copy"), xfail("std_mean"), xfail("topk"), xfail("transpose_copy"), diff --git a/test/distributed/tensor/test_view_ops.py b/test/distributed/tensor/test_view_ops.py index cf3da8ec32900..d8969a43aaca9 100644 --- a/test/distributed/tensor/test_view_ops.py +++ b/test/distributed/tensor/test_view_ops.py @@ -692,6 +692,154 @@ def test_squeeze_(self): ) self.assertEqual(dist_x.placements, [Partial(), Shard(0)]) + # squeeze_ should not trigger any communication + y = torch.randn((1, 4), device=self.device_type) + dist_y = DTensor.from_local(y, mesh_2d, [Partial(), Shard(1)]) + with CommDebugMode() as comm_mode: + torch.ops.aten.squeeze_.dim(dist_y, 0) + self.assertEqual(comm_mode.get_total_counts(), 0) + + @with_comms + def test_squeeze_variants(self): + """Test squeeze.default, squeeze.dim, and squeeze.dims with DTensor.""" + mesh = init_device_mesh(self.device_type, (self.world_size,)) + + # squeeze.dims on sharded tensor - squeeze non-sharded dims + with self.subTest("dims_sharded"): + x = torch.randn(self.world_size, 1, 1, 8, device=self.device_type) + dt = distribute_tensor(x, mesh, [Shard(0)]) + with CommDebugMode() as comm_mode: + result = dt.squeeze((1, 2)) + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([self.world_size, 8])) + self.assertEqual(result.placements, (Shard(0),)) + self.assertEqual(result.to_local().shape, torch.Size([1, 8])) + + # squeeze.dim on sharded tensor - squeeze non-sharded dim + with self.subTest("dim_sharded"): + x = torch.randn(self.world_size, 1, 8, device=self.device_type) + dt = distribute_tensor(x, mesh, [Shard(0)]) + with CommDebugMode() as comm_mode: + result = dt.squeeze(1) + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([self.world_size, 8])) + self.assertEqual(result.placements, (Shard(0),)) + self.assertEqual(result.to_local().shape, torch.Size([1, 8])) + + # squeeze.default on replicated tensor + with self.subTest("default_replicated"): + x = torch.randn(4, 1, 1, 8, device=self.device_type) + dt = distribute_tensor(x, mesh, [Replicate()]) + with CommDebugMode() as comm_mode: + result = dt.squeeze() + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([4, 8])) + self.assertEqual(result.placements, (Replicate(),)) + + # squeeze.dims on replicated tensor + with self.subTest("dims_replicated"): + x = torch.randn(2, 1, 3, 1, device=self.device_type) + dt = distribute_tensor(x, mesh, [Replicate()]) + with CommDebugMode() as comm_mode: + result = dt.squeeze((1, 3)) + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([2, 3])) + self.assertEqual(result.placements, (Replicate(),)) + + # squeeze non-singleton dim is a no-op + with self.subTest("non_singleton_noop"): + x = torch.randn(self.world_size, 4, device=self.device_type) + dt = distribute_tensor(x, mesh, [Shard(0)]) + with CommDebugMode() as comm_mode: + result = dt.squeeze(1) + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([self.world_size, 4])) + self.assertEqual(result.placements, (Shard(0),)) + + # Partial passes through squeeze unchanged + with self.subTest("partial_max_passthrough"): + x = torch.randn(1, 4, device=self.device_type) + dt = DTensor.from_local(x, mesh, [Partial("max")]) + with CommDebugMode() as comm_mode: + result = dt.squeeze(0) + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([4])) + self.assertEqual(result.placements, (Partial("max"),)) + + # Partial("sum") also passes through + with self.subTest("partial_sum_passthrough"): + x = torch.randn(1, device=self.device_type) + dt = DTensor.from_local(x, mesh, [Partial("sum")]) + with CommDebugMode() as comm_mode: + result = dt.squeeze() + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([])) + self.assertEqual(result.placements, (Partial("sum"),)) + + # squeeze must not remove sharded dim (local size 1, global size > 1) + with self.subTest("preserve_sharded_dim_default"): + x = ( + torch.arange(self.world_size * 8, device=self.device_type) + .reshape(self.world_size, 8) + .float() + ) + dt = distribute_tensor(x, mesh, [Shard(0)]) + with CommDebugMode() as comm_mode: + result = dt.squeeze() + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([self.world_size, 8])) + self.assertEqual(result._local_tensor.shape, torch.Size([1, 8])) + self.assertEqual(result.placements, (Shard(0),)) + self.assertEqual(result.full_tensor(), x) + + # same as above but via squeeze.dim (single int arg) + with self.subTest("preserve_sharded_dim_explicit"): + x = ( + torch.arange(self.world_size * 8, device=self.device_type) + .reshape(self.world_size, 8) + .float() + ) + dt = distribute_tensor(x, mesh, [Shard(0)]) + with CommDebugMode() as comm_mode: + result = dt.squeeze(0) + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([self.world_size, 8])) + self.assertEqual(result._local_tensor.shape, torch.Size([1, 8])) + self.assertEqual(result.placements, (Shard(0),)) + self.assertEqual(result.full_tensor(), x) + + # squeeze.dims with mixed singleton/non-singleton dims + with self.subTest("mixed_dims"): + x = torch.randn(1, 4, 1, device=self.device_type) + dt = distribute_tensor(x, mesh, [Replicate()]) + with CommDebugMode() as comm_mode: + result = dt.squeeze((0, 1, 2)) + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([4])) + self.assertEqual(result.placements, (Replicate(),)) + self.assertEqual(result.full_tensor(), x.squeeze()) + + # sharded non-singleton dim preserved, shard index shifts + with self.subTest("shard_index_shift"): + x = torch.randn(1, self.world_size, 1, device=self.device_type) + dt = distribute_tensor(x, mesh, [Shard(1)]) + with CommDebugMode() as comm_mode: + result = dt.squeeze((0, 1, 2)) + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result.shape, torch.Size([self.world_size])) + self.assertEqual(result._local_tensor.shape, torch.Size([1])) + self.assertEqual(result.placements, (Shard(0),)) + self.assertEqual(result.full_tensor(), x.squeeze((0, 2))) + + # S(0) on globally-singleton dim becomes R after squeeze (#174136) + with self.subTest("singleton_shard_becomes_replicate"): + x = torch.randn(1, 4, device=self.device_type) + dt = distribute_tensor(x, mesh, [Shard(0)]) + result = dt.squeeze() + self.assertEqual(result.shape, torch.Size([4])) + self.assertEqual(result.placements, (Replicate(),)) + self.assertEqual(result.full_tensor(), x.squeeze()) + @with_comms def test_storage_offset_slice(self): """ diff --git a/torch/distributed/tensor/_dispatch.py b/torch/distributed/tensor/_dispatch.py index d6bf1ad188d40..d346c8e17e117 100644 --- a/torch/distributed/tensor/_dispatch.py +++ b/torch/distributed/tensor/_dispatch.py @@ -173,6 +173,11 @@ def __init__(self) -> None: aten.bernoulli.default, aten.bernoulli_.float, } + self._squeeze_inplace_ops = { + aten.squeeze_.dim, + aten.squeeze_.default, + aten.squeeze_.dims, + } self._custom_op_handlers = { aten.is_same_size.default: is_same_size_handler, aten.is_pinned.default: is_pinned_handler, @@ -390,7 +395,17 @@ def _dispatch_get_local_results_slow_path( local_results = op_call(*local_tensor_args, **op_info.local_kwargs) else: # normal case, run local sharded op computation - local_results = op_call(*local_tensor_args, **op_info.local_kwargs) + if ( + output_sharding.needs_redistribute + and output_sharding.redistribute_schema is not None + and output_sharding.redistribute_schema.op != op_call + ): + # Op was rewritten (e.g., squeeze.default → squeeze.dims) + local_results = output_sharding.redistribute_schema.op( + *local_tensor_args, **op_info.local_kwargs + ) + else: + local_results = op_call(*local_tensor_args, **op_info.local_kwargs) else: # For a non-participating device (happens on rank that does not belong to @@ -481,11 +496,9 @@ def _dispatch_fast_path_python_tail( if not isinstance(args[0], dtensor.DTensor): raise AssertionError - # NOTE: aten.squeeze_.dim is an inplace op but it also may change - # the inplace argument's tensor meta. Here we choose to special case - # this op because as far as I know this is the only inplace op that - # has such as behavior. We can extend this special case if necessary. - if op_call == aten.squeeze_.dim: + # NOTE: squeeze_ inplace ops may change the tensor's metadata + # (shape/strides). We special-case them to update the spec. + if op_call in self._squeeze_inplace_ops: # update the spec to handle tensor meta changes args[0]._spec = output_spec # use return_and_correct_aliasing to match the outer and the inner @@ -593,6 +606,13 @@ def redistribute_local_args( else: new_local_args.append(arg_spec) + # Append extra non-tensor args from rewritten schema (e.g., dims tuple). + if use_val_from_redistribute_schema: + for i in range( + len(op_info.flat_args_schema), len(flatten_args_schema_to_reshard) + ): + new_local_args.append(flatten_args_schema_to_reshard[i]) + op_info.local_args = tuple(new_local_args) def unwrap_to_op_info( diff --git a/torch/distributed/tensor/_ops/_view_ops.py b/torch/distributed/tensor/_ops/_view_ops.py index e6e1478308e42..1ebc01826526c 100644 --- a/torch/distributed/tensor/_ops/_view_ops.py +++ b/torch/distributed/tensor/_ops/_view_ops.py @@ -465,18 +465,22 @@ def dim_transpose(ndim: int, dim1: int, dim2: int) -> DimMap: return tuple(dimmap) -def dim_squeeze(shape: Shape, dim: int | None = None) -> DimMap: - # FIXME: this is wrong when dim=None and one of the dimensions - # equals size of the mesh. For example squeeze(DTensor(tensor(4), Shard[0])) could - # end up as squeeze(tensor(1)) if we have 4 devices; this would lead to - # removal of a dimension that is not actually a singleton. +def dim_squeeze(shape: Shape, dim: DimsType | None = None) -> DimMap: + # Operates on local shape; sharding_prop rewrites squeeze ops to squeeze.dims + # with only globally-singleton dims before this is called. from torch.fx.experimental.symbolic_shapes import guard_or_true + ndim = len(shape) + if dim is None: + target_dims = set(range(ndim)) + elif isinstance(dim, int): + target_dims = {normalize_dim(dim, ndim)} + else: + target_dims = set(normalize_dims(dim, ndim)) return tuple( InputDim(i) for i, s in enumerate(shape) - if guard_or_true(s > 1) - or (dim is not None and i != normalize_dim(dim, len(shape))) + if guard_or_true(s > 1) or i not in target_dims ) @@ -801,6 +805,7 @@ def reshape_strategy(op_schema: OpSchema) -> StrategyType: register_op_strategy_map(aten.squeeze.default, torch.squeeze) +register_op_strategy_map(aten.squeeze_.default, torch.squeeze) register_op_strategy_map( aten.squeeze_.dim, torch.squeeze, schema_info=RuntimeSchemaInfo(1) ) diff --git a/torch/distributed/tensor/_sharding_prop.py b/torch/distributed/tensor/_sharding_prop.py index a85b19130f69f..6d0040f25b7c9 100644 --- a/torch/distributed/tensor/_sharding_prop.py +++ b/torch/distributed/tensor/_sharding_prop.py @@ -346,6 +346,15 @@ def __init__(self) -> None: aten.select_backward.default: 1, aten.slice_backward.default: 1, } + # squeeze ops that need dim arg rewritten to only globally-singleton dims + self.squeeze_op_to_dims_variant: dict[OpOverload, OpOverload] = { + aten.squeeze.default: aten.squeeze.dims, + aten.squeeze.dim: aten.squeeze.dims, + aten.squeeze.dims: aten.squeeze.dims, + aten.squeeze_.default: aten.squeeze_.dims, + aten.squeeze_.dim: aten.squeeze_.dims, + aten.squeeze_.dims: aten.squeeze_.dims, + } def register_sharding_prop_rule( self, @@ -770,6 +779,15 @@ def propagate_op_sharding_non_cached(self, op_schema: OpSchema) -> OutputShardin needs_redistribute = True use_val_from_redistribute_schema = True + # rewrite squeeze to use only globally-singleton dims + if op_schema.op in self.squeeze_op_to_dims_variant: + schema = suggestion_schema or op_schema + adjusted = self._adjust_squeeze_to_global_singletons(schema) + if adjusted is not None: + suggestion_schema = adjusted + needs_redistribute = True + use_val_from_redistribute_schema = True + # construct output spec for the op if op_schema.return_type_tuple_tensor_like(): # for ops that return multiple tensors and the output_specs is not @@ -957,3 +975,47 @@ def _adjust_shape_and_stride_args( ) return OpSchema(schema.op, tuple(expected_input_schema), schema.kwargs_schema) + + def _adjust_squeeze_to_global_singletons(self, schema: OpSchema) -> OpSchema | None: + """ + Rewrite squeeze ops to squeeze.dims with only globally-singleton dims. + Fixes bug where sharded dims with local size 1 get incorrectly squeezed. + Returns None if no rewrite is needed (already squeeze.dims with correct args). + """ + from torch.fx.experimental.symbolic_shapes import guard_or_false + + input_spec = cast(DTensorSpec, schema.args_schema[0]) + tensor_meta = input_spec.tensor_meta + if tensor_meta is None: + raise RuntimeError("squeeze requires tensor metadata") + global_shape = tensor_meta.shape + ndim = len(global_shape) + + def normalize(d: int) -> int: + return d if d >= 0 else d + ndim + + def is_singleton(d: int) -> bool: + nd = normalize(d) + return 0 <= nd < ndim and guard_or_false(global_shape[nd] == 1) + + # guard_or_false: conservatively keep dims when size is symbolic/unknown + if schema.op in (aten.squeeze.default, aten.squeeze_.default): + target_dims = tuple( + i for i, s in enumerate(global_shape) if guard_or_false(s == 1) + ) + elif schema.op in (aten.squeeze.dim, aten.squeeze_.dim): + dim = normalize(schema.args_schema[1]) # type: ignore[arg-type] + target_dims = (dim,) if is_singleton(dim) else () + else: + dims = cast(Sequence[int], schema.args_schema[1]) + target_dims = tuple( # type: ignore[union-attr] + normalize(d) for d in dims if is_singleton(d) + ) + + dims_variant = self.squeeze_op_to_dims_variant[schema.op] + # Skip rewrite if already targeting the right op with the same dims + if schema.op == dims_variant and len(schema.args_schema) > 1: + existing_dims = schema.args_schema[1] + if existing_dims == target_dims: + return None + return OpSchema(dims_variant, (input_spec, target_dims), {}) From f4c03b2f5fc69248fb4d71e60cefe08c8ff64086 Mon Sep 17 00:00:00 2001 From: dshi7 Date: Mon, 23 Mar 2026 19:22:20 +0000 Subject: [PATCH 0026/1172] Fix store_cache miss for TMA epilogue stores causing NameError (#177990) When epilogue fusion runs with TMA store (mode="tma"), `OpsHandler.store()` skipped `_update_store_cache`. The subsequent epilogue load from the same buffer hit a cache miss, forcing it to emit `tl.load(out_ptr0 + ...)`. `remove_kernel_local_buffers` then stripped out_ptr0 from the kernel signature, causing "NameError: out_ptr0 is not defined" at Triton compile time. The fix is to also update store_cache when mode is 'tma'. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177990 Approved by: https://github.com/eellison --- test/inductor/test_max_autotune.py | 36 ++++++++++++++++++++++++++++++ torch/_inductor/codegen/common.py | 3 ++- 2 files changed, 38 insertions(+), 1 deletion(-) diff --git a/test/inductor/test_max_autotune.py b/test/inductor/test_max_autotune.py index e8c0d3ee12a4a..a1df2eb029cb2 100644 --- a/test/inductor/test_max_autotune.py +++ b/test/inductor/test_max_autotune.py @@ -540,6 +540,41 @@ def mm(a, b): torch.testing.assert_close(c_actual, c_expected, atol=1e-2, rtol=1e-2) + @unittest.skipIf( + not has_triton_tma_device(), "Need device-side TMA support in Triton" + ) + def test_persistent_tma_epilogue_fusion_store_cache(self): + # Regression test: when epilogue fusion runs with TMA store, the + # store_cache must be updated so that a subsequent epilogue load from + # the same buffer hits the cache. Otherwise remove_kernel_local_buffers + # strips the buffer pointer from the kernel signature, causing a + # NameError at Triton compile time. + def f(a, b): + a = a.repeat(8, 8) + b = b.repeat(8, 8) + mm = torch.mm(a, b) + return mm.relu() + + M, N, K = 21, 31, 11 + a = torch.randn(M, K, dtype=torch.float16, device=GPU_TYPE) + b = torch.randn(K, N, dtype=torch.float16, device=GPU_TYPE) + + with config.patch( + { + "max_autotune": True, + "epilogue_fusion": True, + "triton.enable_persistent_tma_matmul": "1", + "triton.native_matmul": False, + "triton.enable_template_tma_store": True, + "triton.disallow_failing_autotune_kernels_TESTING_ONLY": True, + "test_configs.autotune_choice_name_regex": "mm_persistent_tma", + } + ): + actual = torch.compile(f)(a, b) + expected = f(a, b) + + torch.testing.assert_close(actual, expected, atol=1e-2, rtol=1e-2) + @parametrize("dynamic", (False, True)) def test_max_autotune_regular_mm_zero_size_input(self, dynamic: bool): """ @@ -4660,6 +4695,7 @@ class TestMaxAutotuneAsyncPipelined(TestMaxAutotune, TestEpilogueFusionStaticAna "test_non_contiguous_input_mm_plus_mm": "Flaky on trunk", "test_autotune_device_guard": "Flaky on trunk", "test_template_bad_epilogue_fusion": "Benchmarking path is different", + "test_persistent_tma_epilogue_fusion_store_cache": "Epilogue fusion disabled in async pipelining", # Contiguous transform tests - SubgraphChoiceCaller not supported with async pipelining "test_max_autotune_contiguous_transform_mm": "Subgraphs not supported with async pipelining", "test_max_autotune_contiguous_transform_addmm": "Subgraphs not supported with async pipelining", diff --git a/torch/_inductor/codegen/common.py b/torch/_inductor/codegen/common.py index 086e196663e5d..09a1bb1b8cac9 100644 --- a/torch/_inductor/codegen/common.py +++ b/torch/_inductor/codegen/common.py @@ -2851,7 +2851,8 @@ def store( self, name: str, index: sympy.Expr, value: CSEVariable, mode: StoreMode = None ) -> None: self.kernel.store_buffer_names.add(name) - if mode is None: + # Update store cache when mode is None or "tma" + if mode != "atomic_add": self._update_store_cache(name, value) if name not in V.graph.removed_buffers: self.kernel.store(name, index, value, mode=mode) From ce7cd5e2ed922cce6674cf5db2878e326c2666e2 Mon Sep 17 00:00:00 2001 From: Animesh Jain Date: Mon, 23 Mar 2026 09:45:52 -0700 Subject: [PATCH 0027/1172] [dynamo][refactor] Extract vt_identity_compare into object_protocol.py (#178118) Create torch/_dynamo/variables/object_protocol.py, analogous to CPython's Objects/object.c, to hold general PyObject_* algorithm implementations independent of any specific type. Move the identity-comparison logic from BuiltinVariable's inline handle_is closure into vt_identity_compare() in this new module. handle_is now delegates to vt_identity_compare and adjusts the result for is/is-not polarity. Pull Request resolved: https://github.com/pytorch/pytorch/pull/178118 Approved by: https://github.com/guilhermeleobas --- torch/_dynamo/variables/builtin.py | 46 +++------------ torch/_dynamo/variables/object_protocol.py | 65 ++++++++++++++++++++++ 2 files changed, 74 insertions(+), 37 deletions(-) create mode 100644 torch/_dynamo/variables/object_protocol.py diff --git a/torch/_dynamo/variables/builtin.py b/torch/_dynamo/variables/builtin.py index d0fe3c4978758..17e6487046f91 100644 --- a/torch/_dynamo/variables/builtin.py +++ b/torch/_dynamo/variables/builtin.py @@ -850,43 +850,15 @@ def handle_is( left: VariableTracker, right: VariableTracker, ) -> VariableTracker | None: - # VT identity → Python identity - if left is right: - return VariableTracker.build(tx, op.__name__ == "is_") - - # Compare underlying Python objects via hook - left_val = left.get_real_python_backed_value() - right_val = right.get_real_python_backed_value() - - left_known = left_val is not NO_SUCH_SUBOBJ - right_known = right_val is not NO_SUCH_SUBOBJ - - if left_known and right_known: - result = left_val is right_val - return VariableTracker.build( - tx, result if op.__name__ == "is_" else not result - ) - - # One side has a concrete value, the other doesn't — they - # can't be identical (if they were the same object, both - # sides would resolve). - if left_known != right_known: - return VariableTracker.build(tx, op.__name__ != "is_") - - # Mutable containers created during tracing: VT identity - # = Python identity. Already False from `left is right`. - if isinstance(left, (ConstDictVariable, ListVariable)): - return VariableTracker.build(tx, op.__name__ != "is_") - - # Different exception types are never identical - if ( - istype(left, variables.ExceptionVariable) - and istype(right, variables.ExceptionVariable) - and left.exc_type is not right.exc_type - ): - return VariableTracker.build(tx, op.__name__ != "is_") - - return None + from .object_protocol import vt_identity_compare + + result = vt_identity_compare(left, right) + if result is None: + return None + is_same = result.as_python_constant() + return VariableTracker.build( + tx, is_same if op.__name__ == "is_" else not is_same + ) result.append(((VariableTracker, VariableTracker), handle_is)) # type: ignore[arg-type] diff --git a/torch/_dynamo/variables/object_protocol.py b/torch/_dynamo/variables/object_protocol.py new file mode 100644 index 0000000000000..eb6e2bf576c26 --- /dev/null +++ b/torch/_dynamo/variables/object_protocol.py @@ -0,0 +1,65 @@ +""" +Dynamo implementations of CPython's PyObject_* default slot algorithms. + +Analogous to CPython's Objects/object.c, this module holds the general +comparison dispatch machinery that is independent of any specific type. +Per-type richcompare_impl hooks live in their respective VT files. +""" + +from ..utils import istype +from .base import NO_SUCH_SUBOBJ, VariableTracker +from .constant import CONSTANT_VARIABLE_FALSE, CONSTANT_VARIABLE_TRUE + + +def vt_identity_compare( + left: VariableTracker, + right: VariableTracker, +) -> "VariableTracker | None": + """Try to determine Python identity (left is right) at trace time. + + Returns ConstantVariable(True/False) if determinable, else None. + Mirrors the logic in BuiltinVariable's handle_is handler. + """ + if left is right: + return CONSTANT_VARIABLE_TRUE + + left_val = left.get_real_python_backed_value() + right_val = right.get_real_python_backed_value() + left_known = left_val is not NO_SUCH_SUBOBJ + right_known = right_val is not NO_SUCH_SUBOBJ + + if left_known and right_known: + return ( + CONSTANT_VARIABLE_TRUE if left_val is right_val else CONSTANT_VARIABLE_FALSE + ) + + # One side has a concrete backing object, the other doesn't — they can't + # be the same object. + if left_known != right_known: + return CONSTANT_VARIABLE_FALSE + + # Mutable containers created during tracing: VT identity = Python identity. + from .dicts import ConstDictVariable + from .lists import ListVariable + + if isinstance(left, (ConstDictVariable, ListVariable)): + return CONSTANT_VARIABLE_FALSE + + # Different Python types can never be the same object. + try: + if left.python_type() is not right.python_type(): + return CONSTANT_VARIABLE_FALSE + except NotImplementedError: + pass + + # Different exception types are never identical. + from .. import variables + + if ( + istype(left, variables.ExceptionVariable) + and istype(right, variables.ExceptionVariable) + and left.exc_type is not right.exc_type # type: ignore[attr-defined] + ): + return CONSTANT_VARIABLE_FALSE + + return None From e4a5dd60cb72869d7ca9997f2a73179a9d10e640 Mon Sep 17 00:00:00 2001 From: Huy Do Date: Mon, 23 Mar 2026 19:44:40 +0000 Subject: [PATCH 0028/1172] Fix GH / fbcode discrepancy from D96380700 (#178186) This needs to be landed and skipped on diff train. cc @georgehong Pull Request resolved: https://github.com/pytorch/pytorch/pull/178186 Approved by: https://github.com/georgehong, https://github.com/Skylion007 --- buckbuild.bzl | 1 + 1 file changed, 1 insertion(+) diff --git a/buckbuild.bzl b/buckbuild.bzl index 7b1ad284d51cc..fa7b7ba879a90 100644 --- a/buckbuild.bzl +++ b/buckbuild.bzl @@ -905,6 +905,7 @@ def define_buck_targets( ("aten/src", "ATen/ops/*.h"), # ATen Base ("aten/src", "ATen/*.h"), + ("aten/src", "ATen/accelerator/*.h"), ("aten/src", "ATen/cpu/**/*.h"), ("aten/src", "ATen/detail/*.h"), ("aten/src", "ATen/functorch/**/*.h"), From 0ec1530b6e6d694f06babef32205928504e8f8ae Mon Sep 17 00:00:00 2001 From: Fadi Arafeh Date: Mon, 23 Mar 2026 17:02:22 +0000 Subject: [PATCH 0029/1172] [CPU][Inductor] Use VecMask::from for scalar masks in codegen (#178148) Fixes: #178136, https://github.com/vllm-project/vllm/issues/37325 Signed-off-by: Fadi Arafeh Pull Request resolved: https://github.com/pytorch/pytorch/pull/178148 Approved by: https://github.com/malfet, https://github.com/Skylion007 Co-authored-by: Nikita Shulga <2453524+malfet@users.noreply.github.com> --- test/inductor/test_cpu_repro.py | 20 ++++++++++++++++++++ torch/_inductor/codegen/cpp.py | 2 +- 2 files changed, 21 insertions(+), 1 deletion(-) diff --git a/test/inductor/test_cpu_repro.py b/test/inductor/test_cpu_repro.py index 3e64f653f00c2..3b8146cdfd0cb 100644 --- a/test/inductor/test_cpu_repro.py +++ b/test/inductor/test_cpu_repro.py @@ -5077,6 +5077,26 @@ def fn(x): "at::vec::VectorizedN::loadu", 2, exactly=True ).run(code) + @requires_vectorization + def test_indirect_assert_scalar_mask_tail_vec_no_crash(self): + # https://github.com/pytorch/pytorch/issues/178136 + def fn(positions, cache): + x = cache[positions] + y = x[0].clone() + y[..., 1::3] = x[1, ..., 1::3] + y[..., 2::3] = x[2, ..., 2::3] + return y + + positions = torch.tensor([[0, 0], [0, 0], [0, 0]], dtype=torch.int64) + cache = torch.arange(3, dtype=torch.float32).reshape(1, 3) + + with config.patch({"cpp.enable_loop_tail_vec": True}): + expected = fn(positions, cache) + compiled_fn = torch.compile(fn, fullgraph=True) + actual = compiled_fn(positions, cache) + + torch.testing.assert_close(actual, expected) + def test_uint64_pointwise_vec(self): def fn(x): return x * x diff --git a/torch/_inductor/codegen/cpp.py b/torch/_inductor/codegen/cpp.py index 0481c5e662a9e..59ab41efd5dbd 100644 --- a/torch/_inductor/codegen/cpp.py +++ b/torch/_inductor/codegen/cpp.py @@ -3503,7 +3503,7 @@ def indirect_assert(self, var, lower, upper, mask=None): cond = f"{self._get_mask_type(var.dtype)}({cond})" if mask: if not mask.is_vec: - mask = f"{self._get_mask_type(var.dtype)}({mask})" + mask = f"{self._get_mask_type(var.dtype)}::from({mask})" # We need not check when the mask is False cond = f"({cond}) | ~({mask})" if self.tail_size: From 8198cb4aa6816d15a80272f60eaebdbe50e8eee9 Mon Sep 17 00:00:00 2001 From: Phillip Liu Date: Mon, 23 Mar 2026 20:08:45 +0000 Subject: [PATCH 0030/1172] [FR script] Fixed bug which caused FR script to fail in the case of a coalsced collective not scheduled (#177076) (#177076) Summary: FR script assumed that the coalsced collective is always scheduled from all ranks in the pg, but in the case of a subset of ranks didn't schedule a coalsced collective, the FR script would try to remove the coalsced collective in its clean-up stage and fail abruptly. This diff fixed the bug to skip cleaning-up the non-scheduled coalseced collective. Test Plan: `buck2 run //investigations/dr_patternson/analyzers/ai_observability:ai_observability-all-analyzers-cli -- flight_recorder_analyzer --mast_job_name fire-xiaoxuanw-f1047233868 --mast_job_version 0 --mast_job_attempt 2` Confirm no error, confirm the SBDive view has insights https://fburl.com/ai_infra/ppmvoeb8 Reviewed By: fduwjj, YongzhongYang Differential Revision: D96016690 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177076 Approved by: https://github.com/fduwjj --- torch/distributed/flight_recorder/components/builder.py | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/torch/distributed/flight_recorder/components/builder.py b/torch/distributed/flight_recorder/components/builder.py index 87c52d6d4b090..663363e718257 100644 --- a/torch/distributed/flight_recorder/components/builder.py +++ b/torch/distributed/flight_recorder/components/builder.py @@ -311,7 +311,12 @@ def build_collectives( # This extra cleanup is needed because we need to pop all collectives within a coalesced collective. for i, k in idx_map.items(): for _ in range(1, num_coalesced_entries): - all_entries[i].pop(k) + try: + all_entries[i].pop(k) + except IndexError: + # In the case of a missing rank symptom that a rank didn't schedule the coalesced collective, + # we should not fail the analysis script here. + pass else: # Iterate through all the ranks and check if there is a mismatch for the current entry. check_current_entry_match( From cefbd65695576c434c67150c79e86fddeeedea72 Mon Sep 17 00:00:00 2001 From: Divyansh Khanna Date: Mon, 23 Mar 2026 20:11:49 +0000 Subject: [PATCH 0031/1172] Guidelines for testing changes to memory viz (#177551) This should help both engineers and agents for when folks want to add new features to the memory visualizer tool. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177551 Approved by: https://github.com/scotts --- torch/utils/viz/MemoryViz.js | 58 ++++++++++++++++++++++++++++++++++++ 1 file changed, 58 insertions(+) diff --git a/torch/utils/viz/MemoryViz.js b/torch/utils/viz/MemoryViz.js index c0ef85a65d2f3..3c516cf30a40c 100644 --- a/torch/utils/viz/MemoryViz.js +++ b/torch/utils/viz/MemoryViz.js @@ -1,3 +1,61 @@ +/** + * ================================================================================ + * MemoryViz.js - PyTorch Memory Visualization Tool + * ================================================================================ + * + * OVERVIEW: + * --------- + * This file contains the core visualization logic for PyTorch's memory profiler. + * It renders memory allocation timelines, stack traces, and provides interactive + * exploration of memory snapshots captured during model execution. + * + * KEY FEATURES: + * - Multiple visualization tabs/views for different memory analysis perspectives + * - Interactive stack trace display (supports both click and hover modes) + * - Zoom and brush controls for navigating large memory timelines + * - Support for loading memory snapshot files (.pickle format) + * + * ================================================================================ + * TESTING INSTRUCTIONS FOR ENGINEERS & AGENTS + * ================================================================================ + * + * 1. LOCAL TESTING SETUP: + * - Create a simple HTML file that references this JS file: + * + * + * + * MemoryViz Test + * + * + * + * + * + * - Serve locally using: python3 -m http.server 8888 + * - Open http://localhost:8888 in your browser + * + * 2. WHAT TO TEST: + * - Ensure ALL tabs/views render correctly and switch properly + * - Verify BOTH interaction modes work: + * * Click mode: stack traces appear on click + * * Hover mode: stack traces appear on mouseover + * - Test zoom and brush controls for timeline navigation + * - Verify memory allocation blocks are rendered and interactive + * + * 3. TEST DATA REQUIREMENTS: + * - DO NOT just test with small dummy .pickle files + * - Use realistic, decent-sized .pickle files (10-100+ MB range) + * - Large files stress-test rendering performance and memory handling + * - Test with snapshots from real model training/inference runs + * + * 4. COMMON ISSUES TO WATCH FOR: + * - Performance degradation with large snapshots + * - Stack trace popups not appearing or positioning incorrectly + * - Tab switching not updating the visualization properly + * - Zoom/brush state not persisting across interactions + * + * ================================================================================ + */ + 'use strict'; import * as d3 from "https://cdn.jsdelivr.net/npm/d3@7/+esm"; From 09e8f86e7a675fcb3e80ecbb674eac655d3a7b28 Mon Sep 17 00:00:00 2001 From: NikhilAPatel Date: Tue, 17 Mar 2026 14:05:50 -0700 Subject: [PATCH 0032/1172] [CI] Fix periodic inductor CI silently skipping all tests (#177695) Should fix https://github.com/pytorch/pytorch/issues/177084 The daily periodic schedule for inductor-unittest has been running without `mem_leak_check` mode for CUDA jobs. [PR #161536](https://github.com/pytorch/pytorch/pull/161536) renamed the build job from `cuda12.8-py3.10-gcc9-sm86` to `inductor-build`, which removed "cuda" from the job name. `is_cuda_or_rocm_job()` in filter_test_configs.py only checked the job name string, so it started returning False for the CUDA inductor job. Without `mem_leak_check`, the only periodic mode was `rerun_disabled_tests`, which by design skips all non-disabled tests. Since few inductor tests are disabled, nearly all tests were skipped on every periodic run for ~6.5 months. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177695 Approved by: https://github.com/Lucaskabela --- .github/scripts/filter_test_configs.py | 28 +++++++++++++++------ .github/scripts/test_filter_test_configs.py | 20 ++++++++++++++- 2 files changed, 39 insertions(+), 9 deletions(-) diff --git a/.github/scripts/filter_test_configs.py b/.github/scripts/filter_test_configs.py index 5393f50673ef7..9f7f7a3a7bedc 100755 --- a/.github/scripts/filter_test_configs.py +++ b/.github/scripts/filter_test_configs.py @@ -30,19 +30,31 @@ logging.basicConfig(level=logging.INFO) -def is_cuda_or_rocm_job(job_name: str | None) -> bool: - if not job_name: - return False +def is_cuda_or_rocm_job( + job_name: str | None, config: dict[str, Any] | None = None +) -> bool: + if job_name and ("cuda" in job_name or "rocm" in job_name): + return True - return "cuda" in job_name or "rocm" in job_name + # Also check the runner name in the config, since some workflows (e.g. + # inductor-unittest) use job names that don't include "cuda" even though + # they target CUDA runners. + if config: + runner = config.get("runner", "") + if "nvidia.gpu" in runner or "rocm.gpu" in runner: + return True + + return False # Supported modes when running periodically. Only applying the mode when -# its lambda condition returns true -SUPPORTED_PERIODICAL_MODES: dict[str, Callable[[str | None], bool]] = { +# its lambda condition returns true. Each callable receives (job_name, config). +SUPPORTED_PERIODICAL_MODES: dict[ + str, Callable[[str | None, dict[str, Any] | None], bool] +] = { # Memory leak check is only needed for CUDA and ROCm jobs which utilize GPU memory "mem_leak_check": is_cuda_or_rocm_job, - "rerun_disabled_tests": lambda job_name: True, + "rerun_disabled_tests": lambda job_name, config=None: True, } # The link to the published list of disabled jobs @@ -225,7 +237,7 @@ def set_periodic_modes( for config in test_matrix.get("include", []): for mode, cond in SUPPORTED_PERIODICAL_MODES.items(): - if not cond(job_name): + if not cond(job_name, config): continue cfg = config.copy() diff --git a/.github/scripts/test_filter_test_configs.py b/.github/scripts/test_filter_test_configs.py index 26e38828b7865..1463f82f61eda 100755 --- a/.github/scripts/test_filter_test_configs.py +++ b/.github/scripts/test_filter_test_configs.py @@ -376,13 +376,31 @@ def test_set_periodic_modes(self) -> None: scheduled_test_matrix = set_periodic_modes(test_matrix, job_name) expected_modes = [ - m for m, c in SUPPORTED_PERIODICAL_MODES.items() if c(job_name) + m for m, c in SUPPORTED_PERIODICAL_MODES.items() if c(job_name, None) ] self.assertEqual( len(test_matrix["include"]) * len(expected_modes), len(scheduled_test_matrix["include"]), ) + def test_set_periodic_modes_gpu_runner(self) -> None: + """Job name without 'cuda' but runner indicates a CUDA job (e.g. inductor-unittest).""" + test_matrix = yaml.safe_load( + "{include: [" + '{config: "inductor", shard: 1, num_shards: 2, runner: "linux.g5.4xlarge.nvidia.gpu"}, ' + '{config: "inductor", shard: 2, num_shards: 2, runner: "linux.g5.4xlarge.nvidia.gpu"}' + "]}" + ) + scheduled = set_periodic_modes(test_matrix, "inductor-build / build") + + modes_per_config = [ + entry.get("mem_leak_check") or entry.get("rerun_disabled_tests") + for entry in scheduled["include"] + ] + self.assertIn("mem_leak_check", modes_per_config) + self.assertIn("rerun_disabled_tests", modes_per_config) + self.assertEqual(len(scheduled["include"]), 4) + @mock.patch("filter_test_configs.download_json") def test_remove_disabled_jobs(self, mock_download_json: Any) -> None: mock_download_json.return_value = MOCKED_DISABLED_UNSTABLE_JOBS From 357671c76c945c210f82df6c9ba74498b5353f5e Mon Sep 17 00:00:00 2001 From: Masato Shinokawa Date: Mon, 23 Mar 2026 20:30:52 +0000 Subject: [PATCH 0033/1172] Skip quantile q-value validation for tensor subclass-like inputs (#174859) Fixes #173659. `torch.quantile` fails during `torch.export.export` with `DataDependentOutputException: aten.equal.default`. The validation check in `quantile_compute` uses `is_scalar_tensor_true`, which involves a data-dependent op that cannot be evaluated on FakeTensors. The fix guards the validation with `!isTensorSubclassLike(self)`, consistent with the existing pattern used elsewhere in the same function. The validation still runs in normal eager mode. We also guard the validation `.numel() > 0` likewise, which do not work with symbolic shapes. Test ``` python test/export/test_export.py TestExport.test_quantile_export ``` Pull Request resolved: https://github.com/pytorch/pytorch/pull/174859 Approved by: https://github.com/laithsakka, https://github.com/Lucaskabela Co-authored-by: PyTorch MergeBot --- aten/src/ATen/native/Sorting.cpp | 34 ++++++++++--------- test/distributed/tensor/test_dtensor_ops.py | 1 + test/export/test_export.py | 19 +++++++++++ test/functorch/test_aotdispatch.py | 2 -- test/functorch/test_ops.py | 8 ----- test/test_proxy_tensor.py | 4 --- .../testing/_internal/common_ops_unbacked.py | 2 -- 7 files changed, 38 insertions(+), 32 deletions(-) diff --git a/aten/src/ATen/native/Sorting.cpp b/aten/src/ATen/native/Sorting.cpp index eb3c2c93ec8a0..9c1d6c4c4e0bc 100644 --- a/aten/src/ATen/native/Sorting.cpp +++ b/aten/src/ATen/native/Sorting.cpp @@ -206,7 +206,7 @@ QUANTILE_INTERPOLATION_MODE get_quantile_interpolation_mode( } void quantile_checks(const Tensor& self, const Tensor& q) { - TORCH_CHECK(self.numel() > 0, "quantile() input tensor must be non-empty"); + TORCH_SYM_CHECK(self.sym_numel().sym_gt(0), "quantile() input tensor must be non-empty"); TORCH_CHECK(q.dim() <= 1, "quantile() q must be a scalar or 1D tensor"); TORCH_CHECK( self.scalar_type() == kFloat || self.scalar_type() == kDouble, @@ -219,26 +219,26 @@ void quantile_checks(const Tensor& self, const Tensor& q) { "quantile() q tensor must be on the same device as the input tensor"); } -std::vector quantile_output_shape( +std::vector quantile_output_shape( const std::optional original_dim, const Tensor& self, const Tensor& q, const bool keepdim, int64_t wrapped_dim) { // Compute output shape: q_size + reduced_size - std::vector out_shape; + std::vector out_shape; if (original_dim && self.dim() > 0) { - out_shape = self.sizes().vec(); + out_shape = self.sym_sizes().vec(); if (keepdim) { out_shape[wrapped_dim] = 1; } else { out_shape.erase(out_shape.begin() + wrapped_dim); } } else if (keepdim) { - out_shape = std::vector(self.dim(), 1); + out_shape = std::vector(self.dim(), 1); } if (q.dim() > 0) { - out_shape.insert(out_shape.begin(), q.numel()); + out_shape.insert(out_shape.begin(), q.sym_numel()); } return out_shape; @@ -252,11 +252,13 @@ Tensor quantile_compute( const QUANTILE_INTERPOLATION_MODE& interpolation, const bool ignore_nan, int64_t wrapped_dim, - std::vector out_shape) { + std::vector out_shape) { // Checks that all q values are between 0 and 1, inclusive // NOTE: this check is only performed when running on the CPU to avoid // synchronizing an accelerator with the CPU - if (self.device().is_cpu()) { + // The check is also skipped when the actual q values are not available yet + // e.g. with symbolic shapes or during export + if (self.device().is_cpu() && !isTensorSubclassLike(q)) { auto all_q_in_range = q.ge(0).logical_and_(q.le(1)).all(); TORCH_CHECK(at::is_scalar_tensor_true(all_q_in_range), "quantile() q values must be in the range [0, 1]"); @@ -275,18 +277,18 @@ Tensor quantile_compute( // Treat q as a 1D tensor for the following computations if (q.dim() == 0) { - out_shape.insert(out_shape.begin(), q.numel()); + out_shape.insert(out_shape.begin(), 1); } // View input as reduced_size + size of dim to reduce - std::vector in_shape(out_shape.size()); + std::vector in_shape(out_shape.size()); std::copy(out_shape.begin() + 1, out_shape.end(), in_shape.begin()); - in_shape[in_shape.size() - 1] = sorted.size(-1); - sorted = sorted.view(in_shape); + in_shape[in_shape.size() - 1] = sorted.sym_size(-1); + sorted = sorted.view_symint(in_shape); // Ensure converting from int64_t to double won't overflow - TORCH_CHECK( - sorted.size(-1) <= std::pow(2, 24), + TORCH_SYM_CHECK( + sorted.sym_size(-1).sym_le(1 << 24), "quantile() input tensor is too large"); // Convert q in [0, 1] to ranks in [0, reduction_size) @@ -308,7 +310,7 @@ Tensor quantile_compute( } else { // For quantile, compute ranks based on reduction size. If there is nan // set rank to last index so the quantile computed will be nan. - int64_t last_index = sorted.size(-1) - 1; + auto last_index = sorted.sym_size(-1) - 1; std::vector tl = at::broadcast_tensors({q * last_index, sorted.isnan().any(-1, true)}); ranks = at::masked_fill(tl[0], tl[1], last_index); @@ -388,7 +390,7 @@ void quantile_out_impl( int64_t wrapped_dim = at::maybe_wrap_dim(original_dim.value_or(0), self.dim()); auto out_shape = quantile_output_shape(original_dim, self, q, keepdim, wrapped_dim); - resize_output(out, out_shape); + resize_output_symint(out, out_shape); auto quantile = quantile_compute( self, q, original_dim, keepdim, interpolation, ignore_nan, wrapped_dim, std::move(out_shape)); diff --git a/test/distributed/tensor/test_dtensor_ops.py b/test/distributed/tensor/test_dtensor_ops.py index 65145b37c42ff..4ae34dd938ec8 100644 --- a/test/distributed/tensor/test_dtensor_ops.py +++ b/test/distributed/tensor/test_dtensor_ops.py @@ -894,6 +894,7 @@ def assertEqualOnRank(self, x, y, msg=None, *, rank=0): xfail("nonzero_static"), xfail("permute_copy"), xfail("prod"), + xfail("quantile"), xfail("ravel"), xfail("reshape"), xfail("reshape_as"), diff --git a/test/export/test_export.py b/test/export/test_export.py index bcce9c1836b30..9c2fe266fd124 100755 --- a/test/export/test_export.py +++ b/test/export/test_export.py @@ -17570,6 +17570,25 @@ def forward(self, x, y): expected_mask = torch.ones(3, 5, dtype=torch.bool).triu(diagonal=2) self.assertEqual(eager_out, expected_mask) + def test_quantile_export(self): + class QuantilePair(torch.nn.Module): + def __init__(self, noise=0.1): + super().__init__() + self.noise = noise + + def forward(self, x): + q = torch.tensor( + [self.noise, 1.0 - self.noise], + device=x.device, + dtype=x.dtype, + ) + return torch.quantile(x, q, dim=-1, keepdim=True) + + model = QuantilePair(noise=0.1) + x = torch.randn(1, 3200) + ep = export(model, (x,)) + self.assertEqual(ep.module()(x), model(x)) + @unittest.skipIf(not torchdynamo.is_dynamo_supported(), "dynamo isn't support") class TestOneOffModelExportResult(TestCase): diff --git a/test/functorch/test_aotdispatch.py b/test/functorch/test_aotdispatch.py index 86100c8d54b91..c771ab1b77185 100644 --- a/test/functorch/test_aotdispatch.py +++ b/test/functorch/test_aotdispatch.py @@ -8745,8 +8745,6 @@ def fn(x): xfail("nn.functional.gaussian_nll_loss"), xfail("tensor_split"), xfail("corrcoef"), - xfail("quantile"), - xfail("nanquantile"), skip("narrow"), xfail("istft"), xfail("linalg.eig"), diff --git a/test/functorch/test_ops.py b/test/functorch/test_ops.py index 955b1a8ada144..632a4a915a685 100644 --- a/test/functorch/test_ops.py +++ b/test/functorch/test_ops.py @@ -981,9 +981,6 @@ def fn(inp, *args, **kwargs): "masked.softmax", device_type="cpu", ), - xfail( - "nanquantile", device_type="cpu" - ), # vmap not implemented for at::equal. xfail("native_layer_norm"), # vmap: inplace into a regular tensor # got a batched tensor as input while the running_mean or running_var, # which will be updated in place, were not batched. @@ -1035,9 +1032,6 @@ def fn(inp, *args, **kwargs): xfail("normal"), # calls random op xfail("normal", "number_mean"), # calls random op xfail("pca_lowrank"), # calls random op - xfail( - "quantile", device_type="cpu" - ), # Batching rule not implemented for `at::equal` xfail( "scatter_reduce", "prod" ), # vmap (looks like you are calling item/data-dependent) @@ -1180,10 +1174,8 @@ def vjp_of_vjp(*args_and_cotangents): # TODO: implement batching rule skip("_batch_norm_with_update"), xfail("__getitem__", ""), # dynamic error - xfail("nanquantile", device_type="cpu"), # checks q via a .item() call xfail("nn.functional.gaussian_nll_loss"), # checks var for if any value < 0 xfail("narrow"), # .item() call - xfail("quantile", device_type="cpu"), # checks q via a .item() call xfail("view_as_complex"), # Tensor must have a last dimension with stride 1 # required rank 4 tensor to use channels_last format xfail("bfloat16"), diff --git a/test/test_proxy_tensor.py b/test/test_proxy_tensor.py index 8296f386f0977..171c13bbe3430 100644 --- a/test/test_proxy_tensor.py +++ b/test/test_proxy_tensor.py @@ -2029,8 +2029,6 @@ def f(t): xfail('cov'), xfail('nn.functional.gaussian_nll_loss'), xfail('corrcoef'), - xfail('quantile'), - xfail('nanquantile'), # Seems like it's creating a sparse tensor that isn't captured by tensor.is_sparse xfail('sparse.sampled_addmm'), @@ -2061,11 +2059,9 @@ def f(t): xfail('geqrf', ''), # aten.geqrf.default - couldn't find symbolic meta function/decomposition xfail('histogram', ''), # Could not run 'aten::histogram.bin_ct' with arguments from the 'Meta' backend. This c... xfail('histogramdd', ''), # aten._histogramdd_bin_edges.default - couldn't find symbolic meta function/decomposition - xfail('nanquantile', ''), # Could not run 'aten::equal' with arguments from the 'Meta' backend. xfail('nn.functional.binary_cross_entropy', ''), # aten.new_empty.default - couldn't find symbolic meta function/decom... xfail('nn.functional.cross_entropy', ''), # aten.size.default - couldn't find symbolic meta function/decomposition xfail('nn.functional.ctc_loss'), # aten._ctc_loss.Tensor - couldn't find symbolic meta function/decomposition - xfail('quantile', ''), # Could not run 'aten::equal' with arguments from the 'Meta' backend. xfail('max_pool2d_with_indices_backward', ''), # Expected a value of type 'List[int]' for argument 'kernel_size' but... } diff --git a/torch/testing/_internal/common_ops_unbacked.py b/torch/testing/_internal/common_ops_unbacked.py index 56a47a263cfc4..c49e237501257 100644 --- a/torch/testing/_internal/common_ops_unbacked.py +++ b/torch/testing/_internal/common_ops_unbacked.py @@ -124,7 +124,6 @@ def skip(op_name, variant_name="", *, device_type=None, dtypes=None): xfail("masked.var"), xfail("max_pool2d_with_indices_backward"), xfail("multinomial"), - xfail("nanquantile"), xfail("nn.functional.adaptive_avg_pool1d"), xfail("nn.functional.adaptive_avg_pool2d"), xfail("nn.functional.adaptive_avg_pool3d"), @@ -196,7 +195,6 @@ def skip(op_name, variant_name="", *, device_type=None, dtypes=None): xfail("ormqr"), xfail("pca_lowrank"), xfail("pinverse"), - xfail("quantile"), xfail("qr"), xfail("rand_like"), xfail("randint_like"), From 912584c3934fb0d377e14df4c94db54bd97091e5 Mon Sep 17 00:00:00 2001 From: ygshen-ai Date: Mon, 23 Mar 2026 20:31:57 +0000 Subject: [PATCH 0034/1172] fixed:torch.compile Bad import result for types.ModuleType (#177824) Fixes #177682 torch.compile: Bad import result for types.ModuleType subclass in sys.modules. user define "_ConfigModule" is a subclass of types.ModuleType so change from "istype(value, types.ModuleType)" to "isinstance(value, types.ModuleType)" Pull Request resolved: https://github.com/pytorch/pytorch/pull/177824 Approved by: https://github.com/Lucaskabela Co-authored-by: PyTorch MergeBot --- test/dynamo/test_misc.py | 20 ++++++++++++++++++++ torch/_dynamo/symbolic_convert.py | 2 +- 2 files changed, 21 insertions(+), 1 deletion(-) diff --git a/test/dynamo/test_misc.py b/test/dynamo/test_misc.py index a94c1d27b7471..94532e007f8d2 100644 --- a/test/dynamo/test_misc.py +++ b/test/dynamo/test_misc.py @@ -14840,6 +14840,26 @@ def fn(x): self.assertEqual(ref, result) self.assertEqual(saved_ref, saved["grad"]) + def test_import_user_defined_module(self): + # testcase for https://github.com/pytorch/pytorch/issues/177682 + # Bad import result for types.ModuleType subclass in sys.modules + class _ConfigModule(types.ModuleType): + x = 1 + + _ConfigModule.__module__ = __name__ + sys.modules["my_config"] = _ConfigModule("my_config") + + def fn(): + import my_config # noqa: F401 + + return torch.tensor(1) + + compilefn = torch.compile(fn, fullgraph=True, backend="eager") + + ret1 = fn() + ret2 = compilefn() + self.assertEqual(ret1, ret2) + class MiscTestsPyTree(torch._inductor.test_case.TestCase): @parametrize_pytree_module diff --git a/torch/_dynamo/symbolic_convert.py b/torch/_dynamo/symbolic_convert.py index 2ae4fec59c2ab..e7a8a50ff7d18 100644 --- a/torch/_dynamo/symbolic_convert.py +++ b/torch/_dynamo/symbolic_convert.py @@ -2030,7 +2030,7 @@ def IMPORT_NAME(self, inst: Instruction) -> None: self.exec_recorder.add_local_mod(recorded_name, value) # pyrefly: ignore [unbound-name] - if istype(value, (types.ModuleType, DummyModule)): + if isinstance(value, (types.ModuleType, DummyModule)): # pyrefly: ignore [unbound-name, bad-argument-type] self.push(PythonModuleVariable(value, source=source)) else: From e07fa41a2c048d9b8020fbb9401d28ab79868b78 Mon Sep 17 00:00:00 2001 From: Jane Xu Date: Fri, 20 Mar 2026 15:17:38 -0700 Subject: [PATCH 0035/1172] Improve docs on headeronly DISPATCH macros for our stable ABI (#177996) Existing doc referenced a nonexistent macro. This PR takes the opportunity to make it better. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177996 Approved by: https://github.com/mikaylagawarecki --- docs/cpp/source/stable.rst | 84 +++++++++++++++++++++++++++++++++++--- 1 file changed, 79 insertions(+), 5 deletions(-) diff --git a/docs/cpp/source/stable.rst b/docs/cpp/source/stable.rst index c7c60995419da..4ae32f33f1892 100644 --- a/docs/cpp/source/stable.rst +++ b/docs/cpp/source/stable.rst @@ -289,16 +289,90 @@ Dispatch Macros ^^^^^^^^^^^^^^^ Header-only dispatch macros (THO = Torch Header Only) are available for -dtype and device dispatching: +dtype dispatching: .. code-block:: cpp + #include + + THO_DISPATCH_V2( + tensor.scalar_type(), // will be resolved as scalar_t + "my_kernel", + AT_WRAP(([&]() { + // code to specialize with scalar_t + // scalar_t is the resolved C++ type (e.g. float, double) + auto* data = static_cast(tensor.mutable_data_ptr()); + Scalar s(*data); + })), + AT_EXPAND(AT_ALL_TYPES), + AT_EXPAND(AT_COMPLEX_TYPES), + torch::headeronly::ScalarType::Half, + // as many type arguments as needed + ); + +``THO_DISPATCH_V2`` works the same way as ``AT_DISPATCH_V2`` (see +``ATen/Dispatch_v2.h``) but does not require linking against libtorch. +As a result, whereas ``AT_DISPATCH_V2`` would have thrown ``c10::NotImplementedError`` +for unimplemented paths, ``THO_DISPATCH_V2`` will throw ``std::runtime_error``. + +For ease of use, we've also migrated the below AT_* macros representing +collections of types to be header-only and thus have no dependency on libtorch: + +- ``AT_FLOATING_TYPES`` +- ``AT_INTEGRAL_TYPES`` +- ``AT_INTEGRAL_TYPES_V2`` +- ``AT_ALL_TYPES`` +- ``AT_COMPLEX_TYPES`` +- ``AT_ALL_TYPES_AND_COMPLEX`` +- ``AT_FLOAT8_TYPES`` +- ``AT_BAREBONES_UNSIGNED_TYPES`` +- ``AT_QINT_TYPES`` + +If your extension uses our older AT_DISPATCH version 1 infrastructure, +you can also migrate to a header-only libtorch-free world without upgrading +everything to version 2. + +``THO_DISPATCH_SWITCH`` and ``THO_DISPATCH_CASE`` are the header-only +equivalents of ``AT_DISPATCH_SWITCH`` and ``AT_DISPATCH_CASE``. Similarly, +the only user-visible difference is the exception type on an unhandled dtype, +where the ``AT_`` version throws a ``c10::NotImplementedError`` and the ``THO_`` +version throws a ``std::runtime_error``. + +The migration is pretty mechanical: + +- ``AT_DISPATCH_SWITCH`` → ``THO_DISPATCH_SWITCH`` +- ``AT_DISPATCH_CASE`` → ``THO_DISPATCH_CASE`` +- ``AT_PRIVATE_CASE_TYPE_USING_HINT`` → ``THO_PRIVATE_CASE_TYPE_USING_HINT`` +- ``at::ScalarType::X`` → ``torch::headeronly::ScalarType::X`` + +.. code-block:: cpp + + // ---- Before (requires linking against libtorch) ---- + #include + + #define MY_DISPATCH_CASE_FLOATING_TYPES(...) \ + AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \ + AT_DISPATCH_CASE(at::ScalarType::Half, __VA_ARGS__) \ + AT_DISPATCH_CASE(at::ScalarType::BFloat16, __VA_ARGS__) + + #define MY_DISPATCH_FLOATING_TYPES(TYPE, NAME, ...) \ + AT_DISPATCH_SWITCH(TYPE, NAME, \ + MY_DISPATCH_CASE_FLOATING_TYPES(__VA_ARGS__)) + +.. code-block:: cpp + + // ---- After (header-only, no libtorch dependency) ---- #include - THO_DISPATCH_FLOATING_TYPES(tensor.scalar_type(), "my_kernel", [&] { - // scalar_t is the resolved type - auto* data = tensor.data_ptr(); - }); + #define MY_DISPATCH_CASE_FLOATING_TYPES(...) \ + THO_DISPATCH_CASE(torch::headeronly::ScalarType::Float, __VA_ARGS__) \ + THO_DISPATCH_CASE(torch::headeronly::ScalarType::Half, __VA_ARGS__) \ + THO_DISPATCH_CASE(torch::headeronly::ScalarType::BFloat16, __VA_ARGS__) + + #define MY_DISPATCH_FLOATING_TYPES(TYPE, NAME, ...) \ + THO_DISPATCH_SWITCH(TYPE, NAME, \ + MY_DISPATCH_CASE_FLOATING_TYPES(__VA_ARGS__)) + Full API List ^^^^^^^^^^^^^ From 7daf9eb4ce9f0f02ffab29246cfae1032f2a4246 Mon Sep 17 00:00:00 2001 From: Jane Xu Date: Mon, 23 Mar 2026 11:35:41 -0700 Subject: [PATCH 0036/1172] More improvements on docs, e.g., on STD_TORCH_CHECK (#177997) Pull Request resolved: https://github.com/pytorch/pytorch/pull/177997 Approved by: https://github.com/mikaylagawarecki ghstack dependencies: #177996 --- docs/cpp/source/stable.rst | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/docs/cpp/source/stable.rst b/docs/cpp/source/stable.rst index 4ae32f33f1892..ef7ac80428257 100644 --- a/docs/cpp/source/stable.rst +++ b/docs/cpp/source/stable.rst @@ -227,8 +227,9 @@ Header-Only Utilities --------------------- The ``torch::headeronly`` namespace provides header-only versions of common -PyTorch types and utilities. These can be used without linking against libtorch, -making them ideal for maintaining binary compatibility across PyTorch versions. +PyTorch types and utilities. These can be used without linking against libtorch +at all! This portability makes them ideal for maintaining binary compatibility +across PyTorch versions. Error Checking ^^^^^^^^^^^^^^ @@ -241,6 +242,11 @@ Error Checking STD_TORCH_CHECK(condition, "Error message with ", variable, " interpolation"); +Wherever you used ``TORCH_CHECK`` before, you can replace usage with ``STD_TORCH_CHECK`` +to remove the need to link against libtorch. The only difference is that when the +condition check fails, ``TORCH_CHECK`` throws a fancier ``c10::Error`` while +``STD_TORCH_CHECK`` throws a ``std::runtime_error``. + Core Types ^^^^^^^^^^ From dbd870fcde319b9f9e00970c06454dfbaffa176c Mon Sep 17 00:00:00 2001 From: Kurt Mohler Date: Mon, 23 Mar 2026 12:43:47 -0500 Subject: [PATCH 0037/1172] [MPS] Support complex inputs to `scatter` and `gather` (#177794) Pull Request resolved: https://github.com/pytorch/pytorch/pull/177794 Approved by: https://github.com/malfet, https://github.com/Skylion007 --- .../native/mps/operations/ScatterGather.mm | 32 +++++++++++++++++-- .../_internal/common_methods_invocations.py | 16 ---------- torch/testing/_internal/common_mps.py | 4 +++ 3 files changed, 34 insertions(+), 18 deletions(-) diff --git a/aten/src/ATen/native/mps/operations/ScatterGather.mm b/aten/src/ATen/native/mps/operations/ScatterGather.mm index ce65421c71c9d..b9b055b4fee6f 100644 --- a/aten/src/ATen/native/mps/operations/ScatterGather.mm +++ b/aten/src/ATen/native/mps/operations/ScatterGather.mm @@ -10,10 +10,23 @@ #include #include #include +#include #endif namespace at::native { +static Tensor maybe_expand_0_dim(const Tensor& t) { + return t.dim() == 0 ? t.view({1}) : t; +} + +static Tensor expand_index_as_real(const Tensor& index) { + auto index_view = maybe_expand_0_dim(index); + std::vector index_expanded_sizes = index_view.sizes().vec(); + index_expanded_sizes.push_back(2); + auto index_expanded = index_view.unsqueeze(-1).expand(index_expanded_sizes); + return index_expanded; +} + TORCH_IMPL_FUNC(gather_out_mps) (const Tensor& self_arg, int64_t dim, const Tensor& index, bool sparse_grad, const Tensor& output) { using namespace mps; @@ -27,7 +40,14 @@ TORCH_CHECK(!sparse_grad, "sparse_grad not supported in MPS yet") TORCH_CHECK(self.scalar_type() == output.scalar_type(), "gather(): self and output must have the same scalar type"); TORCH_CHECK(dim >= 0 && dim < self.dim(), "gather(): Indexing dim ", dim, " is out of bounds of tensor"); - TORCH_CHECK(!self.is_complex(), "gather(): Yet not supported for complex"); + + if (self.is_complex()) { + auto self_real = at::view_as_real(self); + auto index_expanded = expand_index_as_real(index); + auto output_real = at::view_as_real(maybe_expand_0_dim(output)); + structured_gather_out_mps::impl(self_real, dim, index_expanded, sparse_grad, output_real); + return; + } struct CachedGraph : public MPSCachedGraph { CachedGraph(MPSGraph* graph) : MPSCachedGraph(graph) {} @@ -145,7 +165,15 @@ static void scatter_mps_general(const Tensor& self_arg, TORCH_CHECK(self.scalar_type() == output.scalar_type() && output.scalar_type() == src.scalar_type(), "scatter(): self, src and output must have the same scalar type"); TORCH_CHECK(dim >= 0 && dim < self.dim(), "scatter(): Indexing dim ", dim, " is out of bounds of tensor"); - TORCH_CHECK(!self.is_complex(), "scatter(): Yet not supported for complex"); + + if (self.is_complex()) { + auto self_real = at::view_as_real(self); + auto index_expanded = expand_index_as_real(index); + auto src_real = at::view_as_real(maybe_expand_0_dim(src)); + auto output_real = at::view_as_real(maybe_expand_0_dim(output)); + scatter_mps_general(self_real, dim, index_expanded, src_real, output_real, func_name, reduce); + return; + } struct CachedGraph : public MPSCachedGraph { CachedGraph(MPSGraph* graph) : MPSCachedGraph(graph) {} diff --git a/torch/testing/_internal/common_methods_invocations.py b/torch/testing/_internal/common_methods_invocations.py index db6904ecc0db3..fd46dbcd52f11 100644 --- a/torch/testing/_internal/common_methods_invocations.py +++ b/torch/testing/_internal/common_methods_invocations.py @@ -19558,10 +19558,6 @@ def sample_inputs_abs(op_info, device, dtype, requires_grad, op_kwargs=None, **k supports_forward_ad=True, supports_fwgrad_bwgrad=True, error_inputs_func=error_inputs_gather, - skips=( - # RuntimeError: gather(): Yet not supported for complex - DecorateInfo(unittest.expectedFailure, 'TestCommon', device_type='mps', dtypes=(torch.complex64,)), - ), ), OpInfo('index_fill', dtypes=all_types_and_complex_and(torch.bool, torch.float16, torch.bfloat16, torch.complex32), @@ -19766,9 +19762,6 @@ def sample_inputs_abs(op_info, device, dtype, requires_grad, op_kwargs=None, **k # Compiler issue on ROCm. Regression started in ROCm 6.4. DecorateInfo(unittest.skip('Skipped!'), 'TestCommon', 'test_non_standard_bool_values', dtypes=[torch.bool], active_if=TEST_WITH_ROCM), - # RuntimeError: scatter(): Yet not supported for complex - DecorateInfo(unittest.expectedFailure, 'TestCommon', 'test_dtypes', device_type='mps'), - DecorateInfo(unittest.expectedFailure, 'TestCommon', device_type='mps', dtypes=(torch.complex64,)), )), UnaryUfuncInfo( 'bfloat16', @@ -20582,9 +20575,6 @@ def sample_inputs_abs(op_info, device, dtype, requires_grad, op_kwargs=None, **k # Compiler issue on ROCm. Regression started in ROCm 6.4. DecorateInfo(unittest.skip('Skipped!'), 'TestCommon', 'test_non_standard_bool_values', dtypes=[torch.bool], active_if=TEST_WITH_ROCM), - # RuntimeError: scatter(): Yet not supported for complex - DecorateInfo(unittest.expectedFailure, 'TestCommon', 'test_dtypes', device_type='mps'), - DecorateInfo(unittest.expectedFailure, 'TestCommon', device_type='mps', dtypes=(torch.complex64,)), )), OpInfo('stack', dtypes=all_types_and_complex_and(torch.complex32, torch.bool, torch.float16, torch.bfloat16), @@ -20962,9 +20952,6 @@ def sample_inputs_abs(op_info, device, dtype, requires_grad, op_kwargs=None, **k decorators=( # RuntimeError: view size is not compatible with input tensor's size and stride DecorateInfo(unittest.expectedFailure, "TestMeta", "test_dispatch_symbolic_meta_outplace_all_strides"), - # MPS: gather(): Yet not supported for complex - DecorateInfo(unittest.expectedFailure, 'TestCommon', 'test_dtypes', device_type='mps'), - DecorateInfo(unittest.expectedFailure, 'TestCommon', device_type='mps', dtypes=(torch.complex64,)), )), ShapeFuncInfo('tile', ref=np.tile, @@ -23692,9 +23679,6 @@ def sample_inputs_abs(op_info, device, dtype, requires_grad, op_kwargs=None, **k DecorateInfo(unittest.expectedFailure, 'TestCommon', 'test_python_ref'), - # MPS: gather(): Yet not supported for complex - DecorateInfo(unittest.expectedFailure, 'TestCommon', 'test_dtypes', device_type='mps'), - DecorateInfo(unittest.expectedFailure, 'TestCommon', device_type='mps', dtypes=(torch.complex64,)), ), ), PythonRefInfo( diff --git a/torch/testing/_internal/common_mps.py b/torch/testing/_internal/common_mps.py index 0e7d23468737e..ef99d85f67a34 100644 --- a/torch/testing/_internal/common_mps.py +++ b/torch/testing/_internal/common_mps.py @@ -73,6 +73,7 @@ def mps_ops_modifier( "expand", "expand_as", "expand_copy", + "gather", "flatten", "fill", "full", @@ -155,6 +156,8 @@ def mps_ops_modifier( "rsqrt", "rsub", "scalar_tensor", + "scatter", + "scatter_add", "select", "sgn", "sigmoid", @@ -179,6 +182,7 @@ def mps_ops_modifier( "svd", "t", "t_copy", + "take_along_dim", "tanh", "tan", "tensor_split", From 47a12556640a30801ab88a0322685cd3ab9fa5d9 Mon Sep 17 00:00:00 2001 From: "Liangang,Zhang" Date: Mon, 23 Mar 2026 21:42:13 +0000 Subject: [PATCH 0038/1172] [XPU] Remove @expectedFailureXPU from test_max_pool2d_with_indices_backward5 (#178129) ## Summary Remove the `@expectedFailureXPU` decorator from `test_max_pool2d_with_indices_backward5` since the test now passes on XPU. ## Root Cause PR #167318 added a `scatter_add`-based decomposition for `max_pool2d_with_indices_backward` that correctly handles all configurations (including large kernel sizes like 13x13). This decomposition works correctly on XPU, but the `@expectedFailureXPU` decorator was not removed. Since `@expectedFailureXPU` expects the test to **fail**, a passing test triggers an "unexpected success" error, which is reported as a test failure. ## Fixes - `test/inductor/test_torchinductor.py::GPUTests::test_max_pool2d_with_indices_backward5_xpu` - `test/inductor/test_torchinductor_dynamic_shapes.py::DynamicShapesGPUTests::test_max_pool2d_with_indices_backward5_dynamic_shapes_xpu` - `test/inductor/test_torchinductor_codegen_dynamic_shapes.py::DynamicShapesCodegenGPUTests::test_max_pool2d_with_indices_backward5_dynamic_shapes_xpu` - `test/inductor/test_compile_subprocess.py::GPUTests::test_max_pool2d_with_indices_backward5_xpu` All four test files share the same `CommonTemplate` via `copy_tests()`, so the single decorator removal fixes all four. ## Verification Tested locally on Intel GPU (BMG) with XPU + Triton: the `scatter_add` decomposition produces correct results with max diff ~6.7e-06 vs eager reference. Pull Request resolved: https://github.com/pytorch/pytorch/pull/178129 Approved by: https://github.com/chuanqi129, https://github.com/isuruf, https://github.com/Skylion007 --- test/inductor/test_torchinductor.py | 1 - 1 file changed, 1 deletion(-) diff --git a/test/inductor/test_torchinductor.py b/test/inductor/test_torchinductor.py index d9a9c17c7f490..20c0132094314 100644 --- a/test/inductor/test_torchinductor.py +++ b/test/inductor/test_torchinductor.py @@ -10491,7 +10491,6 @@ def fn(a, b, c): # Correctness is validated by self.common() above. self.assertGreater(torch._inductor.metrics.generated_kernel_count, 0) - @expectedFailureXPU def test_max_pool2d_with_indices_backward5(self): # Large window size - decomposition handles via scatter_add def fn(a, b, c): From d5927f7a285c3f862b94faed9d4c8f24c795279c Mon Sep 17 00:00:00 2001 From: PyTorch MergeBot Date: Mon, 23 Mar 2026 22:02:32 +0000 Subject: [PATCH 0039/1172] Revert "[inductor] Fix pad_mm leaking padded strides to user-visible outputs (#177546)" This reverts commit 66ceb1e4f4ffe3b1863f1ba6c50eb82819f359c0. Reverted https://github.com/pytorch/pytorch/pull/177546 on behalf of https://github.com/eellison due to internal mem regression ([comment](https://github.com/pytorch/pytorch/pull/177546#issuecomment-4114000298)) --- test/inductor/test_pad_mm.py | 17 ----------------- torch/_inductor/compile_fx.py | 11 ----------- torch/_inductor/graph.py | 23 ++++++++++++++++++++--- torch/_inductor/ir.py | 6 +----- 4 files changed, 21 insertions(+), 36 deletions(-) diff --git a/test/inductor/test_pad_mm.py b/test/inductor/test_pad_mm.py index 38affde68448a..728b8635765e1 100644 --- a/test/inductor/test_pad_mm.py +++ b/test/inductor/test_pad_mm.py @@ -650,23 +650,6 @@ def test_masked_mha(B, H, S, D, device, dtype): torch.manual_seed(42) test_masked_mha(B, H, S, D, device, dtype) - @inductor_config.patch(force_shape_pad=True) - def test_pad_mm_output_strides_preserved(self): - """Regression test: pad_mm creates views with padded strides. - User-visible output strides must match eager execution.""" - - def fn(x, y): - return torch.mm(x, y) - - # N=2 gets padded to 4, so the mm output is (3, 4) sliced to (3, 2), - # creating a view with stride (4, 1) instead of the expected (2, 1). - x = torch.randn(3, 5, device=GPU_TYPE) - y = torch.randn(5, 2, device=GPU_TYPE) - expected = fn(x, y) - compiled = torch.compile(fn)(x, y) - self.assertEqual(compiled, expected) - self.assertEqual(compiled.stride(), expected.stride()) - if __name__ == "__main__": if HAS_GPU_AND_TRITON: diff --git a/torch/_inductor/compile_fx.py b/torch/_inductor/compile_fx.py index 288f72efd1249..8b414805ed808 100644 --- a/torch/_inductor/compile_fx.py +++ b/torch/_inductor/compile_fx.py @@ -259,12 +259,6 @@ def get_static_input_idxs(num_fixed: int) -> list[int]: def record_original_output_strides(gm: GraphModule) -> None: output_node = gm.graph.find_nodes(op="output")[0] - - # Don't overwrite strides that were already recorded (e.g., before - # joint_graph_passes which can introduce padded strides via pad_mm). - if "original_output_strides" in output_node.meta: - return - output_strides = [] if not isinstance(output_node.args[0], torch.fx.Node): @@ -2389,11 +2383,6 @@ def compile_fx_forward( for arg in output.args[0] # type: ignore[union-attr] ] - # Record original output strides BEFORE joint_graph_passes, because - # pad_mm (run as part of joint_graph_passes) can introduce views with - # padded strides that would be incorrectly captured as "original". - _recursive_record_original_output_strides(gm) - inputs_devices = get_inputs_devices(example_inputs, gm) gm = _recursive_joint_graph_passes(gm, input_device=next(iter(inputs_devices))) diff --git a/torch/_inductor/graph.py b/torch/_inductor/graph.py index c386b48d6f57f..24dc57b214dda 100644 --- a/torch/_inductor/graph.py +++ b/torch/_inductor/graph.py @@ -1927,9 +1927,26 @@ def maybe_apply_channels_last_stride_order( result.get_size(), torch.channels_last ) if not unbacked_symbols_in_strides and len(strides): - result = ir.ExternKernel.require_exact_strides( - result, strides, allow_padding=allow_padding - ) + # To avoid converting possible view ops to a copy kernel, we use the previous + # require_exact_strides to handle views. But ultimately it's better to require + # the right strides at the tensor definition. + if n.meta["val"]._is_view() or isinstance( + result.data, + ir.BaseView, + ): + result = ir.ExternKernel.require_stride_order( + result, + ir.get_stride_order(strides), + allow_padding=allow_padding, + ) + else: + # Fix for 0-d tensors: if result size is empty, + # strides should also be empty + if len(result.get_size()) == 0 and len(strides) > 0: + strides = [] + result = ir.ExternKernel.require_exact_strides( + result, strides, allow_padding=allow_padding + ) # Realize if (1) any user need inputs realized, or (2) there is # already too many reads and rematerializing can be bad. diff --git a/torch/_inductor/ir.py b/torch/_inductor/ir.py index f315e042dec60..050d316d6c63d 100644 --- a/torch/_inductor/ir.py +++ b/torch/_inductor/ir.py @@ -465,13 +465,9 @@ def significant_strides_equal( shape: Sequence[_IntLike], ) -> bool: """ - Returns true if the strides are equal, ignoring dimensions of size 0 or 1. - If any dimension is size 0, all strides are insignificant since the tensor - is empty. + Returns true if the strides are equal, ignoring dimensions of size 1 . """ assert len(shape) == len(strides1) and len(strides1) == len(strides2) - if any(V.graph.sizevars.statically_known_equals(dim, 0) for dim in shape): - return True for dim, s1, s2 in zip(shape, strides1, strides2): if V.graph.sizevars.statically_known_leq(dim, 1): continue From 348e0eba5f15f6601644016c6b9d1ed0ade07e93 Mon Sep 17 00:00:00 2001 From: Ethan Wee Date: Mon, 23 Mar 2026 22:13:21 +0000 Subject: [PATCH 0040/1172] [CI] Allow configuring pytest reruns via PYTORCH_NUM_PYTEST_RERUNS env var (#178007) The number of pytest reruns for failed tests in run_test.py is currently hardcoded to 2, making it impossible for downstream consumers (e.g. ROCm CI) to adjust or disable reruns without patching the file. This adds a PYTORCH_NUM_PYTEST_RERUNS environment variable that defaults to 2 (preserving existing behavior) and can be set to 0 to disable reruns entirely. ### Verification Temporarily set `PYTORCH_NUM_PYTEST_RERUNS=0` in `test.sh` to confirm the env var takes effect. CI log from [linux-jammy-rocm-py3.10 / test (default, 1, 6)](https://github.com/pytorch/pytorch/actions/runs/23419053468/job/68122075702) shows: ``` + export PYTORCH_NUM_PYTEST_RERUNS=0 + PYTORCH_NUM_PYTEST_RERUNS=0 Executing ['.../python', '-bb', 'test_cuda.py', ..., '-x', '--reruns=0', ...] Executing ['.../python', '-bb', 'test_ops.py', ..., '-x', '--reruns=0', ...] Executing ['.../python', '-bb', 'inductor/test_torchinductor.py', ..., '-x', '--reruns=0', ...] ``` The temporary `test.sh` change has been reverted. Pull Request resolved: https://github.com/pytorch/pytorch/pull/178007 Approved by: https://github.com/jithunnair-amd, https://github.com/jeffdaily --- test/run_test.py | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/test/run_test.py b/test/run_test.py index ff2a617a6eb78..0b2844efda9bf 100755 --- a/test/run_test.py +++ b/test/run_test.py @@ -104,6 +104,7 @@ def upload_adhoc_failure_json(*args, **kwargs): TEST_CONFIG = os.getenv("TEST_CONFIG", "") BUILD_ENVIRONMENT = os.getenv("BUILD_ENVIRONMENT", "") RERUN_DISABLED_TESTS = os.getenv("PYTORCH_TEST_RERUN_DISABLED_TESTS", "0") == "1" +NUM_PYTEST_RERUNS = int(os.getenv("PYTORCH_NUM_PYTEST_RERUNS", "2")) DISTRIBUTED_TEST_PREFIX = "distributed" INDUCTOR_TEST_PREFIX = "inductor" IS_SLOW = "slow" in TEST_CONFIG or "slow" in BUILD_ENVIRONMENT @@ -1266,9 +1267,9 @@ def get_pytest_args(options, is_cpp_test=False, is_distributed_test=False): # flakiness status. Default to 50 re-runs rerun_options = ["--flake-finder", f"--flake-runs={count}"] else: - # When under the normal mode, retry a failed test 2 more times. -x means stop at the first - # failure - rerun_options = ["-x", "--reruns=2"] + # When under the normal mode, retry a failed test NUM_PYTEST_RERUNS more times. + # -x means stop at the first failure. Set PYTORCH_NUM_PYTEST_RERUNS=0 to disable. + rerun_options = ["-x", f"--reruns={NUM_PYTEST_RERUNS}"] pytest_args = [ "-vv", From b5e62dda31721221d9dbfd28d1ee0de92045ec2f Mon Sep 17 00:00:00 2001 From: IvanKobzarev Date: Mon, 23 Mar 2026 09:52:43 -0700 Subject: [PATCH 0041/1172] [inductor] pre comm passes graph spmd check (#177960) Adding graph SPMD check before passes that move communications. Non-SPMD is not supported and applying passes can result in collectives ordering mismatch and result in annoying 10 mins (default NCCL_TIMEOUT=600) hangs. Adding spmd check that is checking with hash the structure of fx nodes and calculates hash of them and check them across ranks. By default this results only in the warning log and tlparse artifact with where the diff is. Optional setting to make it hard failing with RuntimeException. This should improve debugging with agents and without where it silently hangs without any sign on stdout/stderr. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177960 Approved by: https://github.com/eellison --- .../test_aten_comm_compute_reordering.py | 43 ++++ torch/_inductor/codecache.py | 13 + torch/_inductor/config.py | 10 + torch/_inductor/fx_passes/post_grad.py | 9 + torch/_inductor/fx_passes/spmd_check.py | 228 ++++++++++++++++++ 5 files changed, 303 insertions(+) create mode 100644 torch/_inductor/fx_passes/spmd_check.py diff --git a/test/distributed/test_aten_comm_compute_reordering.py b/test/distributed/test_aten_comm_compute_reordering.py index 396c240cdccd3..838b069d0f4c7 100644 --- a/test/distributed/test_aten_comm_compute_reordering.py +++ b/test/distributed/test_aten_comm_compute_reordering.py @@ -1283,6 +1283,49 @@ def func(a, *, ranks): "No-op pad/slice may have been eliminated by remove_noop_ops.", ) + @unittest.skipIf(not HAS_GPU, "Inductor+gpu needs triton and recent GPU arch") + @torch._inductor.config.patch( + { + **get_bucket_patches(), + "aten_distributed_optimizations.enable_overlap_scheduling": True, + "aten_distributed_optimizations.spmd_check": True, + "aten_distributed_optimizations.spmd_mismatch": "error", + } + ) + def test_spmd_verify_crashes_on_mismatch(self): + """Test that spmd_mismatch="error" raises on non-SPMD graphs.""" + + def func(a, *, ranks): + # rank 0 gets (4, 8), rank 1 gets (3, 8) — different node counts + # if pad is eliminated on rank 0 but not rank 1 + full_chunk = (7 + len(ranks) - 1) // len(ranks) + pad_size = full_chunk - a.size(0) + # Only rank 1 will have a real pad op here + if pad_size > 0: + a = torch.nn.functional.pad(a, [0, 0, 0, pad_size]) + ag = _functional_collectives.all_gather_tensor(a, 0, ranks) + return ag + 1 + + with _dynamo_dist_per_rank_init( + self.rank, + self.world_size, + self.backend(device_type), + fake_pg=not at_least_x_gpu(2), + ): + world_size = self.world_size + full_chunk = (7 + world_size - 1) // world_size + local_size = full_chunk if self.rank == 0 else 7 - full_chunk + a = torch.randn(local_size, 8, device=device_type) + ranks = list(range(world_size)) + + func_c = functools.partial(func, ranks=ranks) + compiled = torch.compile(func_c) + + # The graph will differ: rank 0 has no pad, rank 1 has pad. + # With spmd_check_crash_on_mismatch=True, this should raise. + with self.assertRaises(RuntimeError, msg="SPMD graph verification"): + compiled(a) + def get_toy_model(device_type: str): """ diff --git a/torch/_inductor/codecache.py b/torch/_inductor/codecache.py index e247e52882980..a358083753433 100644 --- a/torch/_inductor/codecache.py +++ b/torch/_inductor/codecache.py @@ -500,14 +500,19 @@ def __init__( self, gm: torch.fx.GraphModule, has_user_defined_triton_kernels: bool = False, + device_id_agnostic: bool = False, ) -> None: """ Create an FX graph pickler. If include_non_inlined=True, then pickling will include the _values_ for all Tensors. (Note that any tensors are constants attached as attributes to the GraphModule). Otherwise, pickling will include only the metadata for these tensors. + + If device_id_agnostic=True, device indices in TensorMetadata are normalized + to 0, so that the same graph on different GPUs produces identical bytes. """ self._stream = io.BytesIO() + self._device_id_agnostic = device_id_agnostic super().__init__(self._stream) self.dispatch_table = copyreg.dispatch_table.copy() @@ -540,6 +545,10 @@ def _reduce_fake_tensor( Custom reducer to pickle FakeTensors. """ metadata = extract_tensor_metadata_for_cache_key(t) + if self._device_id_agnostic: + metadata = dataclasses.replace( + metadata, device=torch.device(metadata.device.type, 0) + ) return (_ident, (metadata,)) def _reduce_tensor( @@ -558,6 +567,10 @@ def _reduce_tensor( raise BypassFxGraphCache("mkldnn tensors unpickleable") metadata = extract_tensor_metadata_for_cache_key(t) + if self._device_id_agnostic: + metadata = dataclasses.replace( + metadata, device=torch.device(metadata.device.type, 0) + ) # If this is a non-inlined frozen parameter, we consider the metadata only. if is_frozen_param(t) and not GraphLowering.can_inline_constant(t): diff --git a/torch/_inductor/config.py b/torch/_inductor/config.py index 936e2d090efcb..8a5745af6130d 100644 --- a/torch/_inductor/config.py +++ b/torch/_inductor/config.py @@ -1123,6 +1123,16 @@ class aten_distributed_optimizations: # Prioritize bucketing during overlap scheduling by grouping candidates by bucket key prioritize_bucketing_during_scheduling: bool = True + # Verify FX graphs are identical across ranks before overlap scheduling. + # Detects non-SPMD graphs that would cause NCCL collective ordering + # mismatches and hangs. + spmd_check: bool = True + + # Action on SPMD graph mismatch: "warn" logs a warning, "error" raises + # RuntimeError. "error" fails fast instead of risking silent NCCL hang. + # TODO(ivankobzarev): change default to "error" after real-world testing. + spmd_mismatch: Literal["warn", "error"] = "warn" + def parallel_compile_enabled_internally() -> bool: """ diff --git a/torch/_inductor/fx_passes/post_grad.py b/torch/_inductor/fx_passes/post_grad.py index 7b947937646e6..bbefebf92a237 100644 --- a/torch/_inductor/fx_passes/post_grad.py +++ b/torch/_inductor/fx_passes/post_grad.py @@ -245,6 +245,15 @@ def post_grad_passes(gm: torch.fx.GraphModule, is_inference: bool): pass_name = "custom_backend_passes_" + device GraphTransformObserver(gm, pass_name).apply_gm_pass(custom_backend_pass) + # SPMD verification — before collective reordering passes. + if ( + config.aten_distributed_optimizations.spmd_check + and _needs_spmd_graph_preservation() + ): + from torch._inductor.fx_passes.spmd_check import spmd_check + + spmd_check(gm) + collectives_bucketing: bool = False if config.bucket_reduce_scatters_fx != "none": diff --git a/torch/_inductor/fx_passes/spmd_check.py b/torch/_inductor/fx_passes/spmd_check.py new file mode 100644 index 0000000000000..b51004107a107 --- /dev/null +++ b/torch/_inductor/fx_passes/spmd_check.py @@ -0,0 +1,228 @@ +"""SPMD graph verification for overlap scheduling. + +Verifies all ranks have identical FX graph structure before collective +reordering passes. Non-SPMD graphs cause NCCL collective ordering +mismatches and hangs. +""" + +import hashlib +import logging +from collections import Counter + +import torch +from torch._inductor import config +from torch._logging import trace_structured + + +log = logging.getLogger(__name__) + + +def _compute_hash(gm: torch.fx.GraphModule) -> int | None: + """Compute a structural hash of the graph including tensor metadata. + + Uses FxGraphCachePickler(device_id_agnostic=True) to serialize + (target, val) per call_function node, capturing op targets and + FakeTensor metadata (dtype, shape, stride, etc.) with device indices + normalized to 0. + + Returns None if the graph contains unpicklable objects. + """ + from torch._inductor.codecache import BypassFxGraphCache, FxGraphCachePickler + + try: + pickler = FxGraphCachePickler(gm, device_id_agnostic=True) + data = pickler.dumps( + tuple( + (str(n.target), n.meta.get("val")) + for n in gm.graph.nodes + if n.op == "call_function" + ) + ) + digest = hashlib.blake2b(data, digest_size=8).digest() + return int.from_bytes(digest, "big", signed=True) + except BypassFxGraphCache: + # FxGraphCachePickler can't serialize certain objects: + # mkldnn tensors, BackwardState, torchbind objects, or general + # pickle failures. Skip the SPMD check gracefully. + log.warning("SPMD check: skipping, unpicklable graph objects", exc_info=True) + return None + + +def _build_diag_fingerprint( + gm: torch.fx.GraphModule, +) -> tuple[tuple[str, str | None], ...]: + """Build human-readable fingerprint for mismatch diagnostics. + + Only called on the rare mismatch path. + """ + from torch._inductor.codecache import extract_tensor_metadata_for_cache_key + + entries: list[tuple[str, str | None]] = [] + for n in gm.graph.nodes: + if n.op != "call_function": + continue + target_str = str(n.target) + val = n.meta.get("val") + entries.append( + ( + target_str, + _format_val_metadata(val, extract_tensor_metadata_for_cache_key), + ) + ) + return tuple(entries) + + +def _format_val_metadata(val: object, extract_fn: object) -> str | None: + """Format node val metadata for human-readable diagnostics.""" + if val is None: + return None + if isinstance(val, torch.Tensor): + return str(extract_fn(val)) # type: ignore[operator] + if isinstance(val, (tuple, list)): + parts = [] + for v in val: + if isinstance(v, torch.Tensor): + parts.append(str(extract_fn(v))) # type: ignore[operator] + else: + parts.append(str(type(v).__name__)) + return f"({', '.join(parts)})" + return str(type(val).__name__) + + +def spmd_check(gm: torch.fx.GraphModule) -> bool: + """Verify all ranks have identical FX graph structure (SPMD). + + Computes a structural hash (op targets + tensor metadata including + shapes, dtypes, strides) and compares across ranks. + On mismatch, emits a diagnostic report to stdout, logging, and + trace_structured. + + Returns True if graphs match (SPMD), False on mismatch. + """ + import torch.distributed as dist + + if not dist.is_initialized() or dist.get_world_size() <= 1: + return True + + structure_hash = _compute_hash(gm) + if structure_hash is None: + return True + + from torch._subclasses.fake_tensor import unset_fake_temporarily + from torch.distributed.distributed_c10d import _get_default_group + + pg = _get_default_group() + world_size = dist.get_world_size() + rank = dist.get_rank() + + with unset_fake_temporarily(): + all_hashes: list[int] = [0] * world_size + dist.all_gather_object(all_hashes, structure_hash, pg) + + if all(h == all_hashes[0] for h in all_hashes): + return True + + # Mismatch detected — build and gather diagnostic fingerprints + fingerprint = _build_diag_fingerprint(gm) + with unset_fake_temporarily(): + all_fingerprints: list[tuple[object, ...]] = [() for _ in range(world_size)] + dist.all_gather_object(all_fingerprints, fingerprint, pg) + + report = _build_mismatch_report(all_fingerprints, rank, world_size) + + print(report, flush=True) + log.warning("\n%s", report) + + trace_structured( + "artifact", + metadata_fn=lambda: { + "name": "inductor_spmd_graph_mismatch", + "encoding": "string", + }, + payload_fn=lambda: report, + ) + + if config.aten_distributed_optimizations.spmd_mismatch == "error": + raise RuntimeError( + "SPMD graph verification failed. " + 'Set aten_distributed_optimizations.spmd_mismatch="warn" ' + "to warn instead of fail.\n" + report + ) + + return False + + +def _entry_target(entry: object) -> str: + """Extract the target string from a fingerprint entry.""" + if isinstance(entry, tuple): + return str(entry[0]) + return str(entry) + + +def _entry_metadata(entry: object) -> str: + """Format metadata from a fingerprint entry, if present.""" + if isinstance(entry, tuple) and len(entry) >= 2: + meta = entry[1] + if meta is not None: + return f" meta={meta}" + return "" + + +def _build_mismatch_report( + all_fingerprints: list[tuple[object, ...]], + rank: int, + world_size: int, +) -> str: + """Build diagnostic report for SPMD graph mismatch.""" + lines = [ + "=" * 80, + f"SPMD GRAPH MISMATCH — rank {rank}, world_size={world_size}", + "=" * 80, + ] + + # Node count per rank + counts = [len(t) for t in all_fingerprints] + lines.append("NODE COUNTS PER RANK:") + for r in range(world_size): + marker = " <--" if counts[r] != counts[0] else "" + lines.append(f" rank {r}: {counts[r]} call_function nodes{marker}") + lines.append("") + + # Find entries that differ + ref = all_fingerprints[0] + for r in range(1, world_size): + other = all_fingerprints[r] + if other == ref: + continue + lines.append(f"DIFFS rank 0 vs rank {r}:") + + # Show first few positional differences + max_diffs = 10 + shown = 0 + for i, (a, b) in enumerate(zip(ref, other)): + if a != b and shown < max_diffs: + lines.append(f" node {i}:") + lines.append(f" rank 0: {_entry_target(a)}{_entry_metadata(a)}") + lines.append(f" rank {r}: {_entry_target(b)}{_entry_metadata(b)}") + shown += 1 + + # Also show count-based diffs for op targets + ref_targets = [_entry_target(e) for e in ref] + other_targets = [_entry_target(e) for e in other] + + ref_counts = Counter(ref_targets) + other_counts = Counter(other_targets) + only_ref = ref_counts - other_counts + only_other = other_counts - ref_counts + if only_ref: + lines.append(" Only on rank 0:") + for op, cnt in only_ref.most_common(10): + lines.append(f" {op} (x{cnt})") + if only_other: + lines.append(f" Only on rank {r}:") + for op, cnt in only_other.most_common(10): + lines.append(f" {op} (x{cnt})") + lines.append("") + + lines.append("=" * 80) + return "\n".join(lines) From 994d496ffab451e42d80ad8f77fa442bd0e3b97c Mon Sep 17 00:00:00 2001 From: morrison-turnansky Date: Mon, 23 Mar 2026 22:29:16 +0000 Subject: [PATCH 0042/1172] RFC: Increase Pattern Matcher Observability (#177032) This is a working implementation of the RFC up for discussion at https://dev-discuss.pytorch.org/t/feature-pattern-matcher-observability/3321. At a high level the purpose of this to help the development experience with more debugging and logging tools for PatternMatcher. Critically, this should not effect current functionality of pattern matcher for users. Please refer to RFC for more detail. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177032 Approved by: https://github.com/karthickai Co-authored-by: Parshant Sharma --- test/inductor/test_pattern_matcher.py | 154 ++++++++++++++++++++++++++ torch/_inductor/pattern_matcher.py | 55 ++++++++- 2 files changed, 207 insertions(+), 2 deletions(-) diff --git a/test/inductor/test_pattern_matcher.py b/test/inductor/test_pattern_matcher.py index f2ee6b9a2403c..e5412d1b2a5ac 100644 --- a/test/inductor/test_pattern_matcher.py +++ b/test/inductor/test_pattern_matcher.py @@ -2176,6 +2176,98 @@ def fn(x): specific_record.getMessage(), ) + @make_logging_test() + def test_pattern_match_debug_multiple_nodes(self, records): + def pattern_add(x, y): + return x + y + + def replacement_add(x, y): + return x * y + + def pattern_sub(x, y): + return x - y + + def replacement_sub(x, y): + return x * y + + my_patterns = PatternMatcherPass() + inputs = [ + torch.randn(4, 4, device=GPU_TYPE), + torch.randn(4, 4, device=GPU_TYPE), + ] + register_replacement( + pattern_add, replacement_add, inputs, fwd_only, my_patterns + ) + register_replacement( + pattern_sub, replacement_sub, inputs, fwd_only, my_patterns + ) + + def custom_pass(graph: torch.fx.Graph): + return my_patterns.apply(graph) + + def fn(x, y): + return (x + y) + (x - y) + + x = torch.randn(4, 4, device=GPU_TYPE) + y = torch.randn(4, 4, device=GPU_TYPE) + + # Debug both "add" and "sub" nodes + with unittest.mock.patch.dict( + os.environ, {"TORCHINDUCTOR_PATTERN_MATCH_DEBUG": "add,sub"} + ): + compiled_fn = torch.compile( + fn, options={"post_grad_custom_post_pass": custom_pass} + ) + _ = compiled_fn(x, y) + + self.assertTrue(self.hasRecord(records, "Specific pattern match: add")) + self.assertTrue(self.hasRecord(records, "Specific pattern match: sub")) + + @make_logging_test() + def test_pattern_match_debug_all_nodes(self, records): + def pattern_add(x, y): + return x + y + + def replacement_add(x, y): + return x * y + + def pattern_sub(x, y): + return x - y + + def replacement_sub(x, y): + return x * y + + my_patterns = PatternMatcherPass() + inputs = [ + torch.randn(4, 4, device=GPU_TYPE), + torch.randn(4, 4, device=GPU_TYPE), + ] + register_replacement( + pattern_add, replacement_add, inputs, fwd_only, my_patterns + ) + register_replacement( + pattern_sub, replacement_sub, inputs, fwd_only, my_patterns + ) + + def custom_pass(graph: torch.fx.Graph): + return my_patterns.apply(graph) + + def fn(x, y): + return (x + y) + (x - y) + + x = torch.randn(4, 4, device=GPU_TYPE) + y = torch.randn(4, 4, device=GPU_TYPE) + + with unittest.mock.patch.dict( + os.environ, {"TORCHINDUCTOR_PATTERN_MATCH_DEBUG": "all"} + ): + compiled_fn = torch.compile( + fn, options={"post_grad_custom_post_pass": custom_pass} + ) + _ = compiled_fn(x, y) + self.assertTrue(self.hasRecord(records, "Specific pattern match: add")) + self.assertTrue(self.hasRecord(records, "Specific pattern match: sub")) + def test_gumbel_max_trick(self): counters.clear() @@ -2213,6 +2305,68 @@ def sample(logits, temperature): self.assertTrue(counters["inductor"]["apply_gumbel_max_trick"] == 1) + def test_per_pattern_counter(self): + """Test that per-pattern counters track individual pattern matches""" + with inductor_config.patch(fx_graph_cache=False): + with unittest.mock.patch.dict( + os.environ, {"TORCHINDUCTOR_PATTERN_MATCH_DEBUG": "1"} + ): + counters.clear() + + def fn(x, y): + return torch.bmm(x, y) + + x = torch.randn(4, 10, 10, device=GPU_TYPE) + y = torch.randn(4, 10, 10, device=GPU_TYPE) + + compiled = torch.compile(fn) + compiled(x, y) + + counter_key = "inductor_pattern_matcher_per_pattern" + per_pattern = counters.get(counter_key, None) + + self.assertIsInstance(per_pattern, dict) + self.assertGreater(len(per_pattern), 0) + self.assertIn("CallFunction_aten.bmm.default", per_pattern) + self.assertEqual(per_pattern["CallFunction_aten.bmm.default"], 1) + + def test_per_pattern_counter_accumulation(self): + """Test that per-pattern counters accumulate across compilations""" + with inductor_config.patch(fx_graph_cache=False): + with unittest.mock.patch.dict( + os.environ, {"TORCHINDUCTOR_PATTERN_MATCH_DEBUG": "1"} + ): + counter_key = "inductor_pattern_matcher_per_pattern" + + counters.clear() + + x = torch.randn(2, 10, 10, device=GPU_TYPE) + y = torch.randn(2, 10, 10, device=GPU_TYPE) + + def fn1(a, b): + return torch.bmm(a, b) + + compiled1 = torch.compile(fn1) + compiled1(x, y) + count1 = sum(counters.get(counter_key, {}).values()) + + # Compile second function without clearing counters + def fn2(a, b): + return torch.bmm(a, b) * 2 + + compiled2 = torch.compile(fn2) + compiled2(x, y) + accumulated_count = sum(counters.get(counter_key, {}).values()) + + # Verify accumulation + counters.clear() + torch._dynamo.reset() + compiled2 = torch.compile(fn2) + compiled2(x, y) + count2 = sum(counters.get(counter_key, {}).values()) + + self.assertEqual(accumulated_count, count1 + count2) + if __name__ == "__main__": if IS_LINUX and HAS_GPU: diff --git a/torch/_inductor/pattern_matcher.py b/torch/_inductor/pattern_matcher.py index 3099116bbc7f0..8926174dac0c1 100644 --- a/torch/_inductor/pattern_matcher.py +++ b/torch/_inductor/pattern_matcher.py @@ -90,6 +90,36 @@ backend = os.environ.get("TORCHINDUCTOR_PATTERN_MATCH_BACKEND", "inductor") +_debug_nodes_cache: bool | OrderedSet[str] | None = None +_debug_nodes_env_value_cache: str | None = None + + +def _should_debug_node(node_name: str) -> bool: + def _get_debug_nodes() -> bool | OrderedSet[str]: + global _debug_nodes_cache, _debug_nodes_env_value_cache + + def parse_debug_env(env_value: str | None) -> bool | OrderedSet[str]: + if not env_value: + return False + if env_value == "all": + return True + return OrderedSet(env_value.split(",")) + + current_env = os.environ.get("TORCHINDUCTOR_PATTERN_MATCH_DEBUG") + + # Recompute only if env changed + if current_env != _debug_nodes_env_value_cache or _debug_nodes_cache is None: + _debug_nodes_cache = parse_debug_env(current_env) + _debug_nodes_env_value_cache = current_env + + return _debug_nodes_cache + + debug_nodes = _get_debug_nodes() + if isinstance(debug_nodes, bool): + return debug_nodes + return node_name in debug_nodes + + class SearchFn(Protocol): __name__: str @@ -1150,6 +1180,7 @@ class ReplacementPatternEntry(PatternEntry): """ normalize_args: Callable[..., list[Any]] + pattern_name: str | None = None # Unique identifier for per-pattern telemetry @staticmethod def replace_with_graph( @@ -1454,6 +1485,7 @@ def register_replacement( exclusive_arg_names: Sequence[str] = (), search_fn_pattern: PatternExpr | None = None, skip_duplicates: bool = False, + pattern_name: str | None = None, ) -> bool: """ Create a replacement rule based on example functions that get traced @@ -1587,7 +1619,7 @@ def search_fn_new(*args_new: Any) -> Any: assert node is not None specific_pattern_match = specific_pattern.match(node) - if os.environ.get("TORCHINDUCTOR_PATTERN_MATCH_DEBUG") == node.name: + if _should_debug_node(node.name): log.warning( "Specific pattern match: %s%s %s %s", node, @@ -1657,6 +1689,7 @@ def normalize_args(**kwargs: Any) -> list[Any]: pattern=pattern, extra_check=check_fn, normalize_args=normalize_args, + pattern_name=pattern_name, ) pattern.register(pass_dicts) return pattern.pattern # type: ignore[return-value] @@ -1814,6 +1847,7 @@ def gen_register_replacement( exclusive_arg_names, search_fn_pattern=pat, skip_duplicates=skip_duplicates, + pattern_name=unique_name, ) @@ -2000,6 +2034,10 @@ def wrapper(*args: Any, **kwargs: Any) -> Any: class PatternMatcherPass: + """ + Registry of patterns to match and replace in FX graphs. + """ + def __init__( self, pass_name: str | None = None, @@ -2076,7 +2114,7 @@ def apply(self, gm: torch.fx.GraphModule | torch.fx.Graph) -> int: != 1 ): continue - if os.environ.get("TORCHINDUCTOR_PATTERN_MATCH_DEBUG") == node.name: + if _should_debug_node(node.name): log.warning("%s%s %s %s", node, node.args, m, entry.pattern) if is_match(m) and guard_or_false(entry.extra_check(m)): @@ -2084,6 +2122,19 @@ def apply(self, gm: torch.fx.GraphModule | torch.fx.Graph) -> int: entry.apply(m, graph, node) counters[backend]["pattern_matcher_count"] += 1 counters[backend]["pattern_matcher_nodes"] += len(m.nodes) + + # Track per-pattern counts when debug mode is active + if os.environ.get("TORCHINDUCTOR_PATTERN_MATCH_DEBUG"): + if getattr(entry, "pattern_name", None): + pattern_name = entry.pattern_name + else: + # Fallback: use pattern class name + operation target + pattern_class = entry.pattern.__class__.__name__ + target = str(node.target) + pattern_name = f"{pattern_class}_{target}" + + pattern_key = f"{backend}_pattern_matcher_per_pattern" + counters[pattern_key][pattern_name] += 1 return count def clear(self) -> None: From f9746f67000741cdc55a6e407224b5a92399de28 Mon Sep 17 00:00:00 2001 From: Bin Bao Date: Mon, 23 Mar 2026 10:03:11 -0700 Subject: [PATCH 0043/1172] [inductor] Make CUTLASS non-AOT cpp_wrapper warning only warn once (#178159) Use warnings.warn instead of log.warning so it only warns once. Pull Request resolved: https://github.com/pytorch/pytorch/pull/178159 Approved by: https://github.com/karthickai, https://github.com/mlazos --- torch/_inductor/utils.py | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/torch/_inductor/utils.py b/torch/_inductor/utils.py index adf7c7586660f..3ff803aae1fea 100644 --- a/torch/_inductor/utils.py +++ b/torch/_inductor/utils.py @@ -23,6 +23,7 @@ import textwrap import time import unittest +import warnings from collections.abc import ( Callable, Collection, @@ -2128,9 +2129,9 @@ def use_cutlass_template(layout: Layout, m: int, n: int, k: int) -> bool: # for the compiled CUTLASS .so, similar to how the triton branch uses # static CUfunction + loadKernel for non-AOT mode. if V.graph.cpp_wrapper and not V.graph.aot_mode: - log.warning( + warnings.warn( "CUTLASS backend is not supported with non-AOT cpp_wrapper mode. " - "Skipping CUTLASS backend." + "Skipping CUTLASS backend.", ) return False From c16398fbc63eaf8aff820c242fdde644fc0e3616 Mon Sep 17 00:00:00 2001 From: Daohang Shi Date: Mon, 23 Mar 2026 23:22:44 +0000 Subject: [PATCH 0044/1172] [torchTLX] Rename fused TLX kernels with tlx prefix (#177590) (#177590) Summary: Add a customize_fused_kernel_name hook on InductorChoices and override it in TLXInductorChoices to insert 'tlx' into fused kernel names when the source uses TLX APIs (e.g. 'fused_mm' -> 'fused_tlx_mm'). This replaces the config-based approach from D95602337 with a V.choices hook, minimizing non-fb changes. Authored with Claude. Test Plan: torchTLX unit tests Reviewed By: PaulZhang12 Differential Revision: D96810349 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177590 Approved by: https://github.com/PaulZhang12 --- torch/_inductor/choices.py | 4 ++++ torch/_inductor/codegen/triton.py | 2 ++ 2 files changed, 6 insertions(+) diff --git a/torch/_inductor/choices.py b/torch/_inductor/choices.py index 6d7dae9585d39..c4afe369bb356 100644 --- a/torch/_inductor/choices.py +++ b/torch/_inductor/choices.py @@ -347,6 +347,10 @@ def triton_kernel_kwargs( """Hook to change the kwargs passed to TritonKernel, used to apply fixed configurations""" return kernel_kwargs + def customize_fused_kernel_name(self, fused_name: str, src_code: str) -> str: + """Hook to transform fused kernel names during codegen""" + return fused_name + @staticmethod def should_use_cooperative_reduction(features: SIMDKernelFeatures) -> bool: """Heuristic to decide if a cooperative reduction should be used.""" diff --git a/torch/_inductor/codegen/triton.py b/torch/_inductor/codegen/triton.py index 7b8ac8fb6d80f..7ddeb37599645 100644 --- a/torch/_inductor/codegen/triton.py +++ b/torch/_inductor/codegen/triton.py @@ -6478,6 +6478,8 @@ def define_kernel(self, src_code, node_schedule, kernel): if config.triton.descriptive_names else "" ) + if fused_name: + fused_name = V.choices.customize_fused_kernel_name(fused_name, src_code) kernel_category = get_kernel_category_by_source_code(src_code)[:3] kernel_name = "_".join( ["triton", kernel_category, fused_name, wrapper.next_kernel_suffix()] From 6fcbf6d86540893c33746bf4386f8d3cb1570545 Mon Sep 17 00:00:00 2001 From: Pei Zhang Date: Tue, 24 Mar 2026 00:02:54 +0000 Subject: [PATCH 0045/1172] [python/3.10 removal] Remove CPython 3.10 holdback for caffe2/test/distributed/elastic/rendezvous (#178071) (#178071) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Summary: Remove CPython 3.10 pin for the 8 rendezvous test targets owned by aml_ai_platform oncall, upgrading them to Python 3.12. Changes: - Remove `"python": "3.10"` pin from `caffe2/test/distributed/elastic/rendezvous/PACKAGE` - Remove 8 entries from `CPYTHON310_HOLDBACK_TARGETS_ALLOWLISTS` in `cpython310_allowlist.bzl` - Fix Python 3.12 compatibility issue in `dynamic_rendezvous_test.py`: remove `spec=dist.PrefixStore` from `CustomPrefixStore()` call, since `dist.PrefixStore` is already mocked by `patch.object` decorator and Python 3.12 raises `InvalidSpecError` when spec-ing a Mock with another Mock Test Plan: Ran `buck test fbcode//caffe2/test/distributed/elastic/rendezvous:` — non-etcd tests pass on Python 3.12. Remaining failures are pre-existing (etcd binary missing in sandbox, ASAN heap-use-after-free in C++ teardown, timedelta bug in test_redundancy_transition_to_wait_list_then_join_rendezvous). Reviewed By: d4l3k, dcci Differential Revision: D97058526 Pull Request resolved: https://github.com/pytorch/pytorch/pull/178071 Approved by: https://github.com/Skylion007 --- test/distributed/elastic/rendezvous/dynamic_rendezvous_test.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/distributed/elastic/rendezvous/dynamic_rendezvous_test.py b/test/distributed/elastic/rendezvous/dynamic_rendezvous_test.py index 6340c306adf05..084b262602ff2 100644 --- a/test/distributed/elastic/rendezvous/dynamic_rendezvous_test.py +++ b/test/distributed/elastic/rendezvous/dynamic_rendezvous_test.py @@ -1829,7 +1829,7 @@ def get(self, key): def set(self, key, value): pass - prefix_store = CustomPrefixStore(spec=dist.PrefixStore) + prefix_store = CustomPrefixStore() prefix_store_class_mock.return_value = prefix_store tcp_store = Mock(spec=dist.TCPStore) original_addr = "original_addr" From ab8c7aef8a70d1ef27272a4c90a43ee3479c8c75 Mon Sep 17 00:00:00 2001 From: Anshul Sinha Date: Mon, 23 Mar 2026 12:02:34 -0700 Subject: [PATCH 0046/1172] [dtensor] Add single_dim_strategy infrastructure for foreach/fused ops (#177186) Add infrastructure support in the single_dim_strategy path for list-based ops (foreach, fused, amp_foreach): - single_dim_strategy.py: Add cross_mesh_indices field to _SingleDimStrategyInfo, rename _translate_foreach_op_schema to _translate_list_op_schema with support for fused/amp_foreach prefixes, extend detection to match all list op prefixes - utils.py: Add cross_mesh_indices parameter to expand_to_full_mesh_op_strategy to preserve original mesh for cross-mesh inputs (e.g. state_steps in fused adam) - _sharding_prop.py: Return -1 for List[Tensor] return types in _get_expected_num_tensor_outputs (dynamic count), skip validation when expected is -1, handle output_specs=None in TupleStrategy processing for void-returning inplace list ops Pull Request resolved: https://github.com/pytorch/pytorch/pull/177186 Approved by: https://github.com/wconstab --- .../tensor/_ops/single_dim_strategy.py | 95 +++++++++++++++++-- torch/distributed/tensor/_ops/utils.py | 47 +++++++++ torch/distributed/tensor/_sharding_prop.py | 22 ++++- 3 files changed, 150 insertions(+), 14 deletions(-) diff --git a/torch/distributed/tensor/_ops/single_dim_strategy.py b/torch/distributed/tensor/_ops/single_dim_strategy.py index c64b89a7c156f..4cb138c5b32ae 100644 --- a/torch/distributed/tensor/_ops/single_dim_strategy.py +++ b/torch/distributed/tensor/_ops/single_dim_strategy.py @@ -76,6 +76,10 @@ class _SingleDimStrategyInfo: func: _SingleDimStrategyFunc allow_unbacked_sharding: bool | None = field(default=None) allow_uneven_sharding: bool = field(default=False) + # Positions (in args_schema) of args that may live on a different mesh + # than the op's compute mesh. These args must be Replicate. + # See Note [Multi-mesh args] in expand_to_full_mesh_op_strategy. + different_mesh_args: list[int] | None = field(default=None) # Delegate to func so this can be used interchangeably with a raw # _SingleDimStrategyFunc (e.g. in tests that call strategy functions directly). @@ -282,12 +286,54 @@ def __init__( if isinstance(strategy_fn, _SingleDimStrategyInfo): self.allow_unbacked_sharding = strategy_fn.allow_unbacked_sharding self.allow_uneven_sharding = strategy_fn.allow_uneven_sharding + different_mesh_args = strategy_fn.different_mesh_args func = strategy_fn.func else: self.allow_unbacked_sharding = None self.allow_uneven_sharding = False + different_mesh_args = None func = strategy_fn + # Determine element_mesh from the first OpStrategy arg. For foreach + # per-element schemas the element's inputs may live on a smaller + # sub-mesh than the global compute_mesh. + self.element_mesh: DeviceMesh | None = None + for arg in op_schema.args_schema: + if isinstance(arg, OpStrategy): + self.element_mesh = arg.strategies[0].output_spec.mesh + break + + # Validate that all inputs are on the same mesh (except + # different_mesh_args which are explicitly allowed to differ). + if self.element_mesh is not None: + allowed = set(different_mesh_args or []) + for i, arg in enumerate(op_schema.args_schema): + if isinstance(arg, OpStrategy) and i not in allowed: + arg_mesh = arg.strategies[0].output_spec.mesh + if arg_mesh != self.element_mesh: + raise ValueError( + f"Cannot run {op_schema.op} on inputs with different " + f"meshes: got {self.element_mesh} and {arg_mesh}" + ) + + # Remap different_mesh_args from args_schema positions to + # OpStrategy-only positions. Non-OpStrategy args (e.g. empty lists) + # are filtered out by expand_to_full_mesh_op_strategy, shifting later + # indices. + self.remapped_different_mesh_args: list[int] | None = None + if different_mesh_args is not None: + schema_to_strategy: dict[int, int] = {} + strategy_pos = 0 + for schema_pos, arg in enumerate(op_schema.args_schema): + if isinstance(arg, OpStrategy): + schema_to_strategy[schema_pos] = strategy_pos + strategy_pos += 1 + self.remapped_different_mesh_args = [ + schema_to_strategy[i] + for i in different_mesh_args + if i in schema_to_strategy + ] + if num_inputs is None: num_inputs = _get_num_tensor_inputs(op_schema) self.num_inputs = num_inputs @@ -464,8 +510,10 @@ def expanded_strategy( base_name = op_name.split("::")[1].split(".")[0] is_inplace = base_name.endswith("_") + element_mesh = prepared_strategy.element_mesh or mesh + return expand_to_full_mesh_op_strategy( - mesh, + element_mesh, op_schema, cast(list[PlacementList], prepared_strategy.expanded_strategies), output_tensor_meta=output_tensor_meta, @@ -473,6 +521,7 @@ def expanded_strategy( input_index=prepared_strategy.num_outputs, allow_unbacked_sharding=prepared_strategy.allow_unbacked_sharding, allow_uneven_sharding=prepared_strategy.allow_uneven_sharding, + different_mesh_args=prepared_strategy.remapped_different_mesh_args, ) return expanded_strategy @@ -494,12 +543,14 @@ def _create_expanded_strategy( # Unhashable types (SymInts), skip caching return _create_expanded_strategy_impl(op_schema, output_tensor_meta) - def _translate_foreach_op_schema( - op_schema: OpSchema, output_tensor_meta: Sequence[TensorMeta], index: int - ) -> tuple[OpSchema, TensorMeta]: - """Translate foreach op to per-element version of schema.""" + def _translate_list_op_schema( + op_schema: OpSchema, + output_tensor_meta: Sequence[TensorMeta] | None, + index: int, + ) -> tuple[OpSchema, TensorMeta | None]: + """Translate foreach/fused op to per-element version of schema.""" op_parts = str(op_schema.op).split(".") - base_op_name = op_parts[-2].replace("_foreach_", "") + op_name = op_parts[-2] foreach_variant = op_parts[-1] # select per-element inputs, outputs @@ -509,7 +560,30 @@ def _translate_foreach_op_schema( (op_schema.args_schema, op_schema.kwargs_schema), is_leaf=lambda x: isinstance(x, TupleStrategy), ) - target_output_meta = output_tensor_meta[index] + # For inplace ops, output_tensor_meta is None + target_output_meta = ( + output_tensor_meta[index] if output_tensor_meta is not None else None + ) + + # Strip the prefix to get the base op name and find the per-element op. + # Fused ops (e.g. _fused_adam) have no per-element ATen equivalent, + # so we keep the original op unchanged. + if op_name.startswith("_foreach_"): + base_op_name = op_name.replace("_foreach_", "", 1) + elif op_name.startswith("_amp_foreach_"): + base_op_name = op_name.replace("_amp_foreach_", "", 1) + else: + # Fused ops or unknown: keep original op, no translation + target_op = op_schema.op + op_schema = OpSchema( + target_op, # type: ignore[arg-type] + args_schema=tuple(target_args), + kwargs_schema=op_schema.kwargs_schema, + ) + return op_schema, target_output_meta + + # Strip trailing underscore for inplace ops + base_op_name = base_op_name.removesuffix("_") # figure out target op variant variant_map = { @@ -558,7 +632,7 @@ def expanded_foreach_strategy( child_strategies: list[StrategyType] = [] for tensorlist_i in range(tensorlist_len): - per_index_schema, per_index_output_meta = _translate_foreach_op_schema( + per_index_schema, per_index_output_meta = _translate_list_op_schema( op_schema, output_tensor_meta, # type: ignore[arg-type] tensorlist_i, @@ -575,7 +649,8 @@ def expanded_foreach_strategy( return TupleStrategy(children=child_strategies) # TODO maybe this could be helped by adding a new 'tag' to the OpOverload? - if op_schema.op.name().startswith("aten::_foreach_"): + op_name = op_schema.op.name() + if op_name.startswith(("aten::_foreach_", "aten::_amp_foreach_", "aten::_fused_")): return expanded_foreach_strategy return _create_expanded_strategy(op_schema, output_tensor_meta) @@ -586,6 +661,7 @@ def register_single_dim_strategy( schema_info: RuntimeSchemaInfo | None = None, allow_unbacked_sharding: bool | None = None, allow_uneven_sharding: bool = False, + different_mesh_args: list[int] | None = None, ) -> Callable[[_SingleDimStrategyFunc], _SingleDimStrategyFunc]: """ Registers a single_dim_strategy function for the given op. @@ -634,6 +710,7 @@ def wrapper(impl): func=impl, allow_unbacked_sharding=allow_unbacked_sharding, allow_uneven_sharding=allow_uneven_sharding, + different_mesh_args=different_mesh_args, ) registration_wrapper(info) return impl diff --git a/torch/distributed/tensor/_ops/utils.py b/torch/distributed/tensor/_ops/utils.py index 6bb6309feb1a6..e3afc7bf08ded 100644 --- a/torch/distributed/tensor/_ops/utils.py +++ b/torch/distributed/tensor/_ops/utils.py @@ -367,6 +367,7 @@ def expand_to_full_mesh_op_strategy( [list[DTensorSpec], DTensorSpec | tuple[DTensorSpec | None, ...]], bool ] | None = None, + different_mesh_args: list[int] | None = None, ) -> OpStrategy: """ Convenience function to allow writing a sharding strategy considering only a single mesh dimension, @@ -476,6 +477,52 @@ def expand_to_full_mesh_op_strategy( f"input_specs({len(input_specs)}) != strategies({len(input_args_strategy)}: " f"{len(args_strategy)} args + {len(kwargs_strategy)} kwargs)" ) + + # Note [Multi-mesh args] + # + # Some ops accept args whose DTensor lives on a different DeviceMesh + # than the op's primary compute mesh. We call these "multi-mesh + # args". They arise in fused optimizer ops (e.g. _fused_adam_) + # where *state_steps* is a per-rank scalar counter allocated on a + # smaller sub-mesh (e.g. 1-D DP) while params and grads live on a + # larger mesh (e.g. 2-D DP × TP). + # + # Why must these args be Replicate? + # Sharding implies a specific partitioning of a tensor's data + # across the ranks of a mesh. If a tensor doesn't even *exist* + # on the compute mesh, there is no meaningful way to interpret a + # Shard placement for it. Replicate, on the other hand, is + # mesh-agnostic: every rank already holds the full data, so the + # op can simply read the value regardless of which mesh owns it. + # + # What we do here: + # We preserve the original mesh and Replicate placement for these + # args so the propagator does not try to redistribute them onto + # the compute mesh (which would fail or produce wrong results). + # + # This is distinct from the *element_mesh* handling in + # single_dim_strategy.py, which deals with foreach ops where + # different *elements* in a tensor list may live on different + # sub-meshes (e.g. param group A on 2-D mesh, param group B on + # 1-D mesh). + # TODO: refactor fused_ops handling so that there are no longer + # args on different meshes + if different_mesh_args is not None: + for idx in different_mesh_args: + if idx < len(input_args_strategy): + cross_mesh_input = input_args_strategy[idx] + original_spec = cross_mesh_input.strategies[0].output_spec + if original_spec.mesh != mesh: + if not all(p == Replicate() for p in original_spec.placements): + raise RuntimeError( + f"Cross-mesh input at index {idx} must be Replicate, " + f"but got {original_spec.placements}" + ) + input_specs[idx] = DTensorSpec( + mesh=original_spec.mesh, + placements=original_spec.placements, + tensor_meta=original_spec.tensor_meta, + ) self_spec = input_args_strategy[0].strategies[0].output_spec redistribute_input = self_spec.placements != input_specs[0].placements diff --git a/torch/distributed/tensor/_sharding_prop.py b/torch/distributed/tensor/_sharding_prop.py index 6d0040f25b7c9..28be5fe7a17b7 100644 --- a/torch/distributed/tensor/_sharding_prop.py +++ b/torch/distributed/tensor/_sharding_prop.py @@ -53,14 +53,15 @@ def _length(obj) -> int: return len(obj) -def _get_expected_num_tensor_outputs(op: OpOverload) -> int: +def _get_expected_num_tensor_outputs(op: OpOverload) -> int | None: """ Get the expected number of tensor outputs for an operator based on its schema. Returns: The number of tensor outputs expected. Returns 0 for ops that don't return tensors (e.g., _linalg_check_errors). Returns 1 for single tensor return, and >1 for - tuple returns where each element is a tensor. + tuple returns where each element is a tensor. Returns None for List[Tensor] + returns where the length is unknown at schema time. """ return_types = op._schema.returns if len(return_types) == 0: @@ -71,8 +72,8 @@ def _get_expected_num_tensor_outputs(op: OpOverload) -> int: # Could be single tensor or tuple of tensors return len(return_types) elif isinstance(first_return.type, torch.ListType): - # List[Tensor] - we don't know the length at schema time, treat as 1 - return 1 + # List[Tensor] - we don't know the length at schema time + return None else: # Not a tensor return type return 0 @@ -101,6 +102,16 @@ def _validate_tensor_meta_count( else: actual_outputs = len(tensor_meta) + if expected_outputs is None: + # List[Tensor] return type: length unknown at schema time, but + # tensor_meta must be a list of TensorMeta. + if not isinstance(tensor_meta, list): + raise AssertionError( + f"Tensor meta for {op_schema.op} should be a list[TensorMeta] " + f"(op returns List[Tensor]), but got {type(tensor_meta).__name__}" + ) + return + if actual_outputs != expected_outputs: raise AssertionError( f"Tensor meta count mismatch for {op_schema.op}: " @@ -829,7 +840,8 @@ def propagate_op_sharding_non_cached(self, op_schema: OpSchema) -> OutputShardin raise AssertionError selected_strategy = _select_min_cost_strategy(strategy) selected_strategies.append(selected_strategy) - out_spec_list.append(selected_strategy.output_spec) + if selected_strategy.output_specs is not None: + out_spec_list.append(selected_strategy.output_spec) needs_redistribute = False suggestion_args: list[object] = [] From b723cdba34e3462778f3a342f19d8ef7a9b7f92d Mon Sep 17 00:00:00 2001 From: PyTorch MergeBot Date: Tue, 24 Mar 2026 00:19:22 +0000 Subject: [PATCH 0047/1172] Revert "Make Enums opaque (#176541)" This reverts commit 093f7011ee1b8ce246069a2fd88f4083ac160dbb. Reverted https://github.com/pytorch/pytorch/pull/176541 on behalf of https://github.com/georgehong due to breaking internal builds Enums-as-opaque-by-default conflicts with downstream pytree node registrations. ([comment](https://github.com/pytorch/pytorch/pull/176541#issuecomment-4114545835)) --- test/test_opaque_obj_v2.py | 85 ------------------------ torch/_dynamo/variables/builder.py | 15 ++--- torch/_dynamo/variables/script_object.py | 8 --- torch/_library/infer_schema.py | 12 ++-- torch/_library/opaque_object.py | 22 ++---- 5 files changed, 14 insertions(+), 128 deletions(-) diff --git a/test/test_opaque_obj_v2.py b/test/test_opaque_obj_v2.py index 4457913b0090d..5cafd6c573cdb 100644 --- a/test/test_opaque_obj_v2.py +++ b/test/test_opaque_obj_v2.py @@ -1,7 +1,6 @@ # Owner(s): ["module: custom-operators"] import contextlib -import enum import gc import random import unittest @@ -3346,90 +3345,6 @@ def fn(x): res = opt_fn(x) self.assertEqual(ref, res) - def test_enum_export(self): - class Direction(enum.Enum): - UP = 0 - DOWN = 1 - - class Mod(torch.nn.Module): - def forward(self, x, d): - return x + d.value - - ep = torch.export.export(Mod(), (torch.randn(4, 4), Direction.UP), strict=False) - self.assertEqual( - ep.module()(torch.ones(4, 4), Direction.UP), - torch.ones(4, 4) + Direction.UP.value, - ) - self.assertExpectedInline( - normalize_gm(ep.graph_module.print_readable(False)), - """\ -class GraphModule(torch.nn.Module): - def forward(self, x: "f32[4, 4]", d): - add: "f32[4, 4]" = torch.ops.aten.add.Tensor(x, 0); x = None - return (add,) -""", - ) - - backend = EagerAndRecordGraphs() - opt_fn = torch.compile(Mod(), backend=backend) - x = torch.randn(4, 4) - res = opt_fn(x, Direction.UP) - self.assertEqual( - res, - x + Direction.UP.value, - ) - self.assertExpectedInline( - normalize_gm(backend.graphs[0].print_readable(False)), - """\ -class GraphModule(torch.nn.Module): - def forward(self, L_x_: "f32[4, 4]"): - l_x_ = L_x_ - - add: "f32[4, 4]" = l_x_ + 0; l_x_ = None - return (add,) -""", - ) - - def test_enum_custom_op(self): - def get_color(): - class Color(enum.Enum): - RED = 0 - GREEN = 1 - BLUE = 2 - - return Color - - Color = get_color() - - @torch.library.custom_op("test_enum::add_color", mutates_args=()) - def add_color(x: torch.Tensor, c: Color) -> torch.Tensor: - return x + c.value - - @add_color.register_fake - def _(x, c): - return torch.empty_like(x) - - def fn(x, c): - return add_color(x, c) - - x = torch.randn(4, 4) - ref = fn(x, Color.GREEN) - backend = EagerAndRecordGraphs() - opt_fn = torch.compile(fn, backend=backend, fullgraph=True) - res = opt_fn(x, Color.GREEN) - self.assertEqual(ref, res) - self.assertExpectedInline( - normalize_gm(backend.graphs[0].print_readable(False)), - """\ -class GraphModule(torch.nn.Module): - def forward(self, L_x_: "f32[4, 4]"): - l_x_ = L_x_ - - add_color_default: "f32[4, 4]" = torch.ops.test_enum.add_color.default(l_x_, Color.GREEN); l_x_ = None - return (add_color_default,) -""", - ) - instantiate_parametrized_tests(TestOpaqueObject) diff --git a/torch/_dynamo/variables/builder.py b/torch/_dynamo/variables/builder.py index e53d5e3cc10ea..ce2eaaf535500 100644 --- a/torch/_dynamo/variables/builder.py +++ b/torch/_dynamo/variables/builder.py @@ -1523,17 +1523,14 @@ def build_key_value( # ID_MATCH even if its a global variable. self.install_guards(GuardBuilder.CLASS_MATCH) - if isinstance(value, type) and issubclass(value, enum.Enum): - # Order this before OpaqueObjectClassVariable since it's better - # to use the native UserDefinedEnumClassVariable - return UserDefinedEnumClassVariable( + if is_opaque_type(value): + return OpaqueObjectClassVariable( value, source=self.source, ) - if is_opaque_type(value): - assert not (isinstance(value, type) and issubclass(value, enum.Enum)) - return OpaqueObjectClassVariable( + if isinstance(value, type) and issubclass(value, enum.Enum): + return UserDefinedEnumClassVariable( value, source=self.source, ) @@ -4232,7 +4229,7 @@ def create(tx: "InstructionTranslatorBase", value: Any) -> VariableTracker: if isinstance(value, VariableTracker): # This is always valid to call, and useful for recursive calls. return value - elif is_opaque_value_type(type(value)) and not isinstance(value, enum.Enum): + elif is_opaque_value_type(type(value)): return TorchScriptObjectVariable.create(value, value) elif is_opaque_reference_type(type(value)): # This is for handling opaque objects in custom ops @@ -4240,7 +4237,7 @@ def create(tx: "InstructionTranslatorBase", value: Any) -> VariableTracker: tx.output.fake_mode, value ) return TorchScriptObjectVariable.create( - value, # pyrefly: ignore[bad-argument-type] + value, fake_script_obj, ) # type: ignore[attr-defined] diff --git a/torch/_dynamo/variables/script_object.py b/torch/_dynamo/variables/script_object.py index 1f1c0f75fe603..dc959cb77d4cf 100644 --- a/torch/_dynamo/variables/script_object.py +++ b/torch/_dynamo/variables/script_object.py @@ -18,7 +18,6 @@ by limiting operations to known-safe patterns and failing fast for unsafe usage. """ -import enum import functools import inspect import types @@ -90,10 +89,6 @@ class OpaqueObjectClassVariable(UserDefinedVariable): """ def __init__(self, value: Any, **kwargs: Any) -> None: - assert not (isinstance(value, type) and issubclass(value, enum.Enum)), ( - f"Enum class {value} should use UserDefinedEnumClassVariable, " - "not OpaqueObjectClassVariable" - ) super().__init__(**kwargs) self.value = value @@ -236,9 +231,6 @@ def create( ctor_arg_sources: tuple[Source | None, ...] | None = None, **options: Any, ) -> "TorchScriptObjectVariable": - assert not isinstance(value, enum.Enum), ( - f"Enum {type(value)} should use EnumVariable, not TorchScriptObjectVariable" - ) out = TorchScriptObjectVariable( proxy, value, ctor_args_kwargs, ctor_arg_sources=ctor_arg_sources, **options ) diff --git a/torch/_library/infer_schema.py b/torch/_library/infer_schema.py index ae4ee9bd41a31..90a1d162a10a8 100644 --- a/torch/_library/infer_schema.py +++ b/torch/_library/infer_schema.py @@ -8,11 +8,7 @@ from torch import device, dtype, Tensor, types from torch.utils._exposed_in import exposed_in -from .opaque_object import ( - _resolve_opaque_type_info, - is_opaque_reference_type, - is_opaque_type, -) +from .opaque_object import _OPAQUE_TYPES, is_opaque_reference_type, is_opaque_type # This is used as a negative test for @@ -131,7 +127,7 @@ def unstringify_type(ty: type[object] | str) -> tuple[typing.Any, bool]: schema_type = None if annotation_type not in SUPPORTED_PARAM_TYPES: if is_opaque_type(annotation_type): - schema_type = _resolve_opaque_type_info(annotation_type).class_name # type: ignore[union-attr] + schema_type = _OPAQUE_TYPES[annotation_type].class_name elif annotation_type == torch._C.ScriptObject: error_fn( f"Parameter {name}'s type cannot be inferred from the schema " @@ -304,7 +300,7 @@ def parse_return(annotation, error_fn): if origin is not tuple: if annotation not in SUPPORTED_RETURN_TYPES: if is_opaque_reference_type(annotation): - return _resolve_opaque_type_info(annotation).class_name # type: ignore[union-attr] + return _OPAQUE_TYPES[annotation].class_name error_fn( f"Return has unsupported type {annotation}. " f"The valid types are: {SUPPORTED_RETURN_TYPES}." @@ -323,7 +319,7 @@ def parse_return(annotation, error_fn): def _return_type_str(arg): if ty := SUPPORTED_RETURN_TYPES.get(arg): return ty - return _resolve_opaque_type_info(arg).class_name # type: ignore[union-attr] + return _OPAQUE_TYPES[arg].class_name output_ty = ", ".join(_return_type_str(arg) for arg in args) diff --git a/torch/_library/opaque_object.py b/torch/_library/opaque_object.py index e342b5369f960..cfbe0ae814a33 100644 --- a/torch/_library/opaque_object.py +++ b/torch/_library/opaque_object.py @@ -187,9 +187,7 @@ def register_opaque_type( "registered as a pytree. Opaque objects must be pytree leaves." ) - # Value types store the real object directly during tracing (no - # FakeScriptObject wrapper), so they don't need OpaqueBaseMeta. - if typ != "value" and not isinstance(cls, OpaqueBaseMeta): + if not isinstance(cls, OpaqueBaseMeta): raise TypeError( f"Opaque type {cls} must subclass torch._opaque_base.OpaqueBase " "or 'metaclass=torch._opaque_base.OpaqueBaseMeta'. " @@ -204,8 +202,7 @@ def register_opaque_type( ) if typ == "value": - # Enums use identity-based equality (singletons), which is fine for guarding. - if not issubclass(cls, Enum) and cls.__eq__ is object.__eq__: # type: ignore[comparison-overlap] + if cls.__eq__ is object.__eq__: # type: ignore[comparison-overlap] raise TypeError( f"Value-type opaque object of type {cls} is " "expected to have a non-default `__eq__` " @@ -223,14 +220,13 @@ def register_opaque_type( "for FakeTensor caching." ) - # Enums are special-cased in get_opaque_obj_repr. - if not issubclass(cls, Enum) and not hasattr(cls, "__fx_repr__"): + if not hasattr(cls, "__fx_repr__"): raise TypeError( f"Value-type opaque object of type {cls} is " "expected to have a `__fx_repr__` method " "implementation as we will use this to reconstruct " "the object in the FX codegen. __fx_repr__ should return " - "a tuple of (repr_string, dict[str, type])." + "a tuple of (repr_string, set_of_types)." ) if guard_fn is not None: @@ -250,10 +246,6 @@ def register_opaque_type( torch._C._register_opaque_type(name) -# Enums are always opaque value types. -register_opaque_type(Enum, typ="value") - - def is_opaque_value(value: object) -> TypeIs[OpaqueType]: return is_opaque_type(type(value)) @@ -338,12 +330,6 @@ def get_opaque_obj_repr(obj: Any) -> tuple[str, dict[str, type]]: For example, if repr_string is "Foo(bar=Bar(1))", the dict should be: {"Foo": Foo, "Bar": Bar} """ - - # Enums are special cased - if isinstance(obj, Enum): - cls = type(obj) - return f"{cls.__name__}.{obj.name}", {cls.__name__: cls} - if not hasattr(obj, "__fx_repr__"): raise TypeError( f"Value-type opaque object of type {obj} is " From b584edaecbe9e842bc8385eae272fcebc5f1731c Mon Sep 17 00:00:00 2001 From: Anshul Sinha Date: Mon, 23 Mar 2026 12:02:36 -0700 Subject: [PATCH 0048/1172] [dtensor] Register foreach and fused ops via single_dim_strategy (#177187) Enable foreach and fused ops to use the single_dim_strategy path: - Add foreach ops to existing category lists (binary_additive_ops, binary_mul_ops, etc.) and the pointwise_ops list - Add fused ops to pointwise_ops list for unified registration - Add _is_list_op() helper to detect foreach/fused/amp_foreach ops - Modify _register_single_dim_pointwise to use needs_pytree for list ops and cross_mesh_indices for fused ops - Remove separate register_single_dim_strategy loop for fused ops Tests cover multi-tensor foreach lists, mixed placements across list elements, same-mesh fused adam, and cross-mesh fused adam exercising the cross_mesh_indices code path. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177187 Approved by: https://github.com/wconstab ghstack dependencies: #177186 --- test/distributed/tensor/test_pointwise_ops.py | 318 ++++++++++++++++++ .../distributed/tensor/_ops/_pointwise_ops.py | 192 ++++++++--- 2 files changed, 455 insertions(+), 55 deletions(-) diff --git a/test/distributed/tensor/test_pointwise_ops.py b/test/distributed/tensor/test_pointwise_ops.py index 6a1646d0f93b1..51360d22e63d6 100644 --- a/test/distributed/tensor/test_pointwise_ops.py +++ b/test/distributed/tensor/test_pointwise_ops.py @@ -12,6 +12,7 @@ DeviceMesh, distribute_tensor, DTensor, + init_device_mesh, Partial, Placement, Replicate, @@ -19,6 +20,7 @@ ) from torch.distributed.tensor._ops._math_ops import _NormPartial from torch.distributed.tensor.debug import CommDebugMode +from torch.distributed.tensor.placement_types import _StridedShard from torch.testing._internal.common_utils import ( instantiate_parametrized_tests, parametrize, @@ -918,6 +920,322 @@ def test_inplace_add_replicate_with_partial_avg_requires_comm(self): ) self.assertEqual(d_self.full_tensor(), expected) + @with_comms + @parametrize( + "op_fn,second_arg,placement,expected", + [ + # Binary list ops preserve Partial("sum") + (torch._foreach_add, "list", Partial("sum"), Partial("sum")), + # Scalar mul preserves Partial("sum") (linearity) + (torch._foreach_mul, 2.0, Partial("sum"), Partial("sum")), + # Neg preserves Partial("sum") (linearity) + (torch._foreach_neg, None, Partial("sum"), Partial("sum")), + # Binary list ops preserve Shard(0) + (torch._foreach_add, "list", Shard(0), Shard(0)), + ], + ) + def test_foreach_placement_propagation( + self, op_fn, second_arg, placement, expected + ): + device_mesh = self.build_device_mesh() + comm_mode = CommDebugMode() + + shapes = [(8, 8), (4, 4), (2, 6)] + if placement.is_shard(): + dts = [ + distribute_tensor( + torch.rand(s, device=self.device_type), device_mesh, [placement] + ) + for s in shapes + ] + else: + dts = [ + DTensor.from_local( + torch.rand(s, device=self.device_type), device_mesh, [placement] + ) + for s in shapes + ] + + if second_arg == "list": + if placement.is_shard(): + args = [ + distribute_tensor( + torch.rand(s, device=self.device_type), + device_mesh, + [placement], + ) + for s in shapes + ] + else: + args = [ + DTensor.from_local( + torch.rand(s, device=self.device_type), + device_mesh, + [placement], + ) + for s in shapes + ] + call_args = (dts, args) + elif second_arg is None: + call_args = (dts,) + else: + call_args = (dts, second_arg) + + with comm_mode: + result = op_fn(*call_args) + + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(len(result), len(shapes)) + for r in result: + self.assertIsInstance(r, DTensor) + self.assertEqual(r.placements, (expected,)) + + @with_comms + def test_foreach_mixed_placements(self): + """Each element in a foreach list independently preserves its placement.""" + device_mesh = self.build_device_mesh() + comm_mode = CommDebugMode() + + dt1 = DTensor.from_local( + torch.rand(8, 8, device=self.device_type), + device_mesh, + [Partial("sum")], + ) + ds1 = DTensor.from_local( + torch.rand(8, 8, device=self.device_type), + device_mesh, + [Partial("sum")], + ) + dt2 = distribute_tensor( + torch.rand(8, 4, device=self.device_type), device_mesh, [Shard(0)] + ) + ds2 = distribute_tensor( + torch.rand(8, 4, device=self.device_type), device_mesh, [Shard(0)] + ) + + with comm_mode: + result = torch._foreach_add([dt1, dt2], [ds1, ds2]) + + self.assertEqual(comm_mode.get_total_counts(), 0) + self.assertEqual(result[0].placements, (Partial("sum"),)) + self.assertEqual(result[1].placements, (Shard(0),)) + + @with_comms + @skip_unless_torch_gpu + @parametrize( + "mesh_type,amsgrad", + [ + # All tensors on the same 1D mesh + ("same", False), + # state_steps on a different sub-mesh, with max_exp_avg_sqs present + ("cross", True), + # state_steps on a different sub-mesh, empty max_exp_avg_sqs + # (triggers the different_mesh_args index remapping) + ("cross", False), + ], + ) + def test_fused_adam(self, mesh_type, amsgrad): + if mesh_type == "same": + param_mesh = self.build_device_mesh() + step_mesh = param_mesh + else: + mesh_2d = init_device_mesh( + self.device_type, + (2, self.world_size // 2), + mesh_dim_names=("dp", "tp"), + ) + param_mesh = mesh_2d["tp"] + step_mesh = mesh_2d["dp"] + + d_params = [ + distribute_tensor( + torch.rand(8, 8, device=self.device_type), param_mesh, [Shard(0)] + ) + ] + d_grads = [ + distribute_tensor( + torch.rand(8, 8, device=self.device_type), param_mesh, [Shard(0)] + ) + ] + d_exp_avgs = [ + distribute_tensor( + torch.zeros(8, 8, device=self.device_type), param_mesh, [Shard(0)] + ) + ] + d_exp_avg_sqs = [ + distribute_tensor( + torch.zeros(8, 8, device=self.device_type), param_mesh, [Shard(0)] + ) + ] + d_max_exp_avg_sqs = ( + [ + distribute_tensor( + torch.zeros(8, 8, device=self.device_type), + param_mesh, + [Shard(0)], + ) + ] + if amsgrad + else [] + ) + d_state_steps = [ + distribute_tensor( + torch.tensor(1.0, device=self.device_type), step_mesh, [Replicate()] + ) + ] + + torch._fused_adam_( + d_params, + d_grads, + d_exp_avgs, + d_exp_avg_sqs, + d_max_exp_avg_sqs, + d_state_steps, + lr=0.001, + beta1=0.9, + beta2=0.999, + weight_decay=0.0, + eps=1e-8, + amsgrad=amsgrad, + maximize=False, + ) + + self.assertIsInstance(d_params[0], DTensor) + self.assertEqual(d_params[0].placements, (Shard(0),)) + self.assertEqual(d_params[0].device_mesh, param_mesh) + + @with_comms + @skip_unless_torch_gpu + def test_fused_adamw_strided_shard_cross_mesh(self): + """Reproduce the torchtitan pattern: params with _StridedShard on a 2D + mesh, state_steps with Replicate on a 1D DP sub-mesh.""" + mesh_2d = init_device_mesh( + self.device_type, + (2, self.world_size // 2), + mesh_dim_names=("dp", "tp"), + ) + dp_mesh = mesh_2d["dp"] + placements = [_StridedShard(0, split_factor=2), Shard(0)] + + d_params = [ + DTensor.from_local( + torch.randn(4, 8, device=self.device_type), + mesh_2d, + placements, + run_check=False, + ) + ] + d_grads = [ + DTensor.from_local( + torch.randn(4, 8, device=self.device_type), + mesh_2d, + placements, + run_check=False, + ) + ] + d_exp_avgs = [ + DTensor.from_local( + torch.zeros(4, 8, device=self.device_type), + mesh_2d, + placements, + run_check=False, + ) + ] + d_exp_avg_sqs = [ + DTensor.from_local( + torch.zeros(4, 8, device=self.device_type), + mesh_2d, + placements, + run_check=False, + ) + ] + d_state_steps = [ + DTensor.from_local( + torch.tensor(1.0, device=self.device_type), + dp_mesh, + [Replicate()], + run_check=False, + ) + ] + + torch.ops.aten._fused_adamw_.default( + d_params, + d_grads, + d_exp_avgs, + d_exp_avg_sqs, + [], + d_state_steps, + lr=0.1, + beta1=0.9, + beta2=0.999, + weight_decay=0.01, + eps=1e-8, + amsgrad=False, + maximize=False, + ) + + self.assertIsInstance(d_params[0], DTensor) + self.assertEqual(d_params[0].placements, tuple(placements)) + self.assertEqual(d_params[0].device_mesh, mesh_2d) + + @with_comms + def test_fused_adamw_tensor_lr(self): + """Test _fused_adamw_ with tensor lr kwarg.""" + device_mesh = self.build_device_mesh() + + d_params = [ + distribute_tensor( + torch.rand(8, 8, device=self.device_type), device_mesh, [Shard(0)] + ) + ] + d_grads = [ + distribute_tensor( + torch.rand(8, 8, device=self.device_type), device_mesh, [Shard(0)] + ) + ] + d_exp_avgs = [ + distribute_tensor( + torch.zeros(8, 8, device=self.device_type), device_mesh, [Shard(0)] + ) + ] + d_exp_avg_sqs = [ + distribute_tensor( + torch.zeros(8, 8, device=self.device_type), device_mesh, [Shard(0)] + ) + ] + d_state_steps = [ + distribute_tensor( + torch.tensor(1.0, device=self.device_type), + device_mesh, + [Replicate()], + ) + ] + d_lr = distribute_tensor( + torch.tensor(0.001, device=self.device_type), + device_mesh, + [Replicate()], + ) + + # tensor_lr variant passes lr as a tensor kwarg + torch.ops.aten._fused_adamw_.tensor_lr( + d_params, + d_grads, + d_exp_avgs, + d_exp_avg_sqs, + [], + d_state_steps, + lr=d_lr, + beta1=0.9, + beta2=0.999, + weight_decay=0.01, + eps=1e-8, + amsgrad=False, + maximize=False, + ) + + self.assertIsInstance(d_params[0], DTensor) + self.assertEqual(d_params[0].placements, (Shard(0),)) + instantiate_parametrized_tests(DistElementwiseOpsTest) diff --git a/torch/distributed/tensor/_ops/_pointwise_ops.py b/torch/distributed/tensor/_ops/_pointwise_ops.py index 1e8c5b28a56e0..ff33261cae434 100644 --- a/torch/distributed/tensor/_ops/_pointwise_ops.py +++ b/torch/distributed/tensor/_ops/_pointwise_ops.py @@ -26,7 +26,6 @@ infer_broadcast_dims_map, map_placements_after_broadcast, normalize_dim, - register_op_strategy, ) from torch.distributed.tensor.placement_types import ( _StridedShard, @@ -66,7 +65,12 @@ def _common_pointwise_single_dim_strategy( ) -> Callable[ [OpOverload, ArgsType, KwargsType], list[list[Placement | _ShardingPlaceholder]] ]: - """Factory for single-dim strategies that add partial placement rules.""" + """Factory for single-dim strategies that add partial placement rules. + + Returns strategies shaped [output, *args] only. Tensor kwarg placements + (e.g. ``out``, ``lr``) are appended by the wrapper in + ``_register_single_dim_pointwise``. + """ def strategy( op: OpOverload, @@ -116,46 +120,61 @@ def strategy( return strategy +def _is_list_op(op: OpOverload) -> bool: + """Returns True if op is a foreach, amp_foreach, or fused op.""" + name = op.name() + return name.startswith(("aten::_foreach_", "aten::_amp_foreach_", "aten::_fused_")) + + +# The state_steps arg of fused adam / adamw is a Replicate scalar tensor, which will be put on +# the compute_mesh of an op across all parameter groups, even when not all parameter groups +# are on the same device mesh. This idx will help avoid hitting exceptions or unnecessary +# redistribute during sharding propagation. +_FUSED_OP_SCALAR_IDX = 5 + + def _register_single_dim_pointwise( op: OpOverload, partial_extra_rules: list[list[Placement]] | None = None, static_argnum: int = 0, ) -> None: - strategy_fn = _common_pointwise_single_dim_strategy( + inner_fn = _common_pointwise_single_dim_strategy( partial_extra_rules=partial_extra_rules # pyrefly: ignore[bad-argument-type] ) - # For .out ops, append output placement as the out kwarg placement. - # Strategy functions author [output, *args] without kwargs. The out tensor - # must match the output placement, so we duplicate strategy[0] (output). - # This makes strategies [output, *args, out_kwarg] so _get_num_tensor_inputs - # (which counts the out kwarg) computes num_outputs correctly. - if "out" in op._schema.overload_name: - inner_fn = strategy_fn - - def _out_wrapper( - op: OpOverload, - args: ArgsType, - kwargs: KwargsType, - _fn: Callable = inner_fn, - ) -> list[list[Placement | _ShardingPlaceholder]]: - strategies = _fn(op, args, kwargs) - n_tensor_args = sum(1 for a in args if isinstance(a, TensorMeta)) - n_outputs = sum(1 for r in op._schema.returns if "Tensor" in str(r.type)) - for s in strategies: - if len(s) != n_outputs + n_tensor_args: - raise AssertionError( - f"Strategy length {len(s)} != expected {n_outputs + n_tensor_args} " - f"({n_outputs} output(s) + {n_tensor_args} args) for {op}. " - f"out kwarg will be appended by infra." - ) - return [s + [s[0]] for s in strategies] - strategy_fn = _out_wrapper + # Wrap to append tensor kwarg placements in schema declaration order. + # out = output placement (s[0]); everything else (e.g. lr) = Replicate. + # TODO: move kwargs handling upstream if this works + def strategy_fn( + op: OpOverload, + args: ArgsType, + kwargs: KwargsType, + _fn: Callable = inner_fn, + ) -> list[list[Placement | _ShardingPlaceholder]]: + strategies = _fn(op, args, kwargs) + kw_names = [k for k, v in kwargs.items() if isinstance(v, TensorMeta)] + if not kw_names: + return strategies + return [ + s + [s[0] if name == "out" else Replicate() for name in kw_names] + for s in strategies + ] + + if _is_list_op(op): + schema_info = RuntimeSchemaInfo(needs_pytree=True) + else: + schema_info = RuntimeSchemaInfo(static_argnum, static_kwargkey=["out"]) + # Fused ops (e.g. _fused_adam_) have state_steps on a potentially different + # mesh; see the note in expand_to_full_mesh_op_strategy for details. + different_mesh_args: list[int] | None = None + if op.name().startswith("aten::_fused_"): + different_mesh_args = [_FUSED_OP_SCALAR_IDX] register_single_dim_strategy( op, - schema_info=RuntimeSchemaInfo(static_argnum, static_kwargkey=["out"]), + schema_info=schema_info, allow_uneven_sharding=True, allow_unbacked_sharding=True, + different_mesh_args=different_mesh_args, )(strategy_fn) @@ -171,6 +190,11 @@ def _out_wrapper( aten.sub.Tensor, aten.sub_.Tensor, aten.sub.out, + # foreach variants + aten._foreach_add.List, + aten._foreach_add_.List, + aten._foreach_sub.List, + aten._foreach_sub_.List, ] _BINARY_ADDITIVE_RULES: list[list[Placement]] = [ @@ -191,8 +215,26 @@ def _out_wrapper( _register_single_dim_pointwise(op, _BINARY_ADDITIVE_RULES) # mul: partials propagate through either arg. div: only through numerator. -binary_mul_ops = [aten.mul.Tensor, aten.mul_.Tensor, aten.mul.out] -binary_div_ops = [aten.div.Tensor, aten.div_.Tensor, aten.div.out] +binary_mul_ops = [ + aten.mul.Tensor, + aten.mul_.Tensor, + aten.mul.out, + # foreach variants + aten._foreach_mul.List, + aten._foreach_mul_.List, + aten._foreach_mul.Tensor, + aten._foreach_mul_.Tensor, +] +binary_div_ops = [ + aten.div.Tensor, + aten.div_.Tensor, + aten.div.out, + # foreach variants + aten._foreach_div.List, + aten._foreach_div_.List, + aten._foreach_div.Tensor, + aten._foreach_div_.Tensor, +] # _UNARY_LINEAR_RULES handles the scalar promotion case: Python's __mul__/__truediv__ # promote scalars to 0-dim tensors, so aten.mul.Scalar dispatches as aten.mul.Tensor @@ -220,6 +262,15 @@ def _out_wrapper( aten.div_.Scalar, aten.mul.Scalar, aten.mul_.Scalar, + # foreach variants + aten._foreach_div.Scalar, + aten._foreach_div_.Scalar, + aten._foreach_mul.Scalar, + aten._foreach_mul_.Scalar, + aten._foreach_div.ScalarList, + aten._foreach_div_.ScalarList, + aten._foreach_mul.ScalarList, + aten._foreach_mul_.ScalarList, ] for op in scalar_linear_ops: @@ -293,6 +344,11 @@ def _out_wrapper( aten.nan_to_num.default, aten.nan_to_num_.default, aten.nan_to_num.out, + # foreach variants + aten._foreach_exp.default, + aten._foreach_exp_.default, + aten._foreach_clamp_max_.Scalar, + aten._foreach_clamp_min_.Scalar, ] _NON_DECREASING_RULES: list[list[Placement]] = [ @@ -322,7 +378,14 @@ def _out_wrapper( _register_single_dim_pointwise(op, _NON_INCREASING_RULES) # neg is linear: -(A1 + A2) = -A1 + -A2 -neg_ops = [aten.neg.default, aten.neg_.default, aten.neg.out] +neg_ops = [ + aten.neg.default, + aten.neg_.default, + aten.neg.out, + # foreach variants + aten._foreach_neg.default, + aten._foreach_neg_.default, +] _NEG_RULES: list[list[Placement]] = _UNARY_LINEAR_RULES + _NON_INCREASING_RULES @@ -383,6 +446,8 @@ def _out_wrapper( aten.maximum.default, aten.maximum.out, prims.fmax.default, + # foreach variants + aten._foreach_maximum_.List, ] _MONOTONE_MAX_PRESERVING_BINARY_BASE_RULES: list[list[Placement]] = [ @@ -721,6 +786,45 @@ def _out_wrapper( prims.ne.default, prims.spherical_bessel_j0.default, prims.zeta.default, + # foreach variants + aten._foreach_abs.default, + aten._foreach_abs_.default, + aten._foreach_addcdiv_.Scalar, + aten._foreach_addcdiv_.ScalarList, + aten._foreach_addcdiv_.Tensor, + aten._foreach_addcmul.Scalar, + aten._foreach_addcmul_.Scalar, + aten._foreach_addcmul_.ScalarList, + aten._foreach_addcmul_.Tensor, + aten._foreach_lerp_.Scalar, + aten._foreach_pow.List, + aten._foreach_pow.ScalarList, + aten._foreach_reciprocal_.default, + aten._foreach_sub.Scalar, + aten._foreach_sub_.Scalar, + aten._foreach_sub.ScalarList, + aten._foreach_sub_.ScalarList, + aten._foreach_sqrt.default, + aten._foreach_sqrt_.default, + aten._foreach_zero_.default, + aten._foreach_cos.default, + aten._foreach_cos_.default, + aten._foreach_log.default, + aten._foreach_log_.default, + aten._amp_foreach_non_finite_check_and_unscale_.default, + # foreach linearity variants + aten._foreach_add.Scalar, + aten._foreach_add_.Scalar, + aten._foreach_add_.ScalarList, + # fused optimizer ops + aten._fused_adam_.default, + aten._fused_adam.default, + aten._fused_adam.tensor_lr, + aten._fused_adam_.tensor_lr, + aten._fused_adamw_.default, + aten._fused_adamw.default, + aten._fused_adamw.tensor_lr, + aten._fused_adamw_.tensor_lr, ] @@ -1227,16 +1331,6 @@ def list_linear_pointwise_strategy(op_schema: OpSchema) -> StrategyType: return list_pointwise_strategy(op_schema, linearity=True) -for op in for_each_ops: - register_op_strategy(op, schema_info=RuntimeSchemaInfo(needs_pytree=True))( - list_pointwise_strategy - ) - -for op in for_each_linearity_ops: - register_op_strategy(op, schema_info=RuntimeSchemaInfo(needs_pytree=True))( - list_linear_pointwise_strategy - ) - fused_ops = [ aten._fused_adam_.default, aten._fused_adam.default, @@ -1249,18 +1343,6 @@ def list_linear_pointwise_strategy(op_schema: OpSchema) -> StrategyType: ] -# The state_steps arg of fused adam / adamw is a Replicate scalar tensor, which will be put on -# the compute_mesh of an op across all parameter groups, even when not all parameter groups -# are on the same device mesh. This idx will help avoid hitting exceptions or unnecessary -# redistribute during sharding propagation. -_FUSED_OP_SCALAR_IDX = 5 - -for op in fused_ops: - register_op_strategy(op, schema_info=RuntimeSchemaInfo(needs_pytree=True))( - list_pointwise_strategy - ) - - def register_inductor_prims() -> None: """Register DTensor sharding strategies for inductor prims ops. From c311aed673b92ee8ae87ae6ef805f7e94d6b2fe7 Mon Sep 17 00:00:00 2001 From: Anshul Sinha Date: Mon, 23 Mar 2026 12:02:37 -0700 Subject: [PATCH 0049/1172] [dtensor] Remove dead code from _pointwise_ops.py (#177208) The old strategy registration path (register_op_strategy-based functions and data structures) was superseded by the register_single_dim_strategy infrastructure but never cleaned up. This removes ~550 lines of dead code including: pointwise_strategy, linear_pointwise_strategy, copy_strategy, common_pointwise_strategy, single_mesh_dim_* strategy functions, list_pointwise_strategy, list_linear_pointwise_strategy, for_each_ops/for_each_linearity_ops/fused_ops lists, and their associated helper sets and unused imports. Authored with Claude. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177208 Approved by: https://github.com/Skylion007 ghstack dependencies: #177186, #177187 --- .../distributed/tensor/_ops/_pointwise_ops.py | 559 +----------------- 1 file changed, 5 insertions(+), 554 deletions(-) diff --git a/torch/distributed/tensor/_ops/_pointwise_ops.py b/torch/distributed/tensor/_ops/_pointwise_ops.py index ff33261cae434..b814a3b052f9c 100644 --- a/torch/distributed/tensor/_ops/_pointwise_ops.py +++ b/torch/distributed/tensor/_ops/_pointwise_ops.py @@ -1,41 +1,16 @@ # Copyright (c) Meta Platforms, Inc. and affiliates -import functools -from collections.abc import Callable, Sequence -from typing import cast +from collections.abc import Callable import torch from torch._ops import OpOverload -from torch.distributed.tensor._dtensor_spec import DTensorSpec, TensorMeta -from torch.distributed.tensor._op_schema import ( - ArgsType, - KwargsType, - OpSchema, - OpSpec, - OpStrategy, - RuntimeSchemaInfo, - StrategyType, - TupleStrategy, -) -from torch.distributed.tensor._ops._math_ops import _NormPartial +from torch.distributed.tensor._dtensor_spec import TensorMeta +from torch.distributed.tensor._op_schema import ArgsType, KwargsType, RuntimeSchemaInfo from torch.distributed.tensor._ops.single_dim_strategy import ( _ShardingPlaceholder, register_single_dim_strategy, ) -from torch.distributed.tensor._ops.utils import ( - generate_redistribute_costs, - infer_broadcast_dims_map, - map_placements_after_broadcast, - normalize_dim, -) -from torch.distributed.tensor.placement_types import ( - _StridedShard, - Partial, - Placement, - Replicate, - Shard, -) -from torch.types import _Number -from torch.utils._typing_utils import not_none +from torch.distributed.tensor._ops.utils import infer_broadcast_dims_map +from torch.distributed.tensor.placement_types import Partial, Placement, Replicate aten = torch.ops.aten @@ -276,8 +251,6 @@ def strategy_fn( for op in scalar_linear_ops: _register_single_dim_pointwise(op, _UNARY_LINEAR_RULES, static_argnum=1) -neg_ops = [aten.neg.default, aten.neg_.default] - # Non-decreasing unary ops: f(max(a,b)) = max(f(a),f(b)). # Only ops that are non-decreasing on their ENTIRE domain belong here. # Ops with restricted domains (e.g. log on (0,∞), asin on [-1,1]) do NOT qualify @@ -477,17 +450,6 @@ def strategy_fn( _register_single_dim_pointwise(op, _MONOTONE_MIN_PRESERVING_BINARY_BASE_RULES) -# The linear pointwise ops map, key is op, value is the type of linearity. -# Reconstructed from category lists for the existing registration path. -linear_pointwise_ops: dict[OpOverload, int] = { - aten.to.dtype: 0, - **dict.fromkeys(binary_additive_ops, 1), - **dict.fromkeys(binary_mul_ops, 2), - **dict.fromkeys(binary_div_ops, 2), - **dict.fromkeys(scalar_linear_ops, 0), - **dict.fromkeys(neg_ops, 0), -} - pointwise_ops = [ # please keep the entries below alphabetically sorted aten.__ilshift__.Scalar, @@ -828,520 +790,9 @@ def strategy_fn( ] -# Reconstruct the original linear_pointwise_ops dict for the existing registration path. -linear_pointwise_ops: dict[OpOverload, int] = { - **dict.fromkeys(unary_linear_ops, 0), - **dict.fromkeys(binary_additive_ops, 1), - **dict.fromkeys(binary_mul_ops, 2), - **dict.fromkeys(binary_div_ops, 2), - **dict.fromkeys(scalar_linear_ops, 0), - **dict.fromkeys(neg_ops, 0), -} - - -def pointwise_strategy( - op_schema: OpSchema, - linearity: int = -1, - preserve_partial: str | None = None, -) -> OpStrategy: - """Strategy for pointwise ops on the old registration path.""" - followed_strategy_index = -1 - max_shards = -1 - max_ndim = -1 - - if op_schema.is_inplace_op(): - # inplace op should follow the first arg strategy - followed_strategy = op_schema.args_schema[0] - followed_strategy_index = 0 - elif op_schema.is_out_variant_op(): - # out variant op should follow the out kwarg strategy - followed_strategy = op_schema.kwargs_schema["out"] - # out variant is technically a kwarg for the strategy to follow so it does not - # have an "index", we set it to a reasonably large number just to indicate it's - # not a valid index - followed_strategy_index = 100 - else: - # normal pointwise op, we choose to follow the arg with - # the max shards in case operands needs reshard - # in case of multiple operands with max shard, we take - # the one with the max number of dimensions - for idx, arg_strategy in enumerate(op_schema.args_schema): - if not isinstance(arg_strategy, OpStrategy): - continue - arg_max_shards = arg_strategy.max_num_shards() - arg_max_ndim = arg_strategy.ndim - if (arg_max_shards > max_shards) or ( - arg_max_shards == max_shards and arg_max_ndim > max_ndim - ): - followed_strategy_index = idx - max_shards = arg_max_shards - max_ndim = arg_max_ndim - followed_strategy = op_schema.args_schema[followed_strategy_index] - - if not isinstance(followed_strategy, OpStrategy): - raise AssertionError(f"no strategy to follow for {op_schema}!") - return common_pointwise_strategy( - op_schema.op, - op_schema.args_schema, - followed_strategy, - followed_strategy_index, - linearity, - preserve_partial=preserve_partial, - ) - - -def linear_pointwise_strategy(op_schema: OpSchema) -> StrategyType: - """ - Linear pointwise operators can propagate pending reductions. - For example, c = add(a, b); if a is pending sum, then c will be - pending sum as well without any communication overhead. - - Note that: - 1. Only unary and binary operations are supported, out variant - ops are not supported. - 2. There're multiple types of linearity, refer to the doc of - common_pointwise_strategy for more details. - """ - linearity_type = linear_pointwise_ops.get(op_schema.op, -1) - return pointwise_strategy(op_schema, linearity=linearity_type) - - -def single_mesh_dim_pointwise_strategy( - op: OpOverload, - args_schema: ArgsType, - kwargs_schema: KwargsType, - linearity: int = -1, -) -> list[list[Placement | _ShardingPlaceholder]]: - return single_mesh_dim_common_pointwise_strategy(args_schema, linearity) - - -def single_mesh_dim_linear_pointwise_strategy( - linearity: int = -1, -) -> Callable[ - [OpOverload, ArgsType, KwargsType], list[list[Placement | _ShardingPlaceholder]] -]: - return functools.partial(single_mesh_dim_pointwise_strategy, linearity=linearity) - - -def single_mesh_dim_common_pointwise_strategy( - args_schema: ArgsType, - linearity: int = -1, - scalar_tensor_idx: int | None = None, -) -> list[list[Placement | _ShardingPlaceholder]]: - # TODO rename - tensor_arg_strategies: list[TensorMeta] = [ - arg for arg in args_schema if isinstance(arg, TensorMeta) - ] - common_shape = torch.broadcast_shapes( - *[arg.shape for arg in args_schema if isinstance(arg, TensorMeta)] - ) - placements_list: list[list[Placement | _ShardingPlaceholder]] = [] - for i in range(len(common_shape)): - # Shard output dim i, and then shard the corresponding arguments if they have a corresponding (non broadcast) dim - shard_placements: list[Placement | _ShardingPlaceholder] = [ - _ShardingPlaceholder(i) - ] - for arg in tensor_arg_strategies: - common_dim_to_arg_dim = infer_broadcast_dims_map(common_shape, arg.shape) - if common_dim_to_arg_dim[i] >= 0: - shard_placements.append(_ShardingPlaceholder(common_dim_to_arg_dim[i])) - else: - shard_placements.append(Replicate()) - - placements_list.append(shard_placements) - - if linearity == 0: - # unary op (e.g. to_copy), and also binary ops like mul.scalar - # input, output can be partial - if len(tensor_arg_strategies) != 1: - raise AssertionError("expected single tensor input for linearity==0 op") - placements_list.append([Partial("sum"), Partial("sum")]) - # TODO: do i need to check scalar_tensor_index and assign a replicate to that one, or do i omit a placement for it - # TODO: can mul.scalar work with avg or only sum? i think only sum works. common_pointwise_strategy seems - # to support both. - # TODO: also, i'll be replacing 'Partial(sum)' here with some kind of 'PartialPlaceholder', not yet designed - placements_list.append([Partial("avg"), Partial("avg")]) - - elif linearity == 1: - # binary add ops - # (A1 + B1) + (A2 + B2) == (A1 + A2) + (B1 + B2) - if len(tensor_arg_strategies) != 2: - raise AssertionError("expected two tensor inputs for linearity==1 op") - placements_list.append([Partial("sum"), Partial("sum"), Partial("sum")]) - elif linearity == 2: - # binary mul ops (2 tensor inputs) - # (A * B1) + (A * B2) == A * (B1 + B2) - if len(tensor_arg_strategies) != 2: - raise AssertionError("expected two tensor inputs for linearity==2 op") - placements_list.append([Partial("sum"), Partial("sum"), Replicate()]) - placements_list.append([Partial("sum"), Replicate(), Partial("sum")]) - - # TODO: handle scalar_tensor_idx - return placements_list - - -def copy_strategy(op_schema: OpSchema) -> StrategyType: - """ - Strategy for copy_ that preserves any Partial placement. - - copy_ simply copies data and should preserve whatever Partial placement - the destination has, regardless of the reduce_op type (sum, avg, max, min, etc.). - """ - return pointwise_strategy(op_schema, preserve_partial="all") - - -def common_pointwise_strategy( - op, - args_schema: Sequence[object], - followed_strategy: OpStrategy, - followed_strategy_index: int, - linearity: int = -1, - scalar_tensor_idx: int | None = None, - preserve_partial: str | None = None, -) -> OpStrategy: - """ - Common strategy for pointwise operations. - - Args: - args_schema: Input arguments schema - followed_strategy: Strategy to follow for output placement - followed_strategy_index: Index of the strategy being followed - linearity: depending on the operator, we support different types of linearity - -1: the operation does not support linearity - 0: the unary operation that supports linearity, output propagates partial. - 1: the binary operation supports add linearity, where it requires every operand - to be partial, output propagates partial. - 2: the binary operation supports multiplicative linearity, where it requires - the primary operand to be partial, and the other operands to be replicate, - output propagates partial. - scalar_tensor_idx: Index of the Replicate scalar tensor for which we allow the mesh - to be different from the mesh of followed_strategy - preserve_partial: If set, Partial placements with this reduce_op will be preserved - through the operation (e.g., "max" for torch.maximum, "min" for torch.minimum). - """ - # handle broadcasting - common_shape = torch.broadcast_shapes( - *[arg.shape for arg in args_schema if isinstance(arg, OpStrategy)] - ) - pointwise_strategy = OpStrategy([]) - - for op_spec in followed_strategy.strategies: - spec_to_follow = op_spec.output_spec - - out_placements: list[Placement] = [] - for placement in spec_to_follow.placements: - if isinstance(placement, Shard | _StridedShard): - shard_dim = normalize_dim(placement.dim, len(spec_to_follow.shape)) - common_ndim = len(common_shape) - new_shard_dim = common_ndim - len(spec_to_follow.shape) + shard_dim - if isinstance(placement, _StridedShard): - out_placements.append( - _StridedShard( - new_shard_dim, split_factor=placement.split_factor - ) - ) - else: - out_placements.append(Shard(new_shard_dim)) - elif isinstance(placement, Partial): - is_scalar_arg = any(isinstance(arg, _Number) for arg in args_schema) - propagate_partial = False - - # ordering matters here since NormPartial is a subclass of Partial - if isinstance(placement, _NormPartial): - # explanation for args_schema[1] >= 0 can be found in summary - # https://github.com/pytorch/pytorch/pull/170035 - propagate_partial = ( - op in norm_partial_avoidable_redistribute_ops - and args_schema[1] >= 0 # pyre-ignore[unsupported-operation] - ) - - elif isinstance(placement, Partial): - propagate_partial = not ( - op in p_sum_scalar_redistribute_ops and is_scalar_arg - ) - - # Check if this partial type should be preserved - # preserve_partial="all" preserves any Partial type (used for copy_) - if preserve_partial == "all": - out_placements.append(placement) - elif preserve_partial is not None and placement.is_partial( - preserve_partial - ): - out_placements.append(placement) - # note that only partial-sum and partial-avg are supported for linearity - elif ( - linearity >= 0 - and (placement.is_partial("sum") or placement.is_partial("avg")) - and propagate_partial - ): - # propagate the partial placement - out_placements.append(placement) - else: - # clear the partial placement if op does not support linearity - # by default we just replicate the partial, need to see if this - # is optimal for all cases - out_placements.append(Replicate()) - else: - out_placements.append(placement) - - input_specs: list[DTensorSpec] = [] - redistribute_costs: list[list[float]] = [] - for input_idx, input_arg in enumerate(args_schema): - if isinstance(input_arg, OpStrategy): - input_arg_spec = input_arg.strategies[0].output_spec - - # sanity check that all args that follow the same strategy - # are on the same DeviceMesh - if input_arg.mesh != followed_strategy.mesh: - # For the scalar tensor arg in fused ops, do not follow followed_strategy; - # instead, let the input mesh and the Replicate placements propagate through. - if input_idx == scalar_tensor_idx: - if not all(p == Replicate() for p in input_arg_spec.placements): - raise AssertionError - input_arg_target_spec = DTensorSpec( - mesh=input_arg.mesh, - placements=input_arg_spec.placements, - tensor_meta=input_arg_spec.tensor_meta, - ) - input_specs.append(input_arg_target_spec) - redistribute_costs.append( - generate_redistribute_costs( - input_arg, input_arg_target_spec - ) - ) - continue - else: - raise ValueError( - f"Could not run pointwise computation across different mesh: " - f"Found {input_arg.mesh} and {followed_strategy.mesh}!" - ) - - # every arg follow the out_placements, but need to handle broadcasting - input_arg_dims_map = infer_broadcast_dims_map( - common_shape, input_arg_spec.shape - ) - - # Determine if this input should convert Partial to Replicate based on linearity - should_convert_partial = ( - linearity == 2 - and input_idx - != followed_strategy_index # Don't convert the "followed" strategy - ) - - # For preserve_partial ops, check if non-followed input has incompatible - # Partial type. If so, it must be redistributed to Replicate first. - if ( - preserve_partial is not None - and input_idx != followed_strategy_index - ): - for out_p, in_p in zip(out_placements, input_arg_spec.placements): - if ( - isinstance(out_p, Partial) - and isinstance(in_p, Partial) - and out_p != in_p - ): - should_convert_partial = True - break - - input_target_placements = map_placements_after_broadcast( - tuple(out_placements), - common_shape, - input_arg_dims_map, - partial_to_replicate=should_convert_partial, - ) - - input_arg_target_spec = DTensorSpec( - mesh=followed_strategy.mesh, - placements=input_target_placements, - tensor_meta=input_arg_spec.tensor_meta, - ) - input_specs.append(input_arg_target_spec) - redistribute_costs.append( - generate_redistribute_costs(input_arg, input_arg_target_spec) - ) - - pointwise_strategy.strategies.append( - OpSpec( - output_specs=DTensorSpec( - mesh=followed_strategy.mesh, - placements=tuple(out_placements), - ), - input_specs=input_specs, - redistribute_cost=redistribute_costs, - ) - ) - return pointwise_strategy - - -p_sum_scalar_redistribute_ops = { - aten.add.Tensor, - aten.add_.Tensor, - aten.sub.Tensor, - aten.sub_.Tensor, -} - -norm_partial_avoidable_redistribute_ops = { - aten.div.Scalar, - aten.div_.Scalar, - aten.mul.Scalar, - aten.mul_.Scalar, -} - for op in pointwise_ops: _register_single_dim_pointwise(op) -# TODO: add all for_each ops -for_each_ops = [ - aten._foreach_abs.default, - aten._foreach_abs_.default, - aten._foreach_addcdiv_.Scalar, - aten._foreach_addcdiv_.ScalarList, - aten._foreach_addcdiv_.Tensor, - aten._foreach_addcmul.Scalar, - aten._foreach_addcmul_.Scalar, - aten._foreach_addcmul_.ScalarList, - aten._foreach_addcmul_.Tensor, - aten._foreach_clamp_max_.Scalar, - aten._foreach_clamp_min_.Scalar, - aten._foreach_div_.List, - aten._foreach_div_.Scalar, - aten._foreach_div_.ScalarList, - aten._foreach_div_.Tensor, - aten._foreach_div.List, - aten._foreach_div.Scalar, - aten._foreach_div.ScalarList, - aten._foreach_div.Tensor, - aten._foreach_lerp_.Scalar, - aten._foreach_maximum_.List, - aten._foreach_mul.Scalar, - aten._foreach_mul.ScalarList, - aten._foreach_mul.Tensor, - aten._foreach_mul.List, - aten._foreach_mul_.Scalar, - aten._foreach_mul_.ScalarList, - aten._foreach_mul_.Tensor, - aten._foreach_mul_.List, - aten._foreach_pow.List, - aten._foreach_pow.ScalarList, - aten._foreach_neg.default, - aten._foreach_neg_.default, - aten._foreach_reciprocal_.default, - aten._foreach_sub.Scalar, - aten._foreach_sub_.Scalar, - aten._foreach_sub.List, - aten._foreach_sub_.List, - aten._foreach_sub.ScalarList, - aten._foreach_sub_.ScalarList, - aten._foreach_sqrt.default, - aten._foreach_sqrt_.default, - aten._foreach_zero_.default, - aten._foreach_exp.default, - aten._foreach_exp_.default, - aten._foreach_cos.default, - aten._foreach_cos_.default, - aten._foreach_log.default, - aten._foreach_log_.default, - aten._amp_foreach_non_finite_check_and_unscale_.default, -] - -for_each_linearity_ops = [ - aten._foreach_add.Scalar, - aten._foreach_add_.Scalar, - aten._foreach_add_.ScalarList, - aten._foreach_add.List, - aten._foreach_add_.List, -] - - -def list_pointwise_strategy( - op_schema: OpSchema, linearity: bool = False -) -> StrategyType: - """ - Apply the pointwise strategy to the zipped arguments. For example, if we - run a foreach add of two lists l1 and l2, then we apply the pointwise - strategy on each pair (l1[i], l2[i]). If the first argument is a list but - the second (or later) one is a tensor, then we broadcast the tensor by - replicating it into a list with the length of the first argument. - - Args: - mesh (DeviceMesh): device mesh for pointwise ops - op_schema (OpSchema): schema of the operator to generate strategy for - linearity (bool): specify whether op(a) + op(b) = op(a + b) - - Returns: - OpStrategy: generated strategy - """ - - def args_tuple_strategies( - args_schema: tuple[object, ...], - ) -> list[TupleStrategy | None]: - first_arg = args_schema[0] - if not isinstance(first_arg, TupleStrategy): - raise AssertionError - strategy_len = len(first_arg.children) - tuple_strategies: list[TupleStrategy | None] = [] - for arg_idx, arg in enumerate(args_schema): - if isinstance(arg, TupleStrategy): - # every tuple strategy should have the same length - if len(arg.children) != strategy_len: - raise AssertionError - tuple_strategies.append(arg) - elif isinstance(arg, OpStrategy): - if arg_idx > 0: # implicitly broadcast - tuple_strategies.append( - TupleStrategy([arg for _ in range(strategy_len)]) - ) - else: - raise RuntimeError( - f"list op only supports tuple strategy! {op_schema}" - ) - else: - # insert None as placeholder so that the idx of arg is kept - tuple_strategies.append(None) - return tuple_strategies - - args_strategies = args_tuple_strategies(op_schema.args_schema) - follow_strategy: TupleStrategy = not_none(args_strategies[0]) - list_strategy: list[OpStrategy] = [] - - for child_idx, child_strtgy in enumerate(follow_strategy.children): - if not isinstance(child_strtgy, OpStrategy): - raise AssertionError - args_schema: list[OpStrategy | None] = [ - cast(OpStrategy, arg_strategy.children[child_idx]) if arg_strategy else None - for arg_strategy in args_strategies - ] - pointwise_strategy: OpStrategy = common_pointwise_strategy( - op_schema.op, - args_schema, - child_strtgy, - linearity, - scalar_tensor_idx=( - _FUSED_OP_SCALAR_IDX if op_schema.op in fused_ops else None - ), - ) - list_strategy.append(pointwise_strategy) - return TupleStrategy(list_strategy) - - -def list_linear_pointwise_strategy(op_schema: OpSchema) -> StrategyType: - """ - for each list op stratgy that supports linearity - """ - return list_pointwise_strategy(op_schema, linearity=True) - - -fused_ops = [ - aten._fused_adam_.default, - aten._fused_adam.default, - aten._fused_adam.tensor_lr, - aten._fused_adam_.tensor_lr, - aten._fused_adamw_.default, - aten._fused_adamw.default, - aten._fused_adamw.tensor_lr, - aten._fused_adamw_.tensor_lr, -] - def register_inductor_prims() -> None: """Register DTensor sharding strategies for inductor prims ops. From 0254024ede582e5991553ea608208e6505b2b15a Mon Sep 17 00:00:00 2001 From: Laith Sakka Date: Thu, 19 Mar 2026 11:15:31 -0700 Subject: [PATCH 0050/1172] reduce threshold for calling symp.factor to 50 (#177779) Pull Request resolved: https://github.com/pytorch/pytorch/pull/177779 Approved by: https://github.com/Skylion007 --- torch/fx/experimental/_size_hinting.py | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/torch/fx/experimental/_size_hinting.py b/torch/fx/experimental/_size_hinting.py index 80791e4dccf64..83811367b8d25 100644 --- a/torch/fx/experimental/_size_hinting.py +++ b/torch/fx/experimental/_size_hinting.py @@ -23,6 +23,11 @@ log = logging.getLogger(__name__) +# Maximum number of free symbols in an expression before we skip +# sympy.factor() in optimization_hint process for unbacked. +# Factoring polynomials with many variables is expensive. +SYMPY_FACTOR_MAX_FREE_SYMBOLS = 50 + if TYPE_CHECKING: from torch.fx.experimental.symbolic_shapes import ShapeEnv @@ -297,9 +302,7 @@ def _sub_unbacked_exprs(shape_env: ShapeEnv, expr: sympy.Expr) -> sympy.Expr: new_expr = expr.subs(replacements) if new_expr == expr: break - # Limit sympy.factor() to expressions with <= 200 free symbols, - # as factoring polynomials with many variables is expensive. - if len(new_expr.free_symbols) <= 200: + if len(new_expr.free_symbols) <= SYMPY_FACTOR_MAX_FREE_SYMBOLS: expr = sympy.factor(new_expr) else: expr = new_expr @@ -391,9 +394,10 @@ def _optimization_hint_base( if has_free_unbacked_symbols(expr): # Make sure to substitute with the factored version # e.g. 10*(s0 + u0) instead of 10*s0 + 10*u0 - # Limit sympy.factor() to expressions with <= 200 free symbols, - # as factoring polynomials with many variables is expensive. - if isinstance(original, sympy.Expr) and len(original.free_symbols) <= 200: + if ( + isinstance(original, sympy.Expr) + and len(original.free_symbols) <= SYMPY_FACTOR_MAX_FREE_SYMBOLS + ): original = sympy.factor(original) expr = _sub_unbacked_exprs(shape_env, original) From 6bc21e0397f6ea29517f2bca106a6df87c22c112 Mon Sep 17 00:00:00 2001 From: amdfaa <107946068+amdfaa@users.noreply.github.com> Date: Tue, 24 Mar 2026 00:43:32 +0000 Subject: [PATCH 0051/1172] [ROCm][CI] Update GPU_FLAG to remove network option (#176612) Removed '--network=host' from GPU_FLAG in action.yml. Pull Request resolved: https://github.com/pytorch/pytorch/pull/176612 Approved by: https://github.com/Skylion007, https://github.com/jeffdaily --- .github/actions/setup-rocm/action.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/actions/setup-rocm/action.yml b/.github/actions/setup-rocm/action.yml index 5bb982a4085b1..3c1445b85fa84 100644 --- a/.github/actions/setup-rocm/action.yml +++ b/.github/actions/setup-rocm/action.yml @@ -110,7 +110,7 @@ runs: # This is due to the device files (/dev/kfd & /dev/dri) being owned by video group on bare metal. # This video group ID maps to subgid 1 inside the docker image due to the /etc/subgid entries. # The group name corresponding to group ID 1 can change depending on the OS, so both are necessary. - echo "GPU_FLAG=--device=/dev/mem --device=/dev/kfd $DEVICE_FLAG --group-add video --group-add $render_gid --group-add daemon --group-add bin --cap-add=SYS_PTRACE --security-opt seccomp=unconfined --network=host" >> "${GITHUB_ENV}" + echo "GPU_FLAG=--device=/dev/mem --device=/dev/kfd $DEVICE_FLAG --group-add video --group-add $render_gid --group-add daemon --group-add bin --cap-add=SYS_PTRACE --security-opt seccomp=unconfined" >> "${GITHUB_ENV}" - name: Login to ECR uses: pytorch/pytorch/.github/actions/ecr-login@main From 5327b2623329e5cf1d1833cc9a8d152b632383d2 Mon Sep 17 00:00:00 2001 From: Vasiliy Kuznetsov Date: Tue, 24 Mar 2026 00:50:16 +0000 Subject: [PATCH 0052/1172] automate cherry pick handling for release note scripts (#175742) Summary: Adds two scripts to automate the release note generation process further: * `parse_cherry_picks.py` takes in a tracker issue (like https://github.com/pytorch/pytorch/issues/170119) and a commitlist.csv file and parses the right github hashes of trunk commits from the tracker issue, validates they match an entry in commitlist.csv, and writes the parsed cherry-pick hashes to a file * `remove_cherry_picks.py` takes the output of the script above and creates a new commit list with the cherry picks removed all unexpected behavior is logged for manual inspection Test Plan: ``` // (before this script - generate results/commitlist.csv) $ cd scripts/release_notes // parse cherry picks into 2.10, creates results/cherry_picks_170119.csv $ python parse_cherry_picks.py --commitlist results/commitlist.csv https://github.com/pytorch/pytorch/issues/170119 // create a new commitlist.csv with cherry picks removed, creates results/commitlist_no_cherry_picks.csv $ python remove_cherry_picks.py --commitlist results/commitlist.csv --cherry-picks results/cherry_picks_170119.csv // manually swap the commitlist files $ mv results/commitlist.csv results/commitlist_before_cherry_pick_removal.csv $ mv results/commitlist_no_cherry_picks.csv results/commitlist.csv // inspect dir contents $ ls results cherry_pick_removals.log cherry_picks_170119.csv commitlist.csv commitlist_before_cherry_pick_removal.csv // proceeed with further steps of release notes prep (not shown) ``` Pull Request resolved: https://github.com/pytorch/pytorch/pull/175742 Approved by: https://github.com/atalman --- scripts/release_notes/parse_cherry_picks.py | 319 +++++++++++++++++++ scripts/release_notes/remove_cherry_picks.py | 139 ++++++++ 2 files changed, 458 insertions(+) create mode 100644 scripts/release_notes/parse_cherry_picks.py create mode 100644 scripts/release_notes/remove_cherry_picks.py diff --git a/scripts/release_notes/parse_cherry_picks.py b/scripts/release_notes/parse_cherry_picks.py new file mode 100644 index 0000000000000..5494a4fdba203 --- /dev/null +++ b/scripts/release_notes/parse_cherry_picks.py @@ -0,0 +1,319 @@ +#!/usr/bin/env python3 +"""Parse cherry-pick comments from a PyTorch GitHub issue and extract trunk PR info. + +Usage: + python scripts/release_notes/parse_cherry_picks.py \\ + https://github.com/pytorch/pytorch/issues/170119 \\ + --commitlist scripts/release_notes/results/commitlist.csv + +Outputs a CSV file to scripts/release_notes/results/cherry_picks_.csv. +Columns: comment_id, pr_number, pr_title, commit_sha +Validates that each commit_sha matches an entry in the commitlist and logs +warnings for any mismatches. +""" + +import argparse +import csv +import json +import logging +import re +import subprocess +import sys +from pathlib import Path + + +logging.basicConfig(level=logging.INFO, format="%(levelname)s: %(message)s") +logger = logging.getLogger(__name__) + + +def parse_issue_url(url: str) -> tuple[str, str]: + """Extract owner/repo and issue number from a GitHub issue URL.""" + m = re.match(r"https://github\.com/([^/]+/[^/]+)/issues/(\d+)", url) # @lint-ignore + if not m: + raise ValueError(f"Invalid GitHub issue URL: {url}") + return m.group(1), m.group(2) + + +def gh_api(endpoint: str) -> list | dict: + """Call gh api with pagination and return parsed JSON.""" + result = subprocess.run( + ["gh", "api", endpoint, "--paginate"], + capture_output=True, + text=True, + check=True, + ) + # --paginate may return multiple JSON arrays concatenated; we need to handle that + # gh api --paginate with --jq is cleaner, but let's parse raw output + # When paginating, gh outputs one JSON array per page on separate "lines" + output = result.stdout.strip() + if not output: + return [] + + # Try parsing as a single JSON value first + try: + return json.loads(output) + except json.JSONDecodeError: + pass + + # If that fails, it's multiple JSON arrays concatenated + # Split on ][ boundaries and merge + all_items = [] + decoder = json.JSONDecoder() + pos = 0 + while pos < len(output): + # Skip whitespace + while pos < len(output) and output[pos] in " \t\n\r": + pos += 1 + if pos >= len(output): + break + obj, end_pos = decoder.raw_decode(output, pos) + if isinstance(obj, list): + all_items.extend(obj) + else: + all_items.append(obj) + pos = end_pos + return all_items + + +def fetch_comments(repo: str, issue_number: str) -> list[dict]: + """Fetch all comments from a GitHub issue.""" + endpoint = f"repos/{repo}/issues/{issue_number}/comments?per_page=100" + comments = gh_api(endpoint) + logger.info(f"Fetched {len(comments)} comments from {repo}#{issue_number}") + return comments + + +def fetch_pr_title(repo: str, pr_number: str) -> str: + """Fetch the title of a PR.""" + try: + result = subprocess.run( + ["gh", "api", f"repos/{repo}/pulls/{pr_number}", "--jq", ".title"], + capture_output=True, + text=True, + check=True, + ) + return result.stdout.strip() + except subprocess.CalledProcessError as e: + logger.warning(f"Failed to fetch title for PR #{pr_number}: {e}") + return "" + + +def fetch_landed_commit(repo: str, pr_number: str) -> str: + """Fetch the actual landed commit SHA for a PR. + + PyTorch uses a merge bot that squash-merges outside of GitHub's standard + merge mechanism, so the PR API's merge_commit_sha is unreliable. + + Strategy: + 1. Issue events API: look for the 'closed' event with a commit_id + 2. Commit search API: search for commits containing '(#NNNNN)' in message + """ + # Strategy 1: Issue events API + try: + result = subprocess.run( + [ + "gh", + "api", + f"repos/{repo}/issues/{pr_number}/events?per_page=100", + "--paginate", + "--jq", + '.[] | select(.event == "closed" and .commit_id) | .commit_id', + ], + capture_output=True, + text=True, + check=True, + ) + commit_ids = result.stdout.strip().splitlines() + if commit_ids and commit_ids[0]: + return commit_ids[0] + except subprocess.CalledProcessError: + pass + + # Strategy 2: Commit search API + logger.info(f" Events API had no commit for PR #{pr_number}, trying search API...") + try: + result = subprocess.run( + [ + "gh", + "api", + f'search/commits?q=repo:{repo}+"(#{pr_number})"', + "--jq", + ".items[0].sha", + ], + capture_output=True, + text=True, + check=True, + ) + sha = result.stdout.strip() + if sha and sha != "null": + return sha + except subprocess.CalledProcessError: + pass + + logger.warning(f"Could not find landed commit for PR #{pr_number}") + return "" + + +def extract_trunk_prs(comment_body: str) -> list[dict]: + """Extract trunk PR numbers/commit hashes from a comment body. + + Returns a list of dicts with keys: pr_number (str or ""), raw_commit (str or ""). + Returns an empty list if the comment doesn't contain the trunk PR section. + """ + # Find the "Link to landed trunk PR" section + # Handle variations: "Link to landed trunk PR", "Link to the landed trunk PR" + trunk_pattern = re.compile( + r"Link to (?:the )?landed trunk PR[^:]*:\s*\n(.*?)(?=Li[nn][kt] to (?:the )?release branch PR|Criteria Category:|$)", + re.DOTALL | re.IGNORECASE, + ) + match = trunk_pattern.search(comment_body) + if not match: + return [] + + trunk_section = match.group(1).strip() + + # Check for NA/N/A + if re.match(r"^\*?\s*-?\s*N/?A\s*$", trunk_section, re.IGNORECASE): + return [{"pr_number": "", "raw_commit": ""}] + + results = [] + + # Extract PR URLs + pr_urls = re.findall( + r"https://github\.com/pytorch/pytorch/pull/(\d+)", # @lint-ignore + trunk_section, + ) + for pr_num in pr_urls: + results.append({"pr_number": pr_num, "raw_commit": ""}) + + # Extract raw commit hashes (40-char hex) that aren't part of a URL + # Remove URLs first to avoid matching hashes embedded in URLs + section_no_urls = re.sub(r"https://\S+", "", trunk_section) + raw_commits = re.findall(r"\b([0-9a-f]{40})\b", section_no_urls) + for commit_hash in raw_commits: + results.append({"pr_number": "", "raw_commit": commit_hash}) + + # If we found the section but nothing matched, treat as empty/NA + if not results: + return [{"pr_number": "", "raw_commit": ""}] + + return results + + +def main(): + parser = argparse.ArgumentParser( + description="Parse cherry-pick comments from a PyTorch GitHub issue" + ) + parser.add_argument("issue_url", help="GitHub issue URL to parse") + parser.add_argument( + "-o", + "--output", + default=None, + help="Output CSV file path (default: results/cherry_picks_.csv relative to script)", + ) + parser.add_argument( + "--commitlist", + required=True, + help="Path to commitlist.csv to validate commit hashes against", + ) + args = parser.parse_args() + + repo, issue_number = parse_issue_url(args.issue_url) + if args.output: + output_path = args.output + else: + script_dir = Path(__file__).resolve().parent + results_dir = script_dir / "results" + results_dir.mkdir(exist_ok=True) + output_path = str(results_dir / f"cherry_picks_{issue_number}.csv") + + comments = fetch_comments(repo, issue_number) + + rows = [] + for comment in comments: + comment_id = comment["id"] + body = comment["body"] or "" + + trunk_prs = extract_trunk_prs(body) + if not trunk_prs: + # Comment doesn't contain the trunk PR section — skip + continue + + if len(trunk_prs) > 1: + logger.warning( + f"Comment {comment_id} has {len(trunk_prs)} trunk PRs: " + + ", ".join( + p["pr_number"] or p["raw_commit"] or "N/A" for p in trunk_prs + ) + ) + + for pr_info in trunk_prs: + pr_number = pr_info["pr_number"] + raw_commit = pr_info["raw_commit"] + pr_title = "" + commit_sha = raw_commit # Use raw commit if provided + + if pr_number: + logger.info(f"Fetching info for PR #{pr_number}...") + pr_title = fetch_pr_title(repo, pr_number) + commit_sha = fetch_landed_commit(repo, pr_number) + + rows.append( + { + "comment_id": comment_id, + "pr_number": pr_number, + "pr_title": pr_title, + "commit_sha": commit_sha, + } + ) + + # Validate against commitlist + commitlist_hashes = set() + with open(args.commitlist, newline="") as f: + reader = csv.reader(f) + next(reader) # skip header + for row in reader: + if row: + commitlist_hashes.add(row[0]) + + logger.info(f"Loaded {len(commitlist_hashes)} hashes from {args.commitlist}") + + matched = 0 + mismatched = 0 + skipped = 0 + for row in rows: + sha = row["commit_sha"] + if not sha: + skipped += 1 + continue + # commitlist uses abbreviated hashes; check if any is a prefix of + # our full hash, or vice versa + if any( + sha.startswith(cl_hash) or cl_hash.startswith(sha) + for cl_hash in commitlist_hashes + ): + matched += 1 + else: + mismatched += 1 + logger.warning( + f"Commit {sha[:11]} (PR #{row['pr_number'] or 'N/A'}) " + f"not found in commitlist" + ) + + logger.info( + f"Commitlist validation: {matched} matched, {mismatched} not found, " + f"{skipped} skipped (no hash)" + ) + + # Write CSV + fieldnames = ["comment_id", "pr_number", "pr_title", "commit_sha"] + with open(output_path, "w", newline="") as f: + writer = csv.DictWriter(f, fieldnames=fieldnames) + writer.writeheader() + writer.writerows(rows) + + logger.info(f"Wrote {len(rows)} rows to {output_path}") + + +if __name__ == "__main__": + main() diff --git a/scripts/release_notes/remove_cherry_picks.py b/scripts/release_notes/remove_cherry_picks.py new file mode 100644 index 0000000000000..b7ef72c11d594 --- /dev/null +++ b/scripts/release_notes/remove_cherry_picks.py @@ -0,0 +1,139 @@ +#!/usr/bin/env python3 +"""Remove cherry-picked commits from a commitlist CSV. + +Usage: + python scripts/release_notes/remove_cherry_picks.py \\ + --commitlist scripts/release_notes/results/commitlist.csv \\ + --cherry-picks scripts/release_notes/results/cherry_picks_170119.csv + +Reads the cherry picks CSV (output of parse_cherry_picks.py) and removes +matching rows from the commitlist, writing the result to a new file. +The original commitlist is never modified. +""" + +import argparse +import csv +import logging +from pathlib import Path + + +logging.basicConfig(level=logging.INFO, format="%(levelname)s: %(message)s") +logger = logging.getLogger(__name__) + + +def main(): + parser = argparse.ArgumentParser( + description="Remove cherry-picked commits from a commitlist CSV" + ) + parser.add_argument( + "--commitlist", + required=True, + help="Path to the input commitlist.csv", + ) + parser.add_argument( + "--cherry-picks", + required=True, + help="Path to the cherry picks CSV (output of parse_cherry_picks.py)", + ) + parser.add_argument( + "-o", + "--output", + default=None, + help="Output commitlist path (default: commitlist_no_cherry_picks.csv in same dir as commitlist)", + ) + parser.add_argument( + "--log", + default=None, + help="Path for the removal log file (default: results/cherry_pick_removals.log relative to script)", + ) + args = parser.parse_args() + + commitlist_path = Path(args.commitlist) + if args.output: + output_path = Path(args.output) + else: + output_path = commitlist_path.parent / "commitlist_no_cherry_picks.csv" + + script_dir = Path(__file__).resolve().parent + if args.log: + log_path = Path(args.log) + else: + results_dir = script_dir / "results" + results_dir.mkdir(exist_ok=True) + log_path = results_dir / "cherry_pick_removals.log" + + # Load cherry picks + cherry_picks = [] + with open(args.cherry_picks, newline="") as f: + reader = csv.DictReader(f) + for row in reader: + cherry_picks.append(row) + logger.info(f"Loaded {len(cherry_picks)} cherry pick entries") + + # Read commitlist as raw lines to preserve exact formatting + with open(commitlist_path) as f: + header_line = f.readline() + commitlist_lines = f.readlines() + logger.info(f"Loaded {len(commitlist_lines)} rows from {commitlist_path}") + + # Build a map from abbreviated hash to line for fast lookup + hash_to_line = {} + for line in commitlist_lines: + abbrev_hash = line.split(",", 1)[0] + hash_to_line[abbrev_hash] = line + + # Process cherry picks: find and remove matching rows + removed = 0 + not_found = 0 + skipped = 0 + log_entries = [] + + for cp in cherry_picks: + sha = cp.get("commit_sha", "") + pr_number = cp.get("pr_number", "") + pr_title = cp.get("pr_title", "") + label = f"PR #{pr_number} ({pr_title})" if pr_number else "N/A" + + if not sha: + skipped += 1 + log_entries.append(f"SKIPPED (no hash): {label}") + continue + + # Find matching abbreviated hash via prefix match + matched_hash = None + for abbrev_hash in hash_to_line: + if sha.startswith(abbrev_hash) or abbrev_hash.startswith(sha): + matched_hash = abbrev_hash + break + + if matched_hash: + removed += 1 + log_entries.append(f"REMOVED: {matched_hash} -> {label}") + del hash_to_line[matched_hash] + else: + not_found += 1 + log_entries.append(f"NOT FOUND: {sha[:11]} -> {label}") + logger.warning(f"Commit {sha[:11]} ({label}) not found in commitlist") + + # Write new commitlist + with open(output_path, "w") as f: + f.write(header_line) + for line in commitlist_lines: + abbrev_hash = line.split(",", 1)[0] + if abbrev_hash in hash_to_line: + f.write(line) + + logger.info( + f"Wrote {len(hash_to_line)} rows to {output_path} " + f"(removed {removed}, not found {not_found}, skipped {skipped})" + ) + + # Write log file + with open(log_path, "w") as f: + for entry in log_entries: + f.write(entry + "\n") + logger.info(f"Wrote removal log to {log_path}") + + +if __name__ == "__main__": + main() From 7f1e2a26497d550c40ba6e856bf9a4152cb30d90 Mon Sep 17 00:00:00 2001 From: Vasiliy Kuznetsov Date: Tue, 24 Mar 2026 00:54:40 +0000 Subject: [PATCH 0053/1172] release notes categorize: make torchtext dependency optional (#178144) Summary: Moves torchtext from a required dependency to optional run-time dependency for the part of release note generation which categorizes commits. torchtext is archived and having it as a dependency makes it hard to use recent PyTorch versions. For now, make the dependency optional. In the future we can probably just remove the classifier (only 5 commits need classification for 2.11, faster to just manually do it) or replace it with a more recent general model. Test Plan: ```bash python categorize.py ``` Pull Request resolved: https://github.com/pytorch/pytorch/pull/178144 Approved by: https://github.com/atalman --- scripts/release_notes/categorize.py | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/scripts/release_notes/categorize.py b/scripts/release_notes/categorize.py index 10ee551b74f1a..aa27e0b50577f 100644 --- a/scripts/release_notes/categorize.py +++ b/scripts/release_notes/categorize.py @@ -4,27 +4,25 @@ from pathlib import Path import common - -# Imports for working with classi -from classifier import ( - CategoryConfig, - CommitClassifier, - CommitClassifierInputs, - get_author_map, - get_file_map, - XLMR_BASE, -) from commitlist import CommitList from common import get_commit_data_cache, topics -import torch - class Categorizer: def __init__(self, path, category="Uncategorized", use_classifier: bool = False): self.cache = get_commit_data_cache() self.commits = CommitList.from_existing(path) if use_classifier: + from classifier import ( + CategoryConfig, + CommitClassifier, + get_author_map, + get_file_map, + XLMR_BASE, + ) + + import torch + print("Using a classifier to aid with categorization.") device = "cuda" if torch.cuda.is_available() else "cpu" classifier_config = CategoryConfig(common.categories) @@ -108,6 +106,8 @@ def handle_commit(self, commit, i, total, commits): features = self.features(commit) if self.classifier is not None: + from classifier import CommitClassifierInputs + # Some commits don't have authors: author = features.author if features.author else "Unknown" files = " ".join(features.files_changed) From b1311297ee2770e244d002ed8d58a2671a4e5bc9 Mon Sep 17 00:00:00 2001 From: Eddie Yan Date: Tue, 24 Mar 2026 01:03:07 +0000 Subject: [PATCH 0054/1172] [TEST][cuDNN][TF32] Bump tolerances for `test_Conv2d_naive_groups` (#178161) New TF32 kernel selection on B200 requires looser tolerances Pull Request resolved: https://github.com/pytorch/pytorch/pull/178161 Approved by: https://github.com/Skylion007, https://github.com/mlazos --- test/nn/test_convolution.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/nn/test_convolution.py b/test/nn/test_convolution.py index 786ac2cd74933..bb393f63fc38f 100644 --- a/test/nn/test_convolution.py +++ b/test/nn/test_convolution.py @@ -3581,7 +3581,7 @@ def test_ConvTranspose3d_size_1_kernel(self, device): @dtypes(torch.float) @torch.backends.cudnn.flags(enabled=True, deterministic=True, benchmark=False) @torch.backends.miopen.flags(immediate=True) - @tf32_on_and_off(0.001) + @tf32_on_and_off(0.005) def test_Conv2d_naive_groups(self, device, dtype): # Check that grouped convolutions matches two half convolutions m = nn.Conv2d(4, 4, kernel_size=3, groups=2).to(device, dtype) From ea0a353f21f954897e890e5a4a583472eccea80f Mon Sep 17 00:00:00 2001 From: Ruth C Date: Tue, 24 Mar 2026 02:22:39 +0000 Subject: [PATCH 0055/1172] [inductor] Fix RNG order change caused by reorder_for_locality (#176842) Fixes #175156 ## Problem The `reorder_for_locality` pass in TorchInductor may reorder RNG operations such as `aten.randint`. Since RNG ops consume global RNG state, their execution order affects program semantics. Reordering them can therefore produce results that differ from eager execution. As reported in the issue, two consecutive `torch.randint` calls may produce different results between eager and compiled execution even when the RNG seed is reset. ## Root cause `reorder_for_locality` may move producer nodes using: node.prepend(other_node) RNG operations are currently treated as pure ops and may therefore be reordered, even though they consume RNG state. ## Fix Prevent RNG ops from being reordered during the `reorder_for_locality` pass. ## Tests Added a regression test: test/inductor/test_rng_reordering.py which verifies that multiple `torch.randint` calls produce consistent results between eager and compiled execution. ## Notes This patch focuses on fixing the reported `aten.randint` case. The same approach could potentially be generalized to other RNG ops if necessary. Pull Request resolved: https://github.com/pytorch/pytorch/pull/176842 Approved by: https://github.com/Lucaskabela, https://github.com/karthickai, https://github.com/jansel Co-authored-by: PyTorch MergeBot --- test/inductor/test_deterministic.py | 19 +++++++++++++++++++ torch/_inductor/fx_passes/post_grad.py | 12 ++++++++++++ 2 files changed, 31 insertions(+) diff --git a/test/inductor/test_deterministic.py b/test/inductor/test_deterministic.py index 4a5b5a2829f50..be14ab8a6697f 100644 --- a/test/inductor/test_deterministic.py +++ b/test/inductor/test_deterministic.py @@ -113,6 +113,25 @@ def foo(x): else: self.assertTrue(counters["inductor"]["coordesc_tuning_bench"] > 0) + def test_reorder_for_locality_preserves_randint_order(self): + with inductor_config.patch(fallback_random=True): + + def fn(): + torch.manual_seed(0) + out = torch.randint(0, 100, (4, 1), dtype=torch.int64) + _ = torch.randint(0, 100, (2, 1), dtype=torch.int64) + return out + + compiled = torch.compile(fn, backend="inductor") + + torch.manual_seed(0) + eager = fn() + + torch.manual_seed(0) + compiled_out = compiled() + + torch.testing.assert_close(eager, compiled_out) + @unittest.skipIf(IS_FBCODE, "Skipping run2run determinism test in fbcode") @parametrize("model_name", ["GoogleFnet", "BertForMaskedLM", "DistillGPT2"]) @parametrize("training_or_inference", ["training", "inference"]) diff --git a/torch/_inductor/fx_passes/post_grad.py b/torch/_inductor/fx_passes/post_grad.py index bbefebf92a237..02ec541caf7a9 100644 --- a/torch/_inductor/fx_passes/post_grad.py +++ b/torch/_inductor/fx_passes/post_grad.py @@ -786,6 +786,13 @@ def check(): def check(): return True + def consumes_rng_state(node: torch.fx.Node) -> bool: + return ( + node.op == "call_function" + and isinstance(node.target, torch._ops.OpOverload) + and torch.Tag.nondeterministic_seeded in node.target.tags + ) + def visit(other_node): if ( other_node.op == "call_function" @@ -795,6 +802,11 @@ def visit(other_node): == get_mutation_region_id(graph, other_node) and check() ): + # Ops that consume RNG state are order-sensitive and must not be + # reordered during locality optimization. + if consumes_rng_state(other_node): + return + # move node's producers right before it node.prepend(other_node) From 88c77db9c862573f9d7a8eda58ae735415bc740d Mon Sep 17 00:00:00 2001 From: "Yu, Guangye" Date: Mon, 23 Mar 2026 18:20:37 +0000 Subject: [PATCH 0056/1172] Introduce a unified API to empty the host cache memory (#171270) # Motivation This PR introduces a unified API, `torch.accelerator.empty_host_cache`, for releasing unused host (pinned) memory. This API can also be used prior to graph capture. For example, as discussed in https://github.com/pytorch/pytorch/pull/167507, CUDA supports graph capture and replay involving host memory as well. Therefore, it is important to clean up both the device and host memory before starting graph capture to ensure correct and predictable behavior. Pull Request resolved: https://github.com/pytorch/pytorch/pull/171270 Approved by: https://github.com/eellison --- aten/src/ATen/DeviceAccelerator.cpp | 6 ++++++ aten/src/ATen/DeviceAccelerator.h | 8 ++++++++ build_variables.bzl | 1 + docs/source/accelerator.md | 1 + torch/_C/__init__.pyi.in | 1 + torch/accelerator/__init__.py | 2 ++ torch/accelerator/graphs.py | 1 + torch/accelerator/memory.py | 11 +++++++++++ torch/csrc/DeviceAccelerator.cpp | 9 +++++++++ 9 files changed, 40 insertions(+) diff --git a/aten/src/ATen/DeviceAccelerator.cpp b/aten/src/ATen/DeviceAccelerator.cpp index efab9ec9c5927..fed2e6c4febe8 100644 --- a/aten/src/ATen/DeviceAccelerator.cpp +++ b/aten/src/ATen/DeviceAccelerator.cpp @@ -1,4 +1,5 @@ #include +#include #include #include @@ -136,6 +137,11 @@ c10::DeviceCapability getDeviceCapability(c10::DeviceIndex device_index) { c10::impl::VirtualGuardImpl impl(device_type); return impl.getDeviceCapability({device_type, device_index}); } + +void emptyHostCache() { + const auto device_type = getAccelerator(true).value(); + at::getHostAllocator(device_type)->empty_cache(); +} // NOLINTEND(bugprone-unchecked-optional-access) } // namespace at::accelerator diff --git a/aten/src/ATen/DeviceAccelerator.h b/aten/src/ATen/DeviceAccelerator.h index 81678ba6efc29..90fad38a5faca 100644 --- a/aten/src/ATen/DeviceAccelerator.h +++ b/aten/src/ATen/DeviceAccelerator.h @@ -78,11 +78,19 @@ TORCH_API c10::DeviceIndex maybeExchangeDevice(c10::DeviceIndex device_index); TORCH_API c10::DeviceCapability getDeviceCapability( c10::DeviceIndex device_index); +// Releases all unused device memory currently held by the accelerator's +// device-side caching allocator. The freed memory becomes available for reuse +// by other applications or processes. TORCH_API inline void emptyCache() { const auto device_type = getAccelerator(true).value(); at::getDeviceAllocator(device_type)->emptyCache(); } +// Releases all unused host (pinned) memory currently held by the accelerator's +// host-side caching allocator. The freed memory becomes available for reuse by +// other applications or processes. +TORCH_API void emptyHostCache(); + TORCH_API inline at::CachingDeviceAllocator::DeviceStats getDeviceStats( c10::DeviceIndex device_index) { const auto device_type = getAccelerator(true).value(); diff --git a/build_variables.bzl b/build_variables.bzl index fab555b9c1714..878132c592060 100644 --- a/build_variables.bzl +++ b/build_variables.bzl @@ -1159,6 +1159,7 @@ aten_cpu_source_non_codegen_list = [ "aten/src/ATen/LegacyVmapMode.cpp", "aten/src/ATen/LegacyVmapTransforms.cpp", "aten/src/ATen/core/BackendSelectFallbackKernel.cpp", + "aten/src/ATen/core/CachingHostAllocator.cpp", "aten/src/ATen/core/DeprecatedTypeProperties.cpp", "aten/src/ATen/core/DeprecatedTypePropertiesRegistry.cpp", "aten/src/ATen/core/Dict.cpp", diff --git a/docs/source/accelerator.md b/docs/source/accelerator.md index 53434cdfbfe25..ffa4fcce82555 100644 --- a/docs/source/accelerator.md +++ b/docs/source/accelerator.md @@ -50,6 +50,7 @@ :nosignatures: empty_cache + empty_host_cache get_memory_info max_memory_allocated max_memory_reserved diff --git a/torch/_C/__init__.pyi.in b/torch/_C/__init__.pyi.in index 0bce936f884d3..e81a9301a89f7 100644 --- a/torch/_C/__init__.pyi.in +++ b/torch/_C/__init__.pyi.in @@ -2606,6 +2606,7 @@ def _accelerator_exchangeDevice(device_index: _int) -> _int: ... def _accelerator_maybeExchangeDevice(device_index: _int) -> _int: ... def _accelerator_isAllocatorInitialized() -> _bool: ... def _accelerator_emptyCache() -> None: ... +def _accelerator_emptyHostCache() -> None: ... def _accelerator_getDeviceStats(device_index: _int) -> dict[str, Any]: ... def _accelerator_resetAccumulatedStats(device_index: _int) -> None: ... def _accelerator_resetPeakStats(device_index: _int) -> None: ... diff --git a/torch/accelerator/__init__.py b/torch/accelerator/__init__.py index 0a8be1b0749df..dc59efe33503a 100644 --- a/torch/accelerator/__init__.py +++ b/torch/accelerator/__init__.py @@ -12,6 +12,7 @@ from .graphs import Graph from .memory import ( empty_cache, + empty_host_cache, get_memory_info, max_memory_allocated, max_memory_reserved, @@ -33,6 +34,7 @@ "device_count", "device_index", "empty_cache", + "empty_host_cache", "get_memory_info", "is_available", "max_memory_allocated", diff --git a/torch/accelerator/graphs.py b/torch/accelerator/graphs.py index fc347a24cfe1d..8c86881c3722e 100644 --- a/torch/accelerator/graphs.py +++ b/torch/accelerator/graphs.py @@ -156,6 +156,7 @@ def __enter__(self) -> None: # very expensive, especially when performing multiple graph captures in sequence. gc.collect() torch.accelerator.empty_cache() + torch.accelerator.empty_host_cache() self.capture_begin() def __exit__(self, *exc_info: object) -> None: diff --git a/torch/accelerator/memory.py b/torch/accelerator/memory.py index a326af6dbcc1b..08998d221a58b 100644 --- a/torch/accelerator/memory.py +++ b/torch/accelerator/memory.py @@ -8,6 +8,7 @@ __all__ = [ "empty_cache", + "empty_host_cache", "get_memory_info", "max_memory_allocated", "max_memory_reserved", @@ -31,6 +32,16 @@ def empty_cache() -> None: torch._C._accelerator_emptyCache() +def empty_host_cache() -> None: + r"""Release all unoccupied cached host (pinned) memory currently held by the host caching + allocator so that it can be used by other applications. + + .. note:: This function is a no-op if the memory allocator for the current + :ref:`accelerator ` has not been initialized. + """ + torch._C._accelerator_emptyHostCache() + + def memory_stats(device_index: _device_t = None, /) -> OrderedDict[str, Any]: r"""Return a dictionary of accelerator device memory allocator statistics for a given device index. diff --git a/torch/csrc/DeviceAccelerator.cpp b/torch/csrc/DeviceAccelerator.cpp index 9e5aa1e5eaa69..c3caf43e805b4 100644 --- a/torch/csrc/DeviceAccelerator.cpp +++ b/torch/csrc/DeviceAccelerator.cpp @@ -101,6 +101,15 @@ void initModule(PyObject* module) { m.def("_accelerator_emptyCache", []() { at::accelerator::emptyCache(); }); + m.def("_accelerator_emptyHostCache", []() { + const auto device_type = at::accelerator::getAccelerator(true).value(); + if (torch::utils::is_device_lazy_init_supported(device_type) && + !torch::utils::is_device_initialized(device_type)) { + return; + } + at::accelerator::emptyHostCache(); + }); + m.def("_accelerator_getDeviceStats", [](c10::DeviceIndex device_index) { using c10::CachingAllocator::Stat; using c10::CachingAllocator::StatArray; From 65a02bb78997f509fd49ea2ff1bbcc13b5e7be80 Mon Sep 17 00:00:00 2001 From: Animesh Jain Date: Mon, 23 Mar 2026 15:18:15 -0700 Subject: [PATCH 0057/1172] [dynamo] Adjust more tests to prepare for inline_inbuilt_nn_module deprecation (#178067) Pull Request resolved: https://github.com/pytorch/pytorch/pull/178067 Approved by: https://github.com/choijon5, https://github.com/aditvenk, https://github.com/HPZ10, https://github.com/mlazos --- test/dynamo/test_activation_checkpointing.py | 1 - test/dynamo/test_decorators.py | 3 +- test/dynamo/test_dynamic_shapes.py | 8 --- test/dynamo/test_higher_order_ops.py | 59 +++-------------- test/dynamo/test_misc.py | 6 +- test/dynamo/test_modules.py | 20 ++---- test/dynamo/test_repros.py | 6 +- test/dynamo/test_structured_trace.py | 66 ++------------------ test/dynamo/test_utils.py | 22 +++---- test/functorch/test_control_flow.py | 9 ++- test/inductor/test_cudagraph_trees.py | 18 +----- 11 files changed, 38 insertions(+), 180 deletions(-) diff --git a/test/dynamo/test_activation_checkpointing.py b/test/dynamo/test_activation_checkpointing.py index 7c7717441217c..23fa43efb4864 100644 --- a/test/dynamo/test_activation_checkpointing.py +++ b/test/dynamo/test_activation_checkpointing.py @@ -1871,7 +1871,6 @@ def forward(self, x): @requires_distributed() @requires_cuda_and_triton def test_dynamo_does_not_trace_getattr_as_top_frame(self): - # inline_inbuilt_nn_modules is a proxy to emulate what FSDP tests do. from torch.distributed.algorithms._checkpoint.checkpoint_wrapper import ( CheckpointWrapper, ) diff --git a/test/dynamo/test_decorators.py b/test/dynamo/test_decorators.py index 96d25d9022ebc..273b915c46a57 100644 --- a/test/dynamo/test_decorators.py +++ b/test/dynamo/test_decorators.py @@ -1323,8 +1323,7 @@ def forward(self, a, *args): def _test_mark_static_address(self, guarded): # This test verifies that dynamo properly marks inputs as static # when using the mark_static_address API. - # For both inline_inbuilt_nn_modules True and False, we expect the - # tensor to be present in the buffers attribute of the graph. + # We expect the tensor to be present in the buffers attribute of the graph. compiles_with_buffers = 0 compiles = 0 diff --git a/test/dynamo/test_dynamic_shapes.py b/test/dynamo/test_dynamic_shapes.py index 3bebcfc345b70..a585c65c82157 100644 --- a/test/dynamo/test_dynamic_shapes.py +++ b/test/dynamo/test_dynamic_shapes.py @@ -1,5 +1,4 @@ # Owner(s): ["module: dynamo"] -import unittest import warnings from torch._dynamo import config @@ -85,13 +84,6 @@ def make_dynamic_cls(cls): make_dynamic_cls(test) del test -if TEST_Z3: - if not config.inline_inbuilt_nn_modules: - # TODO model is somehow not being freed when z3 is available - unittest.expectedFailure( - DynamicShapesMiscTests.test_parameter_free_dynamic_shapes # noqa: F821 - ) - # Test takes too long ~700s as of 414a1fd29f04d06e41b7f895368dd1f83a4be29d DynamicShapesExportTests.test_retracibility_dynamic_shapes = slowTest( # noqa: F821 DynamicShapesExportTests.test_retracibility_dynamic_shapes # noqa: F821 diff --git a/test/dynamo/test_higher_order_ops.py b/test/dynamo/test_higher_order_ops.py index e3ca2aae54a2a..0c6089d1eee80 100644 --- a/test/dynamo/test_higher_order_ops.py +++ b/test/dynamo/test_higher_order_ops.py @@ -2583,10 +2583,6 @@ def f(x): # 3 args - 1 for input, and other 2 for the weight and bias self.assertTrue(len(wrap_node.args), 3) - # Check that the linear bias and weight are getattr in the outer graph - if not torch._dynamo.config.inline_inbuilt_nn_modules: - self.assertTrue(len(dict(backend.graphs[0].named_parameters())) == 2) - # Check that the inner function has one op and its a linear op body_function = getattr(backend.graphs[0], wrap_node.args[0].name) self.assertEqual(op_count(body_function), 1) @@ -2715,10 +2711,6 @@ def f(x): wrap_node = find_first_node(backend.graphs[0], wrap) self.assertTrue(len(wrap_node.args), 3) - # Check that the linear bias and weight are getattr in the outer graph - if not torch._dynamo.config.inline_inbuilt_nn_modules: - self.assertTrue(len(dict(backend.graphs[0].named_parameters())) == 2) - # Check that the inner function has one op and its a linear op body_function = getattr(backend.graphs[0], wrap_node.args[0].name) self.assertEqual(op_count(body_function), 1) @@ -4504,10 +4496,9 @@ def wrapper_fn(model, params, inputs, targets): return actual = normalize_gm(wrapped_gm.print_readable(print_output=False)) - if torch._dynamo.config.inline_inbuilt_nn_modules: - self.assertExpectedInline( - actual, - """\ + self.assertExpectedInline( + actual, + """\ class GraphModule(torch.nn.Module): def forward(self, L_model_parameters_weight_: "f32[3, 3]", L_model_parameters_bias_: "f32[3]", L_inputs_: "f32[64, 3]", L_targets_: "f32[64, 3]"): l_model_parameters_weight_ = L_model_parameters_weight_ @@ -4520,22 +4511,7 @@ def forward(self, L_model_parameters_weight_: "f32[3, 3]", L_model_parameters_bi mse_loss: "f32[]" = torch.nn.functional.mse_loss(prediction, l_targets_); prediction = l_targets_ = None return (mse_loss,) """, - ) - else: - self.assertExpectedInline( - actual, - """\ -class GraphModule(torch.nn.Module): - def forward(self, L_inputs_: "f32[64, 3]", L_targets_: "f32[64, 3]"): - l_inputs_ = L_inputs_ - l_targets_ = L_targets_ - - prediction: "f32[64, 3]" = self.model(l_inputs_); l_inputs_ = None - - mse_loss: "f32[]" = torch.nn.functional.mse_loss(prediction, l_targets_); prediction = l_targets_ = None - return (mse_loss,) -""", - ) + ) def test_functional_call_sequential_params_and_buffers(self): # copied from test/test_stateless.py @@ -4566,8 +4542,7 @@ def wrapper_fn(model, params, buffers, inputs): return actual = normalize_gm(wrapped_gm.print_readable(print_output=False)) - if torch._dynamo.config.inline_inbuilt_nn_modules: - expected = """\ + expected = """\ class GraphModule(torch.nn.Module): def forward(self, L_inputs_: "f32[1, 1]", L_model_modules_l1_parameters_weight_: "f32[1, 1]", L_model_modules_l1_parameters_bias_: "f32[1]", L_model_buffers_buffer_: "f32[1]"): l_inputs_ = L_inputs_ @@ -4578,25 +4553,11 @@ def forward(self, L_inputs_: "f32[1, 1]", L_model_modules_l1_parameters_weight_: add: "f32[1, 1]" = linear + l_model_buffers_buffer_; linear = l_model_buffers_buffer_ = None return (add,) """ - # We found Windows/Linux have some empty line difference, empty_line_normalizer will help fix it. - self.assertExpectedInline( - empty_line_normalizer(actual), - empty_line_normalizer(normalize_gm(expected)), - ) - else: - self.assertExpectedInline( - actual, - """\ -class GraphModule(torch.nn.Module): - def forward(self, L_x_: "f32[1, 1]"): - l_x_ = L_x_ - - l__self___l1: "f32[1, 1]" = self.L__self___l1(l_x_); l_x_ = None - l__self___buffer: "f32[1]" = self.L__self___buffer - add: "f32[1, 1]" = l__self___l1 + l__self___buffer; l__self___l1 = l__self___buffer = None - return (add,) -""", - ) + # We found Windows/Linux have some empty line difference, empty_line_normalizer will help fix it. + self.assertExpectedInline( + empty_line_normalizer(actual), + empty_line_normalizer(normalize_gm(expected)), + ) def test_grad(self): counters.clear() diff --git a/test/dynamo/test_misc.py b/test/dynamo/test_misc.py index 94532e007f8d2..a70378e327ca0 100644 --- a/test/dynamo/test_misc.py +++ b/test/dynamo/test_misc.py @@ -7248,11 +7248,7 @@ def body(x): mod = Module() - error_message = "" - if torch._dynamo.config.inline_inbuilt_nn_modules: - error_message = r"Higher Order Operator: torch\.ops\.higher_order\.map_impl" - else: - error_message = "Can't inplace modify module params/buffers" + error_message = r"Higher Order Operator: torch\.ops\.higher_order\.map_impl" with self.assertRaisesRegex( torch._dynamo.exc.UncapturedHigherOrderOpError, error_message diff --git a/test/dynamo/test_modules.py b/test/dynamo/test_modules.py index 054ac7dcb7c8e..28ec48a82b45d 100644 --- a/test/dynamo/test_modules.py +++ b/test/dynamo/test_modules.py @@ -2140,10 +2140,7 @@ def forward(self, x): ): x = torch.randn(*size, requires_grad=True) mod(x) - if torch._dynamo.config.inline_inbuilt_nn_modules: - self.assertEqual(cnts.frame_count, 1) - else: - self.assertEqual(cnts.frame_count, num_submodules) + self.assertEqual(cnts.frame_count, 1) def test_inline_inbuilt_nn_modules(self): size = (10, 10) @@ -2225,10 +2222,7 @@ def forward(self, x): ]: x = torch.randn(size) mod(x) - if torch._dynamo.config.inline_inbuilt_nn_modules: - self.assertEqual(cnts.frame_count, 2) - else: - self.assertEqual(cnts.frame_count, 2 * num_submodules) + self.assertEqual(cnts.frame_count, 2) def test_recursion(self): mod = MockModule() @@ -2800,16 +2794,10 @@ def foo(mod, x): mod = Mod() foo(mod, torch.rand([4])) - if torch._dynamo.config.inline_inbuilt_nn_modules: - self.assertEqual(compiles_without_buffers, 1) - else: - self.assertEqual(compiles_without_buffers, 0) + self.assertEqual(compiles_without_buffers, 1) foo(mod, torch.rand([4], dtype=torch.half)) - if torch._dynamo.config.inline_inbuilt_nn_modules: - self.assertEqual(compiles_without_buffers, 2) - else: - self.assertEqual(compiles_without_buffers, 1) + self.assertEqual(compiles_without_buffers, 2) class Mod2(Mod): def __setattr__(self, name, value): diff --git a/test/dynamo/test_repros.py b/test/dynamo/test_repros.py index 7c5bf178fc44d..66541b9cacfbf 100644 --- a/test/dynamo/test_repros.py +++ b/test/dynamo/test_repros.py @@ -1399,12 +1399,8 @@ def test_reformer_eval(self): def test_reformer_train(self): with torch.enable_grad(): cnt = self._reformer(nopython=False) - expected_op_count = ( - """10""" if torch._dynamo.config.inline_inbuilt_nn_modules else """4""" - ) - self.assertExpectedInline(cnt.frame_count, """1""") - self.assertExpectedInline(cnt.op_count, expected_op_count) + self.assertExpectedInline(cnt.op_count, """10""") def test_longformer_chunk(self): input1 = torch.randn([1, 4096, 1]) diff --git a/test/dynamo/test_structured_trace.py b/test/dynamo/test_structured_trace.py index 526eb88cbda8f..a0230ca562b8b 100644 --- a/test/dynamo/test_structured_trace.py +++ b/test/dynamo/test_structured_trace.py @@ -655,67 +655,9 @@ def forward(self, x): dist.destroy_process_group() - if not torch._dynamo.config.inline_inbuilt_nn_modules: - self.assertExpectedInline( - self.buffer.getvalue(), - """\ -{"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} -{"artifact": {"name": "dynamo_graph_break_reason", "encoding": "string"}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 1, "has_payload": "HASH"} -{"compilation_metrics": "METRICS", "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 1} -{"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 1, "frame_compile_id": 0, "attempt": 0} -{"artifact": {"name": "dynamo_graph_break_reason", "encoding": "string"}, "rank": 0, "frame_id": 1, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 1, "frame_compile_id": 0, "attempt": 1, "has_payload": "HASH"} -{"compilation_metrics": "METRICS", "rank": 0, "frame_id": 1, "frame_compile_id": 0, "attempt": 1} -{"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 2, "frame_compile_id": 0, "attempt": 0} -{"artifact": {"name": "dynamo_graph_break_reason", "encoding": "string"}, "rank": 0, "frame_id": 2, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 2, "frame_compile_id": 0, "attempt": 1, "has_payload": "HASH"} -{"compilation_metrics": "METRICS", "rank": 0, "frame_id": 2, "frame_compile_id": 0, "attempt": 1} -{"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} -{"describe_storage": {"id": 0, "describer_id": "ID", "size": 4194304}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} -{"describe_tensor": {"id": 0, "ndim": 2, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1024, 1024], "is_leaf": true, "stride": [1024, 1], "storage": 0, "view_func": "VIEW_FUNC", "describer_id": "ID"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} -{"describe_source": {"describer_id": "ID", "id": 0, "source": "L['x']"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} -{"dynamo_output_graph": {"sizes": {"l_x_": [1024, 1024], "l__self___layers_0": [1024, 1024], "l__self___layers_1": [1024, 1024]}}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"optimize_ddp_split_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"optimize_ddp_split_child": {"name": "submod_0"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"optimize_ddp_split_child": {"name": "submod_1"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"describe_storage": {"id": 0, "describer_id": "ID", "size": 4194304}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} -{"describe_tensor": {"id": 0, "ndim": 2, "dtype": "torch.float32", "device": "device(type='cuda', index=0)", "size": [1024, 1024], "is_leaf": true, "stride": [1024, 1], "storage": 0, "view_func": "VIEW_FUNC", "describer_id": "ID"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} -{"describe_source": {"describer_id": "ID", "id": 0, "source": "L['x']"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} -{"artifact": {"name": "before_pre_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "after_pre_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"aot_joint_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "torch._functorch.config", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "aot_forward_graph_fw_metadata", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"aot_forward_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"aot_backward_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "fx_graph_runnable", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "before_post_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "inductor_post_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"inductor_output_code": {"filename": "FILENAME", "file_path": "FILENAME"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "fx_graph_cache_miss", "encoding": "json"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "aotautograd_cache_bypass", "encoding": "json"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "before_pre_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "after_pre_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"aot_joint_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "torch._functorch.config", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "aot_forward_graph_fw_metadata", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"aot_forward_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"aot_backward_graph": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "fx_graph_runnable", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "before_post_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "inductor_post_grad_graph", "encoding": "string"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"inductor_output_code": {"filename": "FILENAME", "file_path": "FILENAME"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "fx_graph_cache_miss", "encoding": "json"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"artifact": {"name": "aotautograd_cache_bypass", "encoding": "json"}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} -{"compilation_metrics": "METRICS", "rank": 0, "frame_id": 3, "frame_compile_id": 0, "attempt": 0} -""", # noqa: B950 - ) - else: - self.assertExpectedInline( - self.buffer.getvalue(), - """\ + self.assertExpectedInline( + self.buffer.getvalue(), + """\ {"dynamo_start": {"stack": "STACK"}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 0} {"artifact": {"name": "dynamo_graph_break_reason", "encoding": "string"}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} {"describe_storage": {"id": 0, "describer_id": "ID", "size": 4194304}, "rank": 0, "frame_id": 0, "frame_compile_id": 0, "attempt": 1} @@ -798,7 +740,7 @@ def forward(self, x): {"dynamo_cpp_guards_str": {}, "rank": 0, "frame_id": 4, "frame_compile_id": 0, "attempt": 0, "has_payload": "HASH"} {"compilation_metrics": "METRICS", "rank": 0, "frame_id": 4, "frame_compile_id": 0, "attempt": 0} """, # noqa: B950 - ) + ) self.assertParses() diff --git a/test/dynamo/test_utils.py b/test/dynamo/test_utils.py index 2cd19fb3334d6..721f99bc317d4 100644 --- a/test/dynamo/test_utils.py +++ b/test/dynamo/test_utils.py @@ -84,7 +84,6 @@ def test_larger_multiplier_for_even_smaller_tensor(self): @dynamo_config.patch( { "log_compilation_metrics": True, - "inline_inbuilt_nn_modules": False, } ) def test_graph_break_counting(self): @@ -445,7 +444,6 @@ def backward(grad_output): @dynamo_config.patch( { "log_compilation_metrics": True, - "inline_inbuilt_nn_modules": False, } ) @inductor_config.patch( @@ -688,7 +686,7 @@ def filter_expected(s: str) -> str: 'compile_time_autotune_time_us': None, 'compiler_config': None, 'compliant_custom_ops': set(), - 'config_inline_inbuilt_nn_modules': False, + 'config_inline_inbuilt_nn_modules': True, 'config_suppress_errors': False, 'cuda_version': None, 'cudagraph_skip_reason': None, @@ -707,11 +705,11 @@ def filter_expected(s: str) -> str: 'fail_user_frame_lineno': None, 'frame_key': '1', 'gc_time_us': 0, - 'graph_input_count': 1, - 'graph_node_count': 3, + 'graph_input_count': 3, + 'graph_node_count': 5, 'graph_node_shapes': None, 'graph_op_count': 1, - 'guard_count': 10, + 'guard_count': 31, 'has_guarded_code': True, 'inductor_code_gen_cumulative_compile_time_us': 0, 'inductor_compile_time_s': 0.0, @@ -781,7 +779,7 @@ def filter_expected(s: str) -> str: 'compile_time_autotune_time_us': None, 'compiler_config': None, 'compliant_custom_ops': set(), - 'config_inline_inbuilt_nn_modules': False, + 'config_inline_inbuilt_nn_modules': True, 'config_suppress_errors': False, 'cuda_version': None, 'cudagraph_skip_reason': None, @@ -800,11 +798,11 @@ def filter_expected(s: str) -> str: 'fail_user_frame_lineno': None, 'frame_key': '1', 'gc_time_us': 0, - 'graph_input_count': 1, - 'graph_node_count': 3, + 'graph_input_count': 3, + 'graph_node_count': 5, 'graph_node_shapes': None, 'graph_op_count': 1, - 'guard_count': 10, + 'guard_count': 31, 'has_guarded_code': True, 'inductor_code_gen_cumulative_compile_time_us': 0, 'inductor_compile_time_s': 0.0, @@ -888,7 +886,7 @@ def filter_expected(s: str) -> str: 'compile_time_autotune_time_us': None, 'compiler_config': None, 'compliant_custom_ops': None, - 'config_inline_inbuilt_nn_modules': False, + 'config_inline_inbuilt_nn_modules': True, 'config_suppress_errors': False, 'cuda_version': None, 'cudagraph_skip_reason': None, @@ -981,7 +979,7 @@ def filter_expected(s: str) -> str: 'compile_time_autotune_time_us': None, 'compiler_config': None, 'compliant_custom_ops': None, - 'config_inline_inbuilt_nn_modules': False, + 'config_inline_inbuilt_nn_modules': True, 'config_suppress_errors': False, 'cuda_version': None, 'cudagraph_skip_reason': None, diff --git a/test/functorch/test_control_flow.py b/test/functorch/test_control_flow.py index 72df12ad1822f..4c5a6b0cf0993 100644 --- a/test/functorch/test_control_flow.py +++ b/test/functorch/test_control_flow.py @@ -5910,10 +5910,9 @@ def test_while_loop_simple_with_linear_compile_check_graph(self): torch.compile(fn, backend=backend)(*inp) self.assertEqual(len(backend.graphs), 1) gm = backend.graphs[0] - if torch._dynamo.config.inline_inbuilt_nn_modules: - self.assertExpectedInline( - normalize_gm(gm.print_readable(print_output=False)), - """\ + self.assertExpectedInline( + normalize_gm(gm.print_readable(print_output=False)), + """\ class GraphModule(torch.nn.Module): def forward(self, L_iter_: "i64[]", L_x_: "f32[2, 2]", L_self_buffers_dec_: "i64[]", L_self_modules_linear_parameters_weight_: "f32[2, 2]", L_self_modules_linear_parameters_bias_: "f32[2]"): l_iter_ = L_iter_ @@ -5941,7 +5940,7 @@ def forward(self, child_2: "i64[]", child_3: "f32[2, 2]", l_self_buffers_dec__co child_4: "f32[2, 2]" = torch._C._nn.linear(child_3, l_self_modules_linear_parameters_weight__body_fn, l_self_modules_linear_parameters_bias__body_fn); child_3 = l_self_modules_linear_parameters_weight__body_fn = l_self_modules_linear_parameters_bias__body_fn = None return (child, child_4) """, # noqa: B950 - ) + ) def test_while_loop_nested2_traced(self): fn, inp = WHILE_LOOP_TESTS["nested2"] diff --git a/test/inductor/test_cudagraph_trees.py b/test/inductor/test_cudagraph_trees.py index 3fa2ec8ffc127..40c5f2e8247df 100644 --- a/test/inductor/test_cudagraph_trees.py +++ b/test/inductor/test_cudagraph_trees.py @@ -22,7 +22,7 @@ from torch._inductor.codecache import FxGraphCache from torch._inductor.compile_fx import compile_fx_inner from torch._inductor.cudagraph_trees import cudagraphify_impl as tree_cudagraphify_impl -from torch._inductor.cudagraph_utils import FunctionID, PlaceholderInfo +from torch._inductor.cudagraph_utils import PlaceholderInfo from torch._inductor.test_case import TestCase as InductorTestCase from torch._inductor.utils import run_and_get_code from torch._ops import OpOverload @@ -3303,20 +3303,8 @@ def forward(self, x) -> torch.Tensor: foo.static_tensor = torch.ones((2, 2), device="cuda") foo.goo.linear.bias = torch.nn.Parameter(torch.ones((2,), device="cuda")) - if torch._dynamo.config.inline_inbuilt_nn_modules: - for _ in range(3): - foo(inp) - else: - # Run with specific function id to avoid dynamo recompiling - self.get_manager().run( - [ - foo.goo.linear.weight, - foo.goo.linear.bias, - foo.static_tensor, - inp, - ], - FunctionID(0), - ) + for _ in range(3): + foo(inp) self.assertEqual(self.get_manager().new_graph_id().id, 2) From c032b7d139df31f29e5242904ca17aa8167024c8 Mon Sep 17 00:00:00 2001 From: Bob Ren Date: Tue, 24 Mar 2026 04:54:51 +0000 Subject: [PATCH 0058/1172] [Dynamo] Reuse tracked objects for Triton prune_configs_by (#177874) Fix #177600 ## Summary 1. What is the root cause problem `DynamoTritonHOPifier.wrap_user_defined_obj` called `VariableBuilder._wrap()` directly for Triton autotuner fields. That bypassed `VariableBuilder.__call__()`'s side-effect deduplication, so the same mutable `kernel.configs` list could be tracked twice when a `prune_configs_by` kernel was invoked twice in one trace. 2. What is the proposed fix Wrap those objects through `VariableBuilder.__call__()` instead of `._wrap()` so Dynamo reuses already-tracked mutable objects and still installs the right guards. 3. Why the proposed fix is the right long term fix `VariableBuilder.__call__()` is Dynamo's canonical entrypoint for building tracked objects. Reusing it keeps Triton autotuner wrapping aligned with the rest of Dynamo's mutation-tracking behavior instead of duplicating partial wrapping logic in this special case. ## Testing - Added a regression test covering two `prune_configs_by` Triton kernel calls inside one compiled function. Drafted via Codex, published after manual review by @bobrenjc93 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177874 Approved by: https://github.com/Lucaskabela --- test/inductor/test_triton_kernels.py | 45 ++++++++++++++++++++++++++++ torch/_dynamo/variables/functions.py | 5 +++- 2 files changed, 49 insertions(+), 1 deletion(-) diff --git a/test/inductor/test_triton_kernels.py b/test/inductor/test_triton_kernels.py index df71a8171ed1e..e9113f68b3c11 100644 --- a/test/inductor/test_triton_kernels.py +++ b/test/inductor/test_triton_kernels.py @@ -4885,6 +4885,51 @@ def f(dst, src, add_float, N): self.assertEqual(counter.op_count, 2) + @requires_gpu + @common_utils.parametrize("backend", ["eager", "aot_eager", "inductor"]) + def test_triton_kernel_prune_configs_by_called_twice(self, backend): + def early_config_prune(configs, named_args, **kwargs): + return [configs[0]] + + @triton.autotune( + configs=[ + triton.Config(kwargs={"BLOCK_SIZE": 128}), + triton.Config(kwargs={"BLOCK_SIZE": 256}), + ], + key=["N"], + prune_configs_by={"early_config_prune": early_config_prune}, + ) + @triton.jit + def add_kernel( + x_ptr, + y_ptr, + out_ptr, + N, + BLOCK_SIZE: tl.constexpr, + ): + offsets = tl.program_id(0) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) + mask = offsets < N + x = tl.load(x_ptr + offsets, mask=mask) + y = tl.load(y_ptr + offsets, mask=mask) + tl.store(out_ptr + offsets, x + y, mask=mask) + + def call_kernel(x, y): + out = torch.empty_like(x) + n_elements = x.numel() + grid = lambda meta: (triton.cdiv(n_elements, meta["BLOCK_SIZE"]),) + add_kernel[grid](x, y, out, N=n_elements) + return out + + @torch.compile(fullgraph=True, backend=backend) + def f(x, y): + out = call_kernel(x, y) + return call_kernel(out, y) + + x = torch.randn(1024, device=GPU_TYPE) + y = torch.randn(1024, device=GPU_TYPE) + + self.assertEqual(f(x, y), x + y + y) + # see: https://github.com/triton-lang/triton/blob/67ea999935f4511a535a25bdecb27e79e3c3af41/python/test/unit/language/test_decorator.py#L31 @requires_gpu @common_utils.parametrize("non_strict", [True, False]) diff --git a/torch/_dynamo/variables/functions.py b/torch/_dynamo/variables/functions.py index a8afcd51e5ca5..57f296215093e 100644 --- a/torch/_dynamo/variables/functions.py +++ b/torch/_dynamo/variables/functions.py @@ -3073,9 +3073,12 @@ def wrap_user_defined_obj( from .builder import VariableBuilder assert tx is not None + # Route through VariableBuilder.__call__ so already-tracked mutable + # objects (for example autotuner config lists) are reused instead of + # being registered for mutation twice in the same trace. wrapped_user_obj = VariableBuilder( tx, AttrSource(variable.kernel_source, f"{name}") - )._wrap(user_obj) + )(user_obj) return wrapped_user_obj def maybe_unpack_configs( From e534ad51600e9ef595a181084ba9487c87c3c1c8 Mon Sep 17 00:00:00 2001 From: Bob Ren Date: Tue, 24 Mar 2026 04:55:35 +0000 Subject: [PATCH 0059/1172] Fix detach_ autograd metadata tracking in Dynamo (#177875) ## Summary Fix #176854 ### Root cause Dynamo traces `Tensor.detach_()` into the FX graph, but it keeps using the pre-detach `TensorVariable` metadata for later attribute reads. Queries like `y.requires_grad` therefore observe stale autograd state even though the in-place detach has already been recorded. ### Proposed fix Add a dedicated Dynamo `detach_()` handler for tensor variables. The handler traces the in-place op, runs fake propagation for it, and then synchronizes the tensor metadata from the mutated fake tensor so later reads reflect the detached state. Add a regression test that checks `requires_grad` and `grad_fn` after `detach_()` under `torch.compile(..., backend="eager")`. ### Why this is the right long term fix The bug is in Dynamo's metadata bookkeeping, not in the underlying op execution. Refreshing the tensor variable metadata at the point where Dynamo traces `detach_()` keeps compile-time state aligned with the FX graph side effect and avoids patching over the mismatch later in the pipeline. Drafted via Codex, published after manual review by @bobrenjc93 Pull Request resolved: https://github.com/pytorch/pytorch/pull/177875 Approved by: https://github.com/Lucaskabela --- test/dynamo/test_misc.py | 14 ++++++++++++++ torch/_dynamo/variables/tensor.py | 14 ++++++++++++++ 2 files changed, 28 insertions(+) diff --git a/test/dynamo/test_misc.py b/test/dynamo/test_misc.py index a70378e327ca0..33ffba0c86390 100644 --- a/test/dynamo/test_misc.py +++ b/test/dynamo/test_misc.py @@ -14640,6 +14640,20 @@ def fn(x): self.assertEqual(x_ref.grad, x_test.grad) + def test_detach_inplace_on_intermediate_updates_metadata(self): + def fn(x): + y = x * 2 + y.detach_() + return y + 1, y.requires_grad, y.grad_fn is None + + x = torch.randn(3, 3, requires_grad=True) + ref = fn(x.clone()) + result = torch.compile(fn, backend="eager", fullgraph=True)(x.clone()) + + self.assertEqual(ref, result) + self.assertFalse(result[1]) + self.assertTrue(result[2]) + def test_requires_grad_on_intermediate(self): def fn(x): y = x * 2 diff --git a/torch/_dynamo/variables/tensor.py b/torch/_dynamo/variables/tensor.py index a602cbc875827..dc0d09c2fe057 100644 --- a/torch/_dynamo/variables/tensor.py +++ b/torch/_dynamo/variables/tensor.py @@ -1884,6 +1884,20 @@ def method_requires_grad_( ) return self + def method_detach_(self, tx: "InstructionTranslator") -> "TensorVariable": + from .builder import wrap_fx_proxy + + proxy = tx.output.create_proxy( + "call_method", + "detach_", + (self.as_proxy(),), + {}, + ) + # Run the fake op so the proxy metadata reflects the detached tensor state. + wrap_fx_proxy(tx, proxy) + self.synchronize_attributes(tx) + return self + def method_share_memory_(self) -> NoReturn: unimplemented( gb_type="Unsupported Tensor.share_memory_() call", From fc879df85c4ad676e8f4706ea204e73f1f7653d1 Mon Sep 17 00:00:00 2001 From: Nikita Shulga Date: Fri, 20 Mar 2026 10:13:55 -0700 Subject: [PATCH 0060/1172] [BE] Use `highest_precision_float` primitive throughout tests (#177976) Followup after https://github.com/pytorch/pytorch/pull/177766 Also introduce `highest_precision_complex` Pull Request resolved: https://github.com/pytorch/pytorch/pull/177976 Approved by: https://github.com/manuelcandales --- test/inductor/test_aot_inductor.py | 8 +++++-- test/inductor/test_torchinductor.py | 14 ++++++++----- test/test_indexing.py | 3 ++- test/test_reductions.py | 9 ++++---- test/test_sparse.py | 11 +++++----- torch/testing/_internal/common_dtype.py | 7 +++++++ .../_internal/common_methods_invocations.py | 21 ++++++++++--------- .../_internal/opinfo/definitions/sparse.py | 6 ++++-- 8 files changed, 50 insertions(+), 29 deletions(-) diff --git a/test/inductor/test_aot_inductor.py b/test/inductor/test_aot_inductor.py index aa9c68a26e634..e7130917185e3 100644 --- a/test/inductor/test_aot_inductor.py +++ b/test/inductor/test_aot_inductor.py @@ -61,6 +61,10 @@ e5m2_type, skipCUDAIf, ) +from torch.testing._internal.common_dtype import ( + highest_precision_complex, + highest_precision_float, +) from torch.testing._internal.common_quantization import ( _group_quantize_tensor, skip_if_no_torchvision, @@ -3632,7 +3636,7 @@ def forward(self, x): # Call eval() here so that batch_norm won't update the running stats # Use float64 to avoid numeric difference failure - dtype = torch.float32 if self.device == "mps" else torch.float64 + dtype = highest_precision_float(self.device) model = Model().to(device=self.device, dtype=dtype).eval() example_inputs = (torch.randn(4, 3, 64, 64, device=self.device, dtype=dtype),) self.check_model(model, example_inputs) @@ -5135,7 +5139,7 @@ def forward(self, x0, x1, x2): ) x2 = torch.tensor( 128, - dtype=torch.complex128 if self.device != "mps" else torch.complex64, + dtype=highest_precision_complex(self.device), device=self.device, ) inputs.append(x0) diff --git a/test/inductor/test_torchinductor.py b/test/inductor/test_torchinductor.py index 20c0132094314..1eee817c01e07 100644 --- a/test/inductor/test_torchinductor.py +++ b/test/inductor/test_torchinductor.py @@ -83,7 +83,11 @@ expectedFailureXPU, largeTensorTest, ) -from torch.testing._internal.common_dtype import all_types, get_all_dtypes +from torch.testing._internal.common_dtype import ( + all_types, + get_all_dtypes, + highest_precision_float, +) from torch.testing._internal.common_quantization import ( _dynamically_quantize_per_channel, _group_quantize_tensor_symmetric, @@ -3336,7 +3340,7 @@ def test_round_correctness(self): def fn(a): return torch.round(a) - dtype = torch.float64 if self.device != "mps" else torch.float32 + dtype = highest_precision_float(self.device) self.common( fn, [torch.arange(-10, 10, 0.1, dtype=dtype)], @@ -4650,7 +4654,7 @@ def fn(x, y): x1 = torch.randn(30, device=self.device) x2 = torch.randn(36, device=self.device) - dtype = torch.float64 if self.device != "mps" else torch.float32 + dtype = highest_precision_float(self.device) y = torch.ones(1, dtype=dtype, device=self.device) self.assertEqual(torch.compile(fn)(x1, y), fn(x1, y)) @@ -6535,7 +6539,7 @@ def test_polar(self): def fn(dist, angle): return torch.polar(dist, angle) - dtype = torch.float64 if self.device != "mps" else torch.float32 + dtype = highest_precision_float(self.device) inp = ( torch.tensor([1, 2], dtype=dtype), torch.tensor([np.pi / 2, 5 * np.pi / 4], dtype=dtype), @@ -13684,7 +13688,7 @@ def fn(tensor, index, source): return out device = "cpu" - dtype = torch.double if self.device != "mps" else torch.float32 + dtype = highest_precision_float(self.device) tensor = torch.rand((1,), dtype=dtype, device=device) index = torch.tensor([0], dtype=torch.long, device=device) source = torch.rand((1,), dtype=dtype, device=device) diff --git a/test/test_indexing.py b/test/test_indexing.py index e7ce521a79639..87c3ddbc56e74 100644 --- a/test/test_indexing.py +++ b/test/test_indexing.py @@ -32,6 +32,7 @@ all_types_and, all_types_and_complex_and, all_types_complex_float8_and, + highest_precision_float, ) from torch.testing._internal.common_utils import ( DeterministicGuard, @@ -1217,7 +1218,7 @@ def func1(x, i, v): @onlyNativeDeviceTypes def test_index_put_accumulate_duplicate_indices(self, device): - dtype = torch.float if device.startswith("mps") else torch.double + dtype = highest_precision_float(device) for i in range(1, 512): # generate indices by random walk, this will create indices with # lots of duplicates interleaved with each other diff --git a/test/test_reductions.py b/test/test_reductions.py index a6afa48308557..f50d53d496802 100644 --- a/test/test_reductions.py +++ b/test/test_reductions.py @@ -14,7 +14,8 @@ from torch import inf, nan from torch.testing import make_tensor from torch.testing._internal.common_dtype import ( - all_types_and_complex_and, get_all_math_dtypes, integral_types, complex_types, floating_types_and, + all_types_and_complex_and, get_all_math_dtypes, highest_precision_float, + integral_types, complex_types, floating_types_and, integral_types_and, floating_and_complex_types_and, all_types_and, all_types, ) from torch.testing._internal.common_utils import ( @@ -1605,8 +1606,8 @@ def test_bucketization(self, device): self.assertEqual(torch.bucketize(values_0_el, boundaries), expected_result) # nan input - values_nan = torch.tensor([1.0, float('nan'), 2.0, float('nan')], device=device, dtype=torch.float64) - boundaries = torch.tensor([0.0, 1.0, 2.0, 3.0], device=device, dtype=torch.float64) + values_nan = torch.tensor([1.0, float('nan'), 2.0, float('nan')], device=device, dtype=highest_precision_float(device)) + boundaries = torch.tensor([0.0, 1.0, 2.0, 3.0], device=device, dtype=highest_precision_float(device)) expected_result = torch.tensor([1, 4, 2, 4], device=device) self.assertEqual(torch.searchsorted(boundaries, values_nan), expected_result) expected_result = torch.tensor([2, 4, 3, 4], device=device) @@ -1615,7 +1616,7 @@ def test_bucketization(self, device): # type promotion and non contiguous tensors values_3d_permute = values_3d.permute(2, 1, 0).to(torch.int32) - boundaries_permute = values_3d.permute(2, 1, 0).to(torch.float64) + boundaries_permute = values_3d.permute(2, 1, 0).to(highest_precision_float(device)) expected_result = torch.tensor([[[0, 0], [0, 1]], [[2, 0], [0, 1]], [[2, 0], [0, 0]]], device=device) if self.device_type != 'xla': self.assertWarnsRegex( diff --git a/test/test_sparse.py b/test/test_sparse.py index e53de3cb8ecb8..b444c71131772 100644 --- a/test/test_sparse.py +++ b/test/test_sparse.py @@ -28,7 +28,8 @@ (op_db, reduction_ops, sparse_unary_ufuncs, sparse_masked_reduction_ops, binary_ufuncs) from torch.testing._internal.common_dtype import ( all_types, all_types_and_complex, all_mps_types, all_types_and_complex_and, floating_and_complex_types, - floating_and_complex_types_and, integral_types, floating_types_and, + floating_and_complex_types_and, highest_precision_complex, highest_precision_float, + integral_types, floating_types_and, ) from torch.testing._internal.opinfo.definitions.sparse import validate_sample_input_sparse from torch.testing._internal.opinfo.refs import ( @@ -333,7 +334,7 @@ def _test_print(self, device, dtype, coalesced): if values.dtype == torch.double: dtypes.append(torch.float) else: - dtypes.append(torch.double if values.device != torch.device("mps:0") else torch.float32) + dtypes.append(highest_precision_float(values.device)) for dtype in dtypes: printed.append(f"########## {dtype} ##########") x = sp_tensor.detach().to(dtype) @@ -583,7 +584,7 @@ def fn(x): x.requires_grad_(True) gradcheck(fn, (x,)) - values_types = [torch.double, torch.cdouble] if device != "mps:0" else [torch.float32, torch.complex64] + values_types = [highest_precision_float(device), highest_precision_complex(device)] for value_type in values_types: i = self.index_tensor([ [0, 1, 2, 2], @@ -628,7 +629,7 @@ def fn(x): def test_to_sparse(self, device, dtype, coalesced): shape = [5, 2, 10, 4] max_nnz = 1 - dtypes = [torch.double, torch.cdouble] if device != "mps:0" else [torch.float32, torch.complex64] + dtypes = [highest_precision_float(device), highest_precision_complex(device)] for value_type in dtypes: for dim, dim_sz in enumerate(shape, 1): max_nnz *= dim_sz @@ -963,7 +964,7 @@ def test_Sparse_to_Sparse_copy_(self, device, dtype, coalesced): x1.copy_(x2) self.assertEqual(x1_dtype, x1.dtype) - x2 = x2.to(torch.float64) if device != "mps:0" else x2.to(torch.float32) + x2 = x2.to(highest_precision_float(device)) x1_dtype = x1.dtype x1.copy_(x2) self.assertEqual(x1_dtype, x1.dtype) diff --git a/torch/testing/_internal/common_dtype.py b/torch/testing/_internal/common_dtype.py index c5f3ad3f390dc..6e62009b6b601 100644 --- a/torch/testing/_internal/common_dtype.py +++ b/torch/testing/_internal/common_dtype.py @@ -229,6 +229,13 @@ def highest_precision_float(device): return torch.float64 +def highest_precision_complex(device): + if torch.device(device).type == "mps": + return torch.complex64 + else: + return torch.complex128 + + float_to_corresponding_complex_type_map = { torch.float16: torch.complex32, torch.float32: torch.complex64, diff --git a/torch/testing/_internal/common_methods_invocations.py b/torch/testing/_internal/common_methods_invocations.py index fd46dbcd52f11..e4d225db954b7 100644 --- a/torch/testing/_internal/common_methods_invocations.py +++ b/torch/testing/_internal/common_methods_invocations.py @@ -23,6 +23,7 @@ _dispatch_dtypes, floating_types, floating_types_and, complex_types, floating_and_complex_types, floating_and_complex_types_and, all_types_and_complex_and, all_types_and, all_types_and_complex, integral_types_and, empty_types, complex_types_and, integral_types, custom_types, all_types_complex_float8_and, float8_types, + highest_precision_complex, highest_precision_float, ) from torch.testing._internal.common_device_type import ( @@ -1616,7 +1617,7 @@ def sample_inputs_like_fns(self, device, dtype, requires_grad, **kwargs): ((S,), {'dtype': dtype, 'device': device}), # Hard-code some dtypes/devices. We want to test cases where the # (dtype, device) is different from the input's (dtype, device) - ((S,), {'dtype': torch.double if device != 'mps:0' else torch.float}), + ((S,), {'dtype': highest_precision_float(device)}), ((S,), {'device': 'cpu'}), ((S,), {'dtype': torch.double, 'device': 'cpu'}), ] @@ -1810,7 +1811,7 @@ def sample_inputs_new_fns(self, device, dtype, requires_grad, *, is_strided=Fals ((S,), (2, 3), (7, 8), {'dtype': dtype, 'device': device}), # Hard-code some dtypes/devices. We want to test cases where the # (dtype, device) is different from the input's (dtype, device) - ((S,), (10,), (S,), {'dtype': torch.double if device != 'mps:0' else torch.float}), + ((S,), (10,), (S,), {'dtype': highest_precision_float(device)}), ((S,), (1, 1, 12), (S, L, M), {'device': 'cpu'}), ((S,), (2, 2, 2), (L, M, S), {'dtype': torch.double, 'device': 'cpu'}), ] @@ -1933,7 +1934,7 @@ def sample_inputs_full_like(self, device, dtype, requires_grad, **kwargs): def get_val(dtype): return make_tensor([], dtype=dtype, device="cpu").item() - double_dtype = torch.double if device != "mps:0" else torch.float + double_dtype = highest_precision_float(device) inputs = [ ((), get_val(dtype), {}), ((S, S), get_val(dtype), {}), @@ -4058,8 +4059,8 @@ def sample_inputs_conv1d(op_info, device, dtype, requires_grad, **kwargs): def error_inputs_conv1d(opinfo, device, **kwargs): - dtype = torch.float64 if device != 'mps:0' else torch.float32 - cdtype = torch.complex128 if device != 'mps:0' else torch.complex64 + dtype = highest_precision_float(device) + cdtype = highest_precision_complex(device) make_arg = partial(make_tensor, device=device, dtype=dtype) make_int_arg = partial(make_tensor, device=device, dtype=torch.int64) make_complex_arg = partial(make_tensor, device=device, dtype=cdtype) @@ -4120,8 +4121,8 @@ def error_inputs_conv1d(opinfo, device, **kwargs): def error_inputs_conv2d(opinfo, device, **kwargs): - dtype = torch.float64 if device != 'mps:0' else torch.float32 - cdtype = torch.complex128 if device != 'mps:0' else torch.complex64 + dtype = highest_precision_float(device) + cdtype = highest_precision_complex(device) make_arg = partial(make_tensor, device=device, dtype=dtype) make_int_arg = partial(make_tensor, device=device, dtype=torch.int64) make_complex_arg = partial(make_tensor, device=device, dtype=cdtype) @@ -4259,8 +4260,8 @@ def sample_inputs_conv3d(opinfo, device, dtype, requires_grad, **kwargs): def error_inputs_conv3d(opinfo, device, **kwargs): - dtype = torch.float64 if device != 'mps:0' else torch.float32 - cdtype = torch.complex128 if device != 'mps:0' else torch.complex64 + dtype = highest_precision_float(device) + cdtype = highest_precision_complex(device) make_arg = partial(make_tensor, device=device, dtype=dtype) make_int_arg = partial(make_tensor, device=device, dtype=torch.int64) make_complex_arg = partial(make_tensor, device=device, dtype=cdtype) @@ -9973,7 +9974,7 @@ def __call__(self, opinfo, device, dtype, requires_grad, **kwargs): for num_tensors, ord, out_dtype, intersperse_empty_tensors in product( num_input_tensors, (0, 1, 2, -1, -2, float('inf'), float('-inf')), - (None,) + (torch.complex128,) if dtype in complex_types() else (torch.float64,), + (None,) + (highest_precision_complex(device),) if dtype in complex_types() else (highest_precision_float(device),), (True, False), ): # inf norm and negative norms on empty tensors is not supported by our reference func vector norm: diff --git a/torch/testing/_internal/opinfo/definitions/sparse.py b/torch/testing/_internal/opinfo/definitions/sparse.py index 7f8781ce87b24..c04009f1f2ac9 100644 --- a/torch/testing/_internal/opinfo/definitions/sparse.py +++ b/torch/testing/_internal/opinfo/definitions/sparse.py @@ -4,6 +4,7 @@ import torch from torch.testing import make_tensor # noqa: F401 +from torch.testing._internal.common_dtype import highest_precision_float from torch.testing._internal.opinfo.core import ( # noqa: F401 BinaryUfuncInfo, ErrorInput, @@ -769,8 +770,9 @@ def _sample_inputs_sparse_like_fns( tensor, args=(), kwargs=dict(device=device, dtype=dtype, layout=layout) ) - if dtype is not torch.float64: - yield SampleInput(tensor, args=(), kwargs=dict(dtype=torch.float64)) + hpf = highest_precision_float(device) + if dtype is not hpf: + yield SampleInput(tensor, args=(), kwargs=dict(dtype=hpf)) if torch.cuda.is_available(): other_device = "cuda" if tensor.device.type == "cpu" else "cpu" From 1c9ac8060dec5eeddbfd5eeb0f422aaaf87581a6 Mon Sep 17 00:00:00 2001 From: "Wang, Chuanqi" Date: Tue, 24 Mar 2026 06:19:18 +0000 Subject: [PATCH 0061/1172] [CI] Use absolute action references in setup-xpu for cross-repo compatibility (#178143) As the title Pull Request resolved: https://github.com/pytorch/pytorch/pull/178143 Approved by: https://github.com/scotts, https://github.com/Skylion007 --- .github/actions/setup-xpu/action.yml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/.github/actions/setup-xpu/action.yml b/.github/actions/setup-xpu/action.yml index dfae0a358259e..db96705c12dfe 100644 --- a/.github/actions/setup-xpu/action.yml +++ b/.github/actions/setup-xpu/action.yml @@ -44,7 +44,7 @@ runs: fi - name: Runner diskspace health check - uses: ./.github/actions/diskspace-cleanup + uses: pytorch/pytorch/.github/actions/diskspace-cleanup@main if: always() - name: Runner health check disconnect on failure @@ -67,4 +67,4 @@ runs: echo "GPU_FLAG=--device=/dev/mem --device=/dev/dri --group-add video --group-add $render_gid" >> "${GITHUB_ENV}" - name: Login to ECR - uses: ./.github/actions/ecr-login + uses: pytorch/pytorch/.github/actions/ecr-login@main From 021759768d2680f96ce53a6abcc534197910bace Mon Sep 17 00:00:00 2001 From: "Wang, Chuanqi" Date: Tue, 24 Mar 2026 07:34:42 +0000 Subject: [PATCH 0062/1172] [CD] Update triton xpu release branch naming convention (#177723) As the triton xpu has standalone repo, and the update cadence in pytorch is different with upstream triton. To mitigate the version gap between those 2 kinds of triton, consider patch version into release branch naming for triton xpu. It just impact release wheel build. Pull Request resolved: https://github.com/pytorch/pytorch/pull/177723 Approved by: https://github.com/EikanWang, https://github.com/atalman --- .github/scripts/build_triton_wheel.py | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/.github/scripts/build_triton_wheel.py b/.github/scripts/build_triton_wheel.py index e5ac9c5937dfa..9408ac1beb9ae 100644 --- a/.github/scripts/build_triton_wheel.py +++ b/.github/scripts/build_triton_wheel.py @@ -81,9 +81,16 @@ def build_triton( check_call(["git", "clone", triton_repo, "triton"], cwd=tmpdir) if release: ver, rev, patch = version.split(".") - check_call( - ["git", "checkout", f"release/{ver}.{rev}.x"], cwd=triton_basedir - ) + if device == "xpu": + # XPU uses the patch version in the release branch name + check_call( + ["git", "checkout", f"release/{ver}.{rev}.{patch}"], + cwd=triton_basedir, + ) + else: + check_call( + ["git", "checkout", f"release/{ver}.{rev}.x"], cwd=triton_basedir + ) else: check_call(["git", "fetch", "origin", commit_hash], cwd=triton_basedir) check_call(["git", "checkout", commit_hash], cwd=triton_basedir) From 4e56ea301613b781a30238d6f9793e8e298c0146 Mon Sep 17 00:00:00 2001 From: "Cui, Lily" Date: Tue, 24 Mar 2026 08:09:02 +0000 Subject: [PATCH 0063/1172] [Inductor][X86] Remove deprecated fusion patterns (#173911) These quantization-related fusion patterns have been moved to torchao. Pull Request resolved: https://github.com/pytorch/pytorch/pull/173911 Approved by: https://github.com/Xia-Weiwen, https://github.com/mingfeima, https://github.com/jerryzh168 --- origin.txt | 17786 ++++++++++++++++ test/inductor/test_aot_inductor.py | 4 +- test/inductor/test_cpu_repro.py | 4 +- test/inductor/test_cpu_select_algorithm.py | 97 +- torch/_inductor/fx_passes/mkldnn_fusion.py | 2 - torch/_inductor/fx_passes/quantization.py | 2330 -- .../testing/_internal/common_quantization.py | 227 +- 7 files changed, 18062 insertions(+), 2388 deletions(-) create mode 100644 origin.txt diff --git a/origin.txt b/origin.txt new file mode 100644 index 0000000000000..48353154959e5 --- /dev/null +++ b/origin.txt @@ -0,0 +1,17786 @@ +/home/lily/pytorch/torch/utils/_config_module.py:540: FutureWarning: torch._dynamo.config.skip_code_recursive_on_recompile_limit_hit is deprecated and does not do anything. It will be removed in a future version of PyTorch. + config[key] = copy.deepcopy(getattr(self, key)) +/home/lily/pytorch/torch/_inductor/mkldnn_lowerings.py:1067: UserWarning: To copy construct from a tensor, it is recommended to use sourceTensor.detach().clone() or sourceTensor.detach().clone().requires_grad_(True), rather than torch.tensor(sourceTensor). + torch.tensor(w_zp_tensor, dtype=torch.int32), name=w_zp.get_name() +[ReinterpretView] Created with layout: size=[m_end - m_start, 128], stride=[128, 1], data=StorageBox( + ComputedBuffer(name='buf0', layout=FixedLayout('cpu', torch.uint8, size=[32, 128], stride=[128, 1]), data=Pointwise( + 'cpu', + torch.uint8, + def inner_fn(index): + i0, i1 = index + tmp0 = ops.load(arg8_1, i1 + 128 * i0) + tmp1 = ops.constant(17.890258428104524, torch.float32) + tmp2 = ops.constant(67, torch.float32) + tmp3 = tmp0 * tmp1 + tmp4 = ops.round(tmp3) + tmp5 = tmp4 + tmp2 + tmp6 = ops.constant(0, torch.float32) + tmp7 = ops.constant(127, torch.float32) + tmp8 = ops.maximum(tmp5, tmp6) + tmp9 = ops.minimum(tmp8, tmp7) + tmp10 = ops.to_dtype(tmp9, torch.uint8) + return tmp10 + , + ranges=[32, 128], + origin_node=quantize_per_tensor, + origins=OrderedSet([quantize_per_tensor]), + stack_traces = {, + File "/home/lily/pytorch/torch/_dynamo/functional_export.py", line 218, in forward, + res = self._export_root(*args, **kwargs), + File "/home/lily/pytorch/test/inductor/test_cpu_select_algorithm.py", line 1451, in forward, + res = self.epilogue(self.linear(x) + other), + File "/home/lily/pytorch/torch/testing/_internal/common_quantization.py", line 3304, in forward, + quantize_per_tensor_default = torch.ops.quantized_decomposed.quantize_per_tensor.default(, + , + } + ), _split_size=None, _original_inner_fn=None, _original_ranges=None, _original_reduction_ranges=None) +) +[ReinterpretView] Stack trace: +File "/home/lily/pytorch/test/inductor/test_cpu_select_algorithm.py", line 3315, in + run_tests() +File "/home/lily/pytorch/torch/_inductor/test_case.py", line 14, in run_tests + dynamo_run_tests(needs) +File "/home/lily/pytorch/torch/_dynamo/test_case.py", line 63, in run_tests + run_tests() +File "/home/lily/pytorch/torch/testing/_internal/common_utils.py", line 1402, in run_tests + unittest.main(argv=argv) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/main.py", line 105, in __init__ + self.runTests() +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/main.py", line 281, in runTests + self.result = testRunner.run(self.test) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/runner.py", line 240, in run + test(result) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/suite.py", line 84, in __call__ + return self.run(*args, **kwds) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/suite.py", line 122, in run + test(result) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/suite.py", line 84, in __call__ + return self.run(*args, **kwds) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/suite.py", line 122, in run + test(result) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/case.py", line 690, in __call__ + return self.run(*args, **kwds) +File "/home/lily/pytorch/torch/testing/_internal/common_device_type.py", line 522, in run + super().run(result=result) +File "/home/lily/pytorch/torch/testing/_internal/common_utils.py", line 3553, in run + self._run_custom( +File "/home/lily/pytorch/torch/testing/_internal/common_utils.py", line 3522, in _run_custom + super_run(result=result) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/case.py", line 634, in run + self._callTestMethod(testMethod) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/case.py", line 589, in _callTestMethod + if method() is not None: +File "/home/lily/pytorch/torch/testing/_internal/common_utils.py", line 3370, in wrapper + method(*args, **kwargs) +File "/home/lily/pytorch/torch/testing/_internal/common_device_type.py", line 430, in instantiated_test + result = test(self, **param_kwargs) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/contextlib.py", line 81, in inner + return func(*args, **kwds) +File "/home/lily/pytorch/test/inductor/test_cpu_select_algorithm.py", line 85, in wrapped + return fn(*args, **kwargs) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/unittest/mock.py", line 1396, in patched + return func(*newargs, **newkeywargs) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/contextlib.py", line 81, in inner + return func(*args, **kwds) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/contextlib.py", line 81, in inner + return func(*args, **kwds) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/contextlib.py", line 81, in inner + return func(*args, **kwds) +File "/home/lily/pytorch/torch/utils/_contextlib.py", line 124, in decorate_context + return func(*args, **kwargs) +File "/home/lily/pytorch/test/inductor/test_cpu_select_algorithm.py", line 1500, in test_quantized_linear_with_pointwise_binary + res = cfn(*inputs) +File "/home/lily/pytorch/torch/_dynamo/eval_frame.py", line 468, in __call__ + return super().__call__(*args, **kwargs) +File "/home/lily/pytorch/torch/nn/modules/module.py", line 1778, in _wrapped_call_impl + return self._call_impl(*args, **kwargs) +File "/home/lily/pytorch/torch/nn/modules/module.py", line 1789, in _call_impl + return forward_call(*args, **kwargs) +File "/home/lily/pytorch/torch/_dynamo/eval_frame.py", line 1030, in compile_wrapper + return fn(*args, **kwargs) +File "/home/lily/pytorch/torch/fx/graph_module.py", line 949, in call_wrapped + return self._wrapped_call(self, *args, **kwargs) +File "/home/lily/pytorch/torch/fx/graph_module.py", line 447, in __call__ + return super(self.cls, obj).__call__(*args, **kwargs) # type: ignore[misc] +File "/home/lily/pytorch/torch/nn/modules/module.py", line 1778, in _wrapped_call_impl + return self._call_impl(*args, **kwargs) +File "/home/lily/pytorch/torch/nn/modules/module.py", line 1884, in _call_impl + return inner() +File "/home/lily/pytorch/torch/nn/modules/module.py", line 1832, in inner + result = forward_call(*args, **kwargs) +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 2360, in __call__ + result = self._torchdynamo_orig_backend( +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 2096, in __call__ + result = self._inner_convert( +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 723, in __call__ + result = _compile( +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 1871, in _compile + guarded_code, tracer_output = compile_inner(code, one_graph, hooks) +File "/home/lily/pytorch/torch/_utils_internal.py", line 96, in wrapper_function + return function(*args, **kwargs) +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 1492, in compile_inner + result = _compile_inner(code, one_graph, hooks) +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 1551, in _compile_inner + dynamo_output = compile_frame( +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 1399, in compile_frame + bytecode, tracer_output = transform_code_object(code, transform) +File "/home/lily/pytorch/torch/_dynamo/bytecode_transformation.py", line 1626, in transform_code_object + tracer_output = transformations(instructions, code_options) +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 1371, in transform + tracer_output = trace_frame( +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 340, in _fn + return fn(*args, **kwargs) +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 857, in trace_frame + run_tracer() +File "/home/lily/pytorch/torch/_dynamo/convert_frame.py", line 838, in run_tracer + tracer.run() +File "/home/lily/pytorch/torch/_dynamo/symbolic_convert.py", line 1694, in run + while self.step(): +File "/home/lily/pytorch/torch/_dynamo/symbolic_convert.py", line 1360, in step + self.dispatch_table[inst.opcode](self, inst) +File "/home/lily/pytorch/torch/_dynamo/symbolic_convert.py", line 4933, in RETURN_VALUE + self._return(inst) +File "/home/lily/pytorch/torch/_dynamo/symbolic_convert.py", line 4915, in _return + all_stack_locals_metadata = self.output.compile_subgraph( +File "/home/lily/pytorch/torch/_dynamo/output_graph.py", line 1975, in compile_subgraph + self.compile_and_call_fx_graph(tx, pass2.graph_output_vars(), root) +File "/home/lily/pytorch/torch/_dynamo/output_graph.py", line 2549, in compile_and_call_fx_graph + compiled_fn = self.call_user_compiler(gm, self.example_inputs()) +File "/home/lily/pytorch/torch/_dynamo/output_graph.py", line 2716, in call_user_compiler + return self._call_user_compiler(gm, example_inputs) +File "/home/lily/pytorch/torch/_dynamo/output_graph.py", line 2774, in _call_user_compiler + compiled_fn = compiler_fn(gm, example_inputs) +File "/home/lily/pytorch/torch/_dynamo/repro/after_dynamo.py", line 156, in __call__ + compiled_gm = compiler_fn(gm, example_inputs) +File "/home/lily/pytorch/torch/__init__.py", line 2476, in __call__ + return compile_fx(model_, inputs_, config_patches=all_patches) +File "/home/lily/pytorch/torch/_inductor/compile_fx.py", line 2629, in compile_fx + return _maybe_wrap_and_compile_fx_main( +File "/home/lily/pytorch/torch/_inductor/compile_fx.py", line 2706, in _maybe_wrap_and_compile_fx_main + return _compile_fx_main( +File "/home/lily/pytorch/torch/_inductor/compile_fx.py", line 2903, in _compile_fx_main + return aot_autograd( +File "/home/lily/pytorch/torch/_dynamo/backends/common.py", line 123, in __call__ + cg = aot_module_simplified(gm, example_inputs, **self.kwargs) +File "/home/lily/pytorch/torch/_functorch/aot_autograd.py", line 1171, in aot_module_simplified + compiled_fn, _ = aot_stage2_compile( +File "/home/lily/pytorch/torch/_functorch/_aot_autograd/graph_compile.py", line 369, in aot_stage2_compile + return aot_stage2_inference(aot_state, aot_graph_capture) +File "/home/lily/pytorch/torch/_functorch/_aot_autograd/graph_compile.py", line 447, in aot_stage2_inference + compiled_fw = _aot_stage2b_inference_compile( +File "/home/lily/pytorch/torch/_functorch/_aot_autograd/graph_compile.py", line 413, in _aot_stage2b_inference_compile + return _aot_stage2b_compile_forward_or_inference( +File "/home/lily/pytorch/torch/_functorch/_aot_autograd/graph_compile.py", line 2515, in _aot_stage2b_compile_forward_or_inference + compiled_fw_func = compiler(fw_module, adjusted_flat_args) +File "/home/lily/pytorch/torch/_inductor/compile_fx.py", line 2140, in fw_compiler_freezing + optimized_function = inner_compile( +File "/home/lily/pytorch/torch/_inductor/compile_fx.py", line 824, in compile_fx_inner + return wrap_compiler_debug(_compile_fx_inner, compiler_name="inductor")( +File "/home/lily/pytorch/torch/_dynamo/repro/after_aot.py", line 273, in debug_wrapper + inner_compiled_fn = compiler_fn(gm, example_inputs) +File "/home/lily/pytorch/torch/_inductor/compile_fx.py", line 1020, in _compile_fx_inner + mb_compiled_graph = fx_codegen_and_compile( +File "/home/lily/pytorch/torch/_inductor/compile_fx.py", line 1796, in fx_codegen_and_compile + return scheme.codegen_and_compile(gm, example_inputs, inputs_to_check, graph_kwargs) +File "/home/lily/pytorch/torch/_inductor/compile_fx.py", line 1483, in codegen_and_compile + graph.run(*example_inputs) +File "/home/lily/pytorch/torch/_inductor/graph.py", line 1032, in run + return super().run(*args) +File "/home/lily/pytorch/torch/fx/interpreter.py", line 200, in run + self.env[node] = self.run_node(node) +File "/home/lily/pytorch/torch/_inductor/graph.py", line 1831, in run_node + result = super().run_node(n) +File "/home/lily/pytorch/torch/fx/interpreter.py", line 297, in run_node + return getattr(self, n.op)(n.target, args, kwargs) +File "/home/lily/pytorch/torch/_inductor/graph.py", line 1304, in call_function + return target(*args, **kwargs) +File "/home/lily/pytorch/torch/_inductor/fx_passes/quantization.py", line 657, in qlinear_binary + return L[computation_op](*computation_args) +File "/home/lily/pytorch/torch/_inductor/lowering.py", line 515, in wrapped + out = decomp_fn(*args, **kwargs) +File "/home/lily/pytorch/torch/_inductor/lowering.py", line 515, in wrapped + out = decomp_fn(*args, **kwargs) +File "/home/lily/pytorch/torch/_inductor/mkldnn_lowerings.py", line 1261, in qlinear_binary + CppGemmTemplate.add_choices( +File "/home/lily/pytorch/torch/_inductor/codegen/cpp_gemm_template.py", line 1124, in add_choices + template.maybe_append_choice(choices) +File "/home/lily/pytorch/torch/_inductor/select_algorithm.py", line 2759, in maybe_append_choice + return type(self._wrapped).maybe_append_choice(self, choices, **kwargs) +File "/home/lily/pytorch/torch/_inductor/codegen/common.py", line 2598, in maybe_append_choice + choices.append(self.generate(**kwargs)) +File "/home/lily/pytorch/torch/_inductor/select_algorithm.py", line 2762, in generate + choice_caller = self._wrapped.generate(**kwargs) +File "/home/lily/pytorch/torch/_inductor/codegen/cpp_template.py", line 54, in generate + code = kernel.render(self, **kwargs) +File "/home/lily/pytorch/torch/_inductor/codegen/cpp_template_kernel.py", line 52, in render + template.render(kernel=self, **kwargs), self.render_hooks +File "/home/lily/pytorch/torch/_inductor/codegen/cpp_gemm_template.py", line 1668, in render + return self._template_from_string(template_str).render(**options) +File "/home/lily/miniforge3/envs/lily_pytorch/lib/python3.12/site-packages/jinja2/environment.py", line 1293, in render + return self.environment.concat(self.root_render_func(ctx)) # type: ignore +File "