From b7508e337938137a699e486d8997646980acfc58 Mon Sep 17 00:00:00 2001 From: Christian Sigg Date: Thu, 14 Dec 2023 17:16:15 +0100 Subject: [PATCH] Fix inline ptx escaping for predicates. (#1264) * Fix inline ptx escaping for predicates. Prevents `error: invalid % escape in inline assembly string` when compiling with clang. * More double-quoting. --- include/cute/arch/cluster_sm90.hpp | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/include/cute/arch/cluster_sm90.hpp b/include/cute/arch/cluster_sm90.hpp index b034b2cd..40b9e2c2 100644 --- a/include/cute/arch/cluster_sm90.hpp +++ b/include/cute/arch/cluster_sm90.hpp @@ -181,11 +181,11 @@ CUTE_HOST_DEVICE uint32_t elect_one_sync() uint32_t laneid = 0; asm volatile( "{\n" - ".reg .b32 %rx;\n" - ".reg .pred %px;\n" - " elect.sync %rx|%px, %2;\n" - "@%px mov.s32 %1, 1;\n" - " mov.s32 %0, %rx;\n" + ".reg .b32 %%rx;\n" + ".reg .pred %%px;\n" + " elect.sync %%rx|%%px, %2;\n" + "@%%px mov.s32 %1, 1;\n" + " mov.s32 %0, %%rx;\n" "}\n" : "+r"(laneid), "+r"(pred) : "r"(0xFFFFFFFF)); @@ -211,11 +211,11 @@ elect_one_leader_sync() uint32_t laneid = 0; asm volatile( "{\n" - ".reg .b32 %rx;\n" - ".reg .pred %px;\n" - " elect.sync %rx|%px, %2;\n" - "@%px mov.s32 %1, 1;\n" - " mov.s32 %0, %rx;\n" + ".reg .b32 %%rx;\n" + ".reg .pred %%px;\n" + " elect.sync %%rx|%%px, %2;\n" + "@%%px mov.s32 %1, 1;\n" + " mov.s32 %0, %%rx;\n" "}\n" : "+r"(laneid), "+r"(pred) : "r"(0xFFFFFFFF));