[SYCLomatic][ASM] Support migration of PTX instruction "fence" with thread level on sope '.cta', '.gpu' and '.sys' (#2471)

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

Author: tomflinda

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -1219,6 +1219,43 @@ class SYCLGen : public SYCLGenBase {
     return SYCLGenError();
   }
 
+  bool handle_fence(const InlineAsmInstruction *Inst) override {
+    if (Inst->getNumInputOperands() != 0)
+      return SYCLGenError();
+
+    OS() << MapNames::getClNamespace() << "atomic_fence(";
+    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";
+    } else if (Inst->hasAttr(InstAttr::sc) && Inst->hasAttr(InstAttr::sys)) {
+      OS() << MapNames::getClNamespace() << "memory_order::seq_cst,"
+           << MapNames::getClNamespace() << "memory_scope::system";
+    } else if (Inst->hasAttr(InstAttr::acq_rel) &&
+               Inst->hasAttr(InstAttr::cta)) {
+      OS() << MapNames::getClNamespace() << "memory_order::acq_rel,"
+           << MapNames::getClNamespace() << "memory_scope::work_group";
+
+    } else if (Inst->hasAttr(InstAttr::acq_rel) &&
+               Inst->hasAttr(InstAttr::gpu)) {
+      OS() << MapNames::getClNamespace() << "memory_order::acq_rel,"
+           << MapNames::getClNamespace() << "memory_scope::device";
+    } else if (Inst->hasAttr(InstAttr::acq_rel) &&
+               Inst->hasAttr(InstAttr::sys)) {
+      OS() << MapNames::getClNamespace() << "memory_order::acq_rel,"
+           << MapNames::getClNamespace() << "memory_scope::system";
+    } else {
+      SYCLGenError();
+    }
+
+    OS() << ')';
+    endstmt();
+    return SYCLGenSuccess();
+  }
+
   bool handle_sub(const InlineAsmInstruction *Inst) override {
     return HandleAddSub(Inst);
   }

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

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

clang/lib/DPCT/SrcAPI/APINames_ASM.inc

@@ -67,7 +67,8 @@ 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", "fence", false, NO_FLAG, P1, "Comment")
+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")

clang/test/dpct/asm/membar_fence.cu

@@ -0,0 +1,31 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2
+// RUN: dpct --format-range=none -out-root %T/membar_fence %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck %s --match-full-lines --input-file %T/membar_fence/membar_fence.dp.cpp
+// RUN: %if build_lit %{icpx -c -fsycl %T/membar_fence/membar_fence.dp.cpp -o %T/membar_fence/membar_fence.dp.o %}
+
+// clang-format off
+#include <cuda_runtime.h>
+#include <cstdint>
+
+__global__ void fence() {
+  // CHECK: sycl::atomic_fence(sycl::memory_order::seq_cst,sycl::memory_scope::work_group);
+  asm volatile("fence.sc.cta; " : : : "memory");
+
+  // CHECK: sycl::atomic_fence(sycl::memory_order::seq_cst,sycl::memory_scope::device);
+  asm volatile("fence.sc.gpu; " : : : "memory");
+
+  // CHECK: sycl::atomic_fence(sycl::memory_order::seq_cst,sycl::memory_scope::system);
+  asm volatile("fence.sc.sys; " : : : "memory");
+
+  // CHECK: sycl::atomic_fence(sycl::memory_order::acq_rel,sycl::memory_scope::work_group);
+  asm volatile("fence.acq_rel.cta; " : : : "memory");
+
+  // CHECK: sycl::atomic_fence(sycl::memory_order::acq_rel,sycl::memory_scope::device);
+  asm volatile("fence.acq_rel.gpu; " : : : "memory");
+
+  // CHECK: sycl::atomic_fence(sycl::memory_order::acq_rel,sycl::memory_scope::system);
+  asm volatile("fence.acq_rel.sys; " : : : "memory");
+}
+
+// clang-format off