From 52fb43f30fd43008ce29c21a677299582409c299 Mon Sep 17 00:00:00 2001 From: Kevin Tong Date: Wed, 10 Jul 2024 11:35:26 -0400 Subject: [PATCH] fix mbarrier invalidate (#1494) --- include/cutlass/arch/barrier.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/cutlass/arch/barrier.h b/include/cutlass/arch/barrier.h index 402ef50cbb..b08ee546d3 100644 --- a/include/cutlass/arch/barrier.h +++ b/include/cutlass/arch/barrier.h @@ -360,7 +360,7 @@ struct ClusterBarrier { uint32_t smem_addr = cute::cast_smem_ptr_to_uint(smem_ptr); asm volatile( "{\n\t" - "mbarrier.ival.shared::cta.b64 [%0]; \n\t" + "mbarrier.inval.shared::cta.b64 [%0]; \n\t" "}" : : "r"(smem_addr));