[SYCLomatic][ASM] Remove debug purpose instruction "brkpt" (#2702)

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

Author: tomflinda

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -2841,6 +2841,23 @@ class SYCLGen : public SYCLGenBase {
 
     return SYCLGenError();
   }
+
+  bool handle_brkpt(const InlineAsmInstruction *Inst) override {
+
+    if (Inst->getNumInputOperands() != 0)
+      return SYCLGenError();
+
+    auto CommonStr = llvm::Twine("")
+                         .concat("\"")
+                         .concat(GAS->getAsmString()->getString())
+                         .concat("\"")
+                         .str();
+
+    report(Diagnostics::FUNC_CALL_REMOVED, true, CommonStr,
+           "this instruction is typically used for debugging. You may need to "
+           "rewrite the code.");
+    return SYCLGenSuccess();
+  }
 };
 
 /// Clean the special character in identifier.

clang/lib/DPCT/SrcAPI/APINames_ASM.inc

@@ -49,7 +49,7 @@ ENTRY("bfind", "bfind", false, NO_FLAG, P1, "Comment")
 ENTRY("bmsk", "bmsk", false, NO_FLAG, P1, "Comment")
 ENTRY("bra", "bra", false, NO_FLAG, P1, "Comment")
 ENTRY("brev", "brev", true, NO_FLAG, P1, "Successful")
-ENTRY("brkpt", "brkpt", false, NO_FLAG, P1, "Comment")
+ENTRY("brkpt", "brkpt", true, NO_FLAG, P1, "Comment")
 ENTRY("brx", "brx", false, NO_FLAG, P1, "Comment")
 ENTRY("call", "call", false, NO_FLAG, P1, "Comment")
 ENTRY("clz", "clz", true, NO_FLAG, P1, "Successful")

clang/test/dpct/asm/brkpt.cu

@@ -0,0 +1,21 @@
+// 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/brkpt %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck %s --match-full-lines --input-file %T/brkpt/brkpt.dp.cpp
+// RUN: %if build_lit %{icpx -c -fsycl %T/brkpt/brkpt.dp.cpp -o %T/brkpt/brkpt.dp.o %}
+
+// clang-format off
+#include <cstdint>
+#include <cuda_runtime.h>
+
+// CHECK:inline void cp_async_wait_all() {
+// CHECK-NEXT: /*
+// CHECK-NEXT: DPCT1026:{{[0-9]+}}: The call to "brkpt;" was removed because this instruction is typically used for debugging. You may need to rewrite the code.
+// CHECK-NEXT: */
+// CHECK-NEXT:   // PTX breakpoint instruction
+// CHECK-NEXT:}
+__device__ inline void cp_async_wait_all() {
+ asm("brkpt;");  // PTX breakpoint instruction
+}
+
+// clang-format on