Fixed illegal PTX syntax (#1225)

NVIDIA · Dec 1, 2023 · 4a1709e · 4a1709e
1 parent bef1fbc
commit 4a1709e
Showing 1 changed file with 2 additions and 3 deletions.
diff --git a/include/cutlass/arch/barrier.h b/include/cutlass/arch/barrier.h
@@ -267,13 +267,12 @@ struct ClusterBarrier {
   static void arrive(ValueType const* smem_ptr) {
 #if CUDA_BARRIER_ENABLED
     uint32_t smem_addr = cute::cast_smem_ptr_to_uint(smem_ptr);
-    uint64_t state = 0;
     asm volatile(
         "{\n\t"
-        "mbarrier.arrive.shared.b64 %1, [%0];\n\t"
+        "mbarrier.arrive.shared.b64 _, [%0];\n\t"
         "}"
         :
-        : "r"(smem_addr), "l"(state));
+        : "r"(smem_addr));
 #elif defined(__CUDA_ARCH__)
     asm volatile ("brkpt;\n" ::);
 #endif