From a81a2e54ed71fe28f52d8599313756bc8bd7c33e Mon Sep 17 00:00:00 2001 From: henrylhtsang Date: Wed, 17 Sep 2025 17:59:10 -0700 Subject: [PATCH] [submodule] CUTLASS upgrade to 4.2.0 and change cutlass to cutlass_cppgen (#163092) Pull Request resolved: https://github.com/pytorch/pytorch/pull/163092 Approved by: https://github.com/drisspg, https://github.com/Skylion007 --- test/inductor/test_cutlass_backend.py | 2 +- test/inductor/test_cutlass_evt.py | 2 +- third_party/cutlass | 2 +- .../codegen/cuda/cutlass_lib_extensions/evt_extensions.py | 6 +++--- torch/_inductor/codegen/cuda/cutlass_utils.py | 6 +++--- 5 files changed, 9 insertions(+), 9 deletions(-) diff --git a/test/inductor/test_cutlass_backend.py b/test/inductor/test_cutlass_backend.py index d0618886660..149e32a2433 100644 --- a/test/inductor/test_cutlass_backend.py +++ b/test/inductor/test_cutlass_backend.py @@ -257,7 +257,7 @@ class TestCutlassBackend(TestCase): if config.is_fbcode(): import python_cutlass else: - import cutlass as python_cutlass # noqa: F401 + import cutlass_cppgen as python_cutlass # noqa: F401 import cutlass_library # noqa: F401 def test_cutlass_key(self): diff --git a/test/inductor/test_cutlass_evt.py b/test/inductor/test_cutlass_evt.py index cae9558d2ec..66f03762fa1 100644 --- a/test/inductor/test_cutlass_evt.py +++ b/test/inductor/test_cutlass_evt.py @@ -36,7 +36,7 @@ if try_import_cutlass(): if config.is_fbcode(): import python_cutlass # type: ignore[import-untyped, import-not-found] # noqa: F401 else: - import cutlass as python_cutlass # type: ignore[import-untyped, import-not-found] # noqa: F401 + import cutlass_cppgen as python_cutlass # type: ignore[import-untyped, import-not-found] # noqa: F401 CutlassTensor = python_cutlass.backend.evt.ir.tensor.Tensor BIAS_CODE = """def example_epilogue(accum, C, aux, bias): diff --git a/third_party/cutlass b/third_party/cutlass index e51efbfe18f..57e3cfb47a2 160000 --- a/third_party/cutlass +++ b/third_party/cutlass @@ -1 +1 @@ -Subproject commit e51efbfe18fe4f4cbb66ab814c55bf4aa0185491 +Subproject commit 57e3cfb47a2d9e0d46eb6335c3dc411498efa198 diff --git a/torch/_inductor/codegen/cuda/cutlass_lib_extensions/evt_extensions.py b/torch/_inductor/codegen/cuda/cutlass_lib_extensions/evt_extensions.py index 605b93dff59..a2beb9ecfc4 100644 --- a/torch/_inductor/codegen/cuda/cutlass_lib_extensions/evt_extensions.py +++ b/torch/_inductor/codegen/cuda/cutlass_lib_extensions/evt_extensions.py @@ -38,7 +38,7 @@ if try_import_cutlass(): if config.is_fbcode(): import python_cutlass # type: ignore[import-untyped, import-not-found] # noqa: F401 else: - import cutlass as python_cutlass # type: ignore[import-untyped, import-not-found] # noqa: F401 + import cutlass_cppgen as python_cutlass # type: ignore[import-untyped, import-not-found] # noqa: F401 from torch._inductor.codegen.cuda import cuda_env from torch._inductor.utils import IndentedBuffer @@ -174,7 +174,7 @@ non-contiguous layout, received stride: {stride} and shape: {shape}" def is_nested_visitor_type(t: type) -> bool: return ".".join([t.__module__, t.__qualname__]) in { "python_cutlass.backend.c_types.visitor_factory..VisitorType", - "cutlass.backend.c_types.visitor_factory..VisitorType", + "cutlass_cppgen.backend.c_types.visitor_factory..VisitorType", } buffer = IndentedBuffer() @@ -235,7 +235,7 @@ non-contiguous layout, received stride: {stride} and shape: {shape}" # Once again, need to check for local class type for stride tuple if str(arg_ty) in { ".TupleType'>", - ".TupleType'>", + ".TupleType'>", }: DEFAULT_STRIDE_LEN = 3 assert len(node.get_layout().stride) <= DEFAULT_STRIDE_LEN diff --git a/torch/_inductor/codegen/cuda/cutlass_utils.py b/torch/_inductor/codegen/cuda/cutlass_utils.py index 7ca33ea779c..bdbbbe58a3b 100644 --- a/torch/_inductor/codegen/cuda/cutlass_utils.py +++ b/torch/_inductor/codegen/cuda/cutlass_utils.py @@ -43,7 +43,7 @@ def move_cutlass_compiled_cache() -> None: if config.is_fbcode(): import python_cutlass # type: ignore[import-not-found] else: - import cutlass as python_cutlass # type: ignore[import-not-found] # noqa: F401 + import cutlass_cppgen as python_cutlass # type: ignore[import-not-found] # noqa: F401 # Check if the CACHE_FILE attribute exists in python_cutlass and if the file exists if not hasattr(python_cutlass, "CACHE_FILE") or not os.path.exists( @@ -118,7 +118,7 @@ def try_import_cutlass() -> bool: tmp_cutlass_full_path = os.path.abspath(os.path.join(cache_dir(), "torch_cutlass")) dst_link_library = path_join(tmp_cutlass_full_path, "cutlass_library") - dst_link_cutlass = path_join(tmp_cutlass_full_path, "cutlass") + dst_link_cutlass = path_join(tmp_cutlass_full_path, "cutlass_cppgen") dst_link_pycute = path_join(tmp_cutlass_full_path, "pycute") # mock modules to import cutlass @@ -156,7 +156,7 @@ def try_import_cutlass() -> bool: ) try: - import cutlass # noqa: F401, F811 + import cutlass_cppgen # noqa: F401, F811 import cutlass_library.generator # noqa: F401 import cutlass_library.library # noqa: F401 import cutlass_library.manifest # noqa: F401