[SYCLomatic][ASM] Support migration of PTX instruction "membar" with level on .cta .gl and sys (#2476)

Signed-off-by: chenwei.sun <[email protected]>

Author: tomflinda

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -1227,7 +1227,6 @@ class SYCLGen : public SYCLGenBase {
     if (Inst->hasAttr(InstAttr::sc) && Inst->hasAttr(InstAttr::cta)) {
       OS() << MapNames::getClNamespace() << "memory_order::seq_cst,"
            << MapNames::getClNamespace() << "memory_scope::work_group";
-
     } else if (Inst->hasAttr(InstAttr::sc) && Inst->hasAttr(InstAttr::gpu)) {
       OS() << MapNames::getClNamespace() << "memory_order::seq_cst,"
            << MapNames::getClNamespace() << "memory_scope::device";
@@ -1248,7 +1247,7 @@ class SYCLGen : public SYCLGenBase {
       OS() << MapNames::getClNamespace() << "memory_order::acq_rel,"
            << MapNames::getClNamespace() << "memory_scope::system";
     } else {
-      SYCLGenError();
+      return SYCLGenError();
     }
 
     OS() << ')';
@@ -1260,6 +1259,28 @@ class SYCLGen : public SYCLGenBase {
     return HandleAddSub(Inst);
   }
 
+  bool handle_membar(const InlineAsmInstruction *Inst) override {
+    if (Inst->getNumInputOperands() != 0)
+      return SYCLGenError();
+
+    OS() << MapNames::getClNamespace() << "atomic_fence("
+         << MapNames::getClNamespace() << "memory_order::seq_cst,";
+
+    if (Inst->hasAttr(InstAttr::cta)) {
+      OS() << MapNames::getClNamespace() << "memory_scope::work_group";
+    } else if (Inst->hasAttr(InstAttr::gl)) {
+      OS() << MapNames::getClNamespace() << "memory_scope::device";
+    } else if (Inst->hasAttr(InstAttr::sys)) {
+      OS() << MapNames::getClNamespace() << "memory_scope::system";
+    } else {
+      return SYCLGenError();
+    }
+
+    OS() << ')';
+    endstmt();
+    return SYCLGenSuccess();
+  }
+
   StringRef GetWiderTypeAsString(const InlineAsmBuiltinType *Type) const {
     switch (Type->getKind()) {
     case InlineAsmBuiltinType::s16:

clang/lib/DPCT/RulesAsm/Parser/AsmTokenKinds.def

@@ -403,6 +403,7 @@ MODIFIER(down, ".down")
 MODIFIER(idx, ".idx")
 MODIFIER(bfly, ".bfly")
 MODIFIER(sc, ".sc")
+MODIFIER(gl, ".gl")
 
 #undef LINKAGE
 #undef TARGET

clang/lib/DPCT/SrcAPI/APINames_ASM.inc

@@ -67,8 +67,7 @@ ENTRY("dp4a", "dp4a", true, NO_FLAG, P1, "Successful")
 ENTRY("elect", "elect", false, NO_FLAG, P1, "Comment")
 ENTRY("ex2", "ex2", true, NO_FLAG, P1, "Successful")
 ENTRY("exit", "exit", false, NO_FLAG, P1, "Comment")
-ENTRY("fence", "true", false, NO_FLAG, P1,
-      "Only support thread fence on scope '.cta', '.gpu' and '.sys'")
+ENTRY("fence", "true", false, NO_FLAG, P1, "Only support thread fence on scope '.cta', '.gpu' and '.sys'")
 ENTRY("fma", "fma", true, NO_FLAG, P1, "Partial")
 ENTRY("fns", "fns", false, NO_FLAG, P1, "Comment")
 ENTRY("getctarank", "getctarank", false, NO_FLAG, P1, "Comment")
@@ -87,7 +86,7 @@ ENTRY("mapa", "mapa", false, NO_FLAG, P1, "Comment")
 ENTRY("match", "match", false, NO_FLAG, P1, "Comment")
 ENTRY("max", "max", true, NO_FLAG, P1, "Successful")
 ENTRY("mbarrier", "mbarrier", false, NO_FLAG, P1, "Comment")
-ENTRY("membar", "membar", false, NO_FLAG, P1, "Comment")
+ENTRY("membar", "membar", true, NO_FLAG, P1, "only support level on .cta .gl and sys)
 ENTRY("min", "min", true, NO_FLAG, P1, "Successful")
 ENTRY("mma", "mma", false, NO_FLAG, P1, "Comment")
 ENTRY("mov", "mov", true, NO_FLAG, P1, "Successful")

clang/test/dpct/asm/membar_fence.cu

@@ -28,4 +28,16 @@ __global__ void fence() {
   asm volatile("fence.acq_rel.sys; " : : : "memory");
 }
 
+__global__ void membar() {
+
+  // CHECK: sycl::atomic_fence(sycl::memory_order::seq_cst,sycl::memory_scope::work_group);
+  asm volatile("membar.cta;":::"memory"); 
+
+  // CHECK: sycl::atomic_fence(sycl::memory_order::seq_cst,sycl::memory_scope::device);
+  asm volatile("membar.gl;":::"memory"); 
+
+  // CHECK: sycl::atomic_fence(sycl::memory_order::seq_cst,sycl::memory_scope::system);
+  asm volatile("membar.sys;":::"memory");
+}
+
 // clang-format off