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

Signed-off-by: chenwei.sun <chenwei.sun@intel.com>

Author: tomflinda

clang/lib/DPCT/Diagnostics/Diagnostics.inc

@@ -298,8 +298,8 @@ DEF_WARNING(JOINT_MATRIX_SHAPE, 1135, HIGH_LEVEL, "Please check if joint_matrix
 DEF_COMMENT(JOINT_MATRIX_SHAPE, 1135, HIGH_LEVEL, "Please check if joint_matrix implementations support the combination of data type and matrix shape type in the target hardware.")
 DEF_WARNING(UNSUPPORTED_EXTMEM_WIN_HANDLE, 1136, HIGH_LEVEL, "SYCL Bindless Images extension only supports importing external resource memory using NT handle on Windows. If assert(%0.get_win32_handle()) fails, you may need to adjust the code to use (%0.get_win32_handle()).")
 DEF_COMMENT(UNSUPPORTED_EXTMEM_WIN_HANDLE, 1136, HIGH_LEVEL, "SYCL Bindless Images extension only supports importing external resource memory using NT handle on Windows. If assert({0}.get_win32_handle()) fails, you may need to adjust the code to use ({0}.get_win32_handle()).")
-DEF_WARNING(ASYNC_COPY_DEVICE_WARN, 1137, LOW_LEVEL, "ASM instruction \"cp.async\" is asynchronous copy, current it is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.")
-DEF_COMMENT(ASYNC_COPY_DEVICE_WARN, 1137, LOW_LEVEL, "ASM instruction \"cp.async\" is asynchronous copy, current it is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.")
+DEF_WARNING(ASYNC_COPY_DEVICE_WARN, 1137, LOW_LEVEL, "ASM instruction \"%0\" is asynchronous copy, current it is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.")
+DEF_COMMENT(ASYNC_COPY_DEVICE_WARN, 1137, LOW_LEVEL, "ASM instruction \"{0}\" is asynchronous copy, current it is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.")
 
 // clang-format on
 

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -37,7 +37,7 @@ using namespace clang::dpct;
 namespace {
 
 inline bool SYCLGenError() { return true; }
-inline bool SYCLGenSuccess() {return false; }
+inline bool SYCLGenSuccess() { return false; }
 
 /// This is used to handle all the AST nodes (except specific instructions, Eg.
 /// mov/setp), and generate functionally equivalent SYCL code.
@@ -2177,8 +2177,8 @@ class SYCLGen : public SYCLGenBase {
     if (emitStmt(Inst->getOutputOperand()))
       return SYCLGenError();
     OS() << " = ";
-    OS() << MapNames::getDpctNamespace() << "bfe_safe<" << TypeStr << ">(" << Op[0]
-         << ", " << Op[1] << ", " << Op[2] << ')';
+    OS() << MapNames::getDpctNamespace() << "bfe_safe<" << TypeStr << ">("
+         << Op[0] << ", " << Op[1] << ", " << Op[2] << ')';
     endstmt();
     insertHeader(HeaderType::HT_DPCT_Math);
     return SYCLGenSuccess();
@@ -2380,8 +2380,8 @@ class SYCLGen : public SYCLGenBase {
     if (DpctGlobalInfo::useIntelDeviceMath() && !RD.empty()) {
       insertHeader(HeaderType::HT_SYCL_Math);
       OS() << MapNames::getClNamespace() << "ext::intel::math::"
-           << (T->getKind() == InlineAsmBuiltinType::f32 ? 'f' : 'd')
-           << "rcp_" << RD << '(' << Op[0] << ')';
+           << (T->getKind() == InlineAsmBuiltinType::f32 ? 'f' : 'd') << "rcp_"
+           << RD << '(' << Op[0] << ')';
     } else {
       OS() << "1 / " << Op[0];
     }
@@ -2773,9 +2773,7 @@ class SYCLGen : public SYCLGenBase {
     return SYCLGenSuccess();
   }
 
-  bool handle_cp(const InlineAsmInstruction *Inst) override {
-    if (Inst->getNumInputOperands() != 3 || Inst->getNumTypes() != 1)
-      return SYCLGenError();
+  bool HandleCopyOperation(const InlineAsmInstruction *Inst) {
 
     llvm::SaveAndRestore<const InlineAsmInstruction *> Store(CurrInst);
     CurrInst = Inst;
@@ -2809,9 +2807,40 @@ class SYCLGen : public SYCLGenBase {
     OS() << CommonBody("3");
     endstmt();
 
-    report(Diagnostics::ASYNC_COPY_DEVICE_WARN, true);
+    auto OpStr =
+        llvm::Twine(Inst->getOpcodeID()->getName()).concat(".async").str();
+    report(Diagnostics::ASYNC_COPY_DEVICE_WARN, true, OpStr);
+    return SYCLGenSuccess();
+  }
+
+  bool HandleCopyWait(const InlineAsmInstruction *Inst) {
+    auto CommonStr = llvm::Twine("")
+                         .concat("\"")
+                         .concat(GAS->getAsmString()->getString())
+                         .concat("\"")
+                         .str();
+
+    report(
+        Diagnostics::FUNC_CALL_REMOVED, true, CommonStr,
+        "current \"cp.async\" is migrated to synchronous copy operation. You "
+        "may need to adjust the code to tune the performance.");
     return SYCLGenSuccess();
   }
+
+  bool handle_cp(const InlineAsmInstruction *Inst) override {
+    if (Inst->getNumInputOperands() == 3 && Inst->getNumTypes() == 1 &&
+        Inst->hasAttr(InstAttr::async))
+      return HandleCopyOperation(Inst);
+
+    if (Inst->getNumInputOperands() == 1 && Inst->hasAttr(InstAttr::async) &&
+        (Inst->hasAttr(InstAttr::commit_group) ||
+         Inst->hasAttr(InstAttr::wait_group) ||
+         Inst->hasAttr(InstAttr::wait_all))) {
+      return HandleCopyWait(Inst);
+    }
+
+    return SYCLGenError();
+  }
 };
 
 /// Clean the special character in identifier.

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

@@ -400,6 +400,9 @@ MODIFIER(wide, ".wide")
 MODIFIER(sync, ".sync")
 MODIFIER(async, ".async")
 MODIFIER(cg, ".cg")
+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/test/dpct/asm/cp.cu

@@ -6,7 +6,6 @@
 
 // clang-format off
 #include <cstdint>
-#include <cstdint>
 #include <cuda_runtime.h>
 
 // CHECK: inline void cp_async4(void *smem_ptr, const void *glob_ptr) {
@@ -67,4 +66,34 @@ __device__ inline void cp_async4_pred(void *smem_ptr, const void *glob_ptr,
                "r"(smem), "l"(glob_ptr), "n"(BYTES));
 }
 
+// CHECK:inline void cp_async_commit_group() {
+// CHECK-NEXT:  /*
+// CHECK-NEXT:  DPCT1026:{{[0-9]+}}: The call to "cp.async.commit_group;" was removed because current "cp.async" is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.
+// 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 current "cp.async" is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.
+// 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 current "cp.async" is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.
+// CHECK-NEXT:  */
+// CHECK-EMPTY:
+// CHECK-NEXT:}
+__device__ inline void cp_async_wait_all() {
+asm volatile("cp.async.wait_all;");
+}
+
 // clang-format on