From 2c1b4cdc70b2fab1a79f32fa91dee342ff7ee14a Mon Sep 17 00:00:00 2001 From: jasl Date: Wed, 15 Oct 2025 07:08:53 +0800 Subject: [PATCH 1/2] Upgrade to CUTLASS 4.2.1 --- 3rdparty/cutlass | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/3rdparty/cutlass b/3rdparty/cutlass index b2dd65dc864e..f3fde58372d3 160000 --- a/3rdparty/cutlass +++ b/3rdparty/cutlass @@ -1 +1 @@ -Subproject commit b2dd65dc864e09688245b316ac46c4a6cd07e15c +Subproject commit f3fde58372d33e9a5650ba7b80fc48b3b49d40c8 From 6ef9539f6134b04f082f8edb731d3c995a17ef46 Mon Sep 17 00:00:00 2001 From: jasl Date: Thu, 16 Oct 2025 03:45:51 +0800 Subject: [PATCH 2/2] Fix test: mbarrier.try_wait requires .target sm_90 or higher --- tests/python/tir-base/test_tir_ptx_cp_async.py | 2 +- .../tir-transform/test_tir_transform_inject_ptx_async_copy.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/python/tir-base/test_tir_ptx_cp_async.py b/tests/python/tir-base/test_tir_ptx_cp_async.py index f3255bd257c6..3d61a90106b5 100644 --- a/tests/python/tir-base/test_tir_ptx_cp_async.py +++ b/tests/python/tir-base/test_tir_ptx_cp_async.py @@ -94,7 +94,7 @@ def ptx_cp_async_barrier( B[tx, i] = A_shared[tx, i] -@tvm.testing.requires_cuda_compute_version(8) +@tvm.testing.requires_cuda_compute_version(9) def test_ptx_cp_async_barrier(): f = ptx_cp_async_barrier diff --git a/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py b/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py index aa4f5138a17f..0855afcfd64a 100644 --- a/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py +++ b/tests/python/tir-transform/test_tir_transform_inject_ptx_async_copy.py @@ -214,7 +214,7 @@ def ptx_global_to_shared_copy_fp32x1_barrier( B[tx, i] = A_shared[tx, i] -@tvm.testing.requires_cuda +@tvm.testing.requires_cuda_compute_version(9) def test_inject_async_copy_barrier(): dtype = "float32" vec_size = 1