[SYCLomatic][PTX] Enable to migrate "cp.async.commit_group", "cp.async.wait_group" and "cp.async.wait_all"

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -2712,6 +2712,27 @@ class SYCLGen : public SYCLGenBase {
     insertHeader(HeaderType::HT_DPCT_Atomic);
     return SYCLGenSuccess();
   }
+
+  bool handle_cp(const InlineAsmInstruction *Inst) override {
+    if (Inst->getNumInputOperands() == 0 && Inst->hasAttr(InstAttr::async) &&
+        (Inst->hasAttr(InstAttr::commit_group) ||
+         Inst->hasAttr(InstAttr::wait_group) ||
+         Inst->hasAttr(InstAttr::wait_all))) {
+      auto CommonStr = llvm::Twine("")
+                           .concat("\"")
+                           .concat(GAS->getAsmString()->getString())
+                           .concat("\"")
+                           .str();
+
+      report(Diagnostics::FUNC_CALL_REMOVED, true, CommonStr,
+             "SYCL currently has a functionality gap for asynchronous copy "
+             "operations. Synchronous copy operations are used on the SYCL "
+             "side and may cause performance issues. You may need to adjust "
+             "the code.");
+      return SYCLGenSuccess();
+    }
+    return SYCLGenError();
+  }
 };
 
 /// Clean the special character in identifier.

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

@@ -397,6 +397,10 @@ MODIFIER(clamp, ".clamp")
 MODIFIER(wrap, ".wrap")
 MODIFIER(wide, ".wide")
 MODIFIER(sync, ".sync")
+MODIFIER(async, ".async")
+MODIFIER(commit_group, ".commit_group")
+MODIFIER(wait_group, ".wait_group")
+MODIFIER(wait_all, ".wait_all")
 MODIFIER(warp, ".warp")
 MODIFIER(up, ".up")
 MODIFIER(down, ".down")

clang/lib/DPCT/SrcAPI/APINames_ASM.inc

@@ -56,7 +56,7 @@ ENTRY("clz", "clz", true, NO_FLAG, P1, "Successful")
 ENTRY("cnot", "cnot", true, NO_FLAG, P1, "Successful")
 ENTRY("copysign", "copysign", true, NO_FLAG, P1, "Successful")
 ENTRY("cos", "cos", true, NO_FLAG, P1, "Successful")
-ENTRY("cp", "cp", false, NO_FLAG, P1, "Comment")
+ENTRY("cp", "cp", true, NO_FLAG, P1, "Partial")
 ENTRY("createpolicy", "createpolicy", false, NO_FLAG, P1, "Comment")
 ENTRY("cvt", "cvt", true, NO_FLAG, P1, "Partial")
 ENTRY("cvta", "cvta", false, NO_FLAG, P1, "Comment")

clang/test/dpct/asm/cp.cu

@@ -0,0 +1,42 @@
+// 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/cp %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck %s --match-full-lines --input-file %T/cp/cp.dp.cpp
+// RUN: %if build_lit %{icpx -c -fsycl %T/cp/cp.dp.cpp -o %T/cp/cp.dp.o %}
+
+// clang-format off
+#include <cuda_runtime.h>
+#include <cstdint>
+
+// CHECK:inline void cp_async_commit_group() {
+// CHECK-NEXT:  /*
+// CHECK-NEXT:  DPCT1026:{{[0-9]+}}: The call to "cp.async.commit_group;" was removed because SYCL currently has a functionality gap for asynchronous copy operations. Synchronous copy operations are used on the SYCL side and may cause performance issues. You may need to adjust the code.
+// CHECK-NEXT:  */
+// CHECK-EMPTY:
+// CHECK-NEXT:}
+__device__ inline void cp_async_commit_group() {
+asm volatile("cp.async.commit_group;" ::);
+}
+
+
+// CHECK:inline void cp_async_wait_group() {
+// CHECK-NEXT:  /*
+// CHECK-NEXT:  DPCT1026:{{[0-9]+}}: The call to "cp.async.wait_group 0;" was removed because SYCL currently has a functionality gap for asynchronous copy operations. Synchronous copy operations are used on the SYCL side and may cause performance issues. You may need to adjust the code.
+// CHECK-NEXT:  */
+// CHECK-EMPTY:
+// CHECK-NEXT:}
+__device__ inline void cp_async_wait_group() {
+asm volatile("cp.async.wait_group 0;");
+}
+
+// CHECK:inline void cp_async_wait_all() {
+// CHECK-NEXT:  /*
+// CHECK-NEXT:  DPCT1026:{{[0-9]+}}: The call to "cp.async.wait_all;" was removed because SYCL currently has a functionality gap for asynchronous copy operations. Synchronous copy operations are used on the SYCL side and may cause performance issues. You may need to adjust the code.
+// CHECK-NEXT:  */
+// CHECK-EMPTY:
+// CHECK-NEXT:}
+__device__ inline void cp_async_wait_all() {
+asm volatile("cp.async.wait_all;");
+}
+
+// clang-format on