[SYCLomatic][ASM] Support migration of PTX instruction prmt.b32 (#2708)

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

Author: tomflinda

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -2858,6 +2858,45 @@ class SYCLGen : public SYCLGenBase {
            "rewrite the code.");
     return SYCLGenSuccess();
   }
+
+  bool handle_prmt(const InlineAsmInstruction *Inst) override {
+    if (Inst->getNumInputOperands() != 3 || Inst->getNumTypes() != 1)
+      return SYCLGenError();
+
+    if (emitStmt(Inst->getOutputOperand()))
+      return SYCLGenError();
+    OS() << " = " << MapNames::getDpctNamespace()
+         << "byte_level_permute_custom(";
+
+    llvm::SaveAndRestore<const InlineAsmInstruction *> Save(CurrInst);
+    CurrInst = Inst;
+    std::string Op[3];
+    if (tryEmitAllInputOperands(Op, Inst))
+      return SYCLGenError();
+
+    OS() << Op[0] << ", ";
+    OS() << Op[1] << ", ";
+    OS() << Op[2] << ", ";
+    if (Inst->hasAttr(InstAttr::f4e)) {
+      OS() << "1";
+    } else if (Inst->hasAttr(InstAttr::b4e)) {
+      OS() << "2";
+    } else if (Inst->hasAttr(InstAttr::rc8)) {
+      OS() << "3";
+    } else if (Inst->hasAttr(InstAttr::ecl)) {
+      OS() << "4";
+    } else if (Inst->hasAttr(InstAttr::ecr)) {
+      OS() << "5";
+    } else if (Inst->hasAttr(InstAttr::rc16)) {
+      OS() << "6";
+    } else {
+      OS() << "0";
+    }
+    OS() << ")";
+
+    endstmt();
+    return SYCLGenSuccess();
+  }
 };
 
 /// Clean the special character in identifier.

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

@@ -412,7 +412,12 @@ MODIFIER(sc, ".sc")
 MODIFIER(gl, ".gl")
 MODIFIER(L1, ".L1")
 MODIFIER(L2, ".L2")
-
+MODIFIER(f4e, ".f4e")
+MODIFIER(b4e, ".b4e")
+MODIFIER(rc8, ".rc8")
+MODIFIER(ecl, ".ecl")
+MODIFIER(ecr, ".ecr")
+MODIFIER(rc16, ".rc16")
 
 #undef LINKAGE
 #undef TARGET

clang/lib/DPCT/SrcAPI/APINames_ASM.inc

@@ -101,7 +101,7 @@ ENTRY("or", "or", true, NO_FLAG, P1, "Successful")
 ENTRY("pmevent", "pmevent", false, NO_FLAG, P1, "Comment")
 ENTRY("popc", "popc", true, NO_FLAG, P1, "Successful")
 ENTRY("prefetch", "prefetch", true, NO_FLAG, P1, "Partial")
-ENTRY("prmt", "prmt", false, NO_FLAG, P1, "Comment")
+ENTRY("prmt", "prmt", true, NO_FLAG, P1, "Successful")
 ENTRY("rcp", "rcp", true, NO_FLAG, P1, "Successful")
 ENTRY("red", "red", true, NO_FLAG, P1, "Partial")
 ENTRY("redux", "redux", false, NO_FLAG, P1, "Comment")

clang/runtime/dpct-rt/include/dpct/util.hpp

@@ -187,6 +187,48 @@ inline unsigned int byte_level_permute(unsigned int a, unsigned int b,
   return ret;
 }
 
+/// \param [in] low32 The 4 bytes to construct the 8 bytes value as low 32 bits.
+/// \param [in] high32 The 4 bytes to construct the 8 bytes value as high 32
+/// bits.
+/// \param [in] sel The selector value. It is used to generate selectors which are
+/// used to fetch byte value from \p low32 and \p high32 to construct result
+/// value.
+/// \param [in] mode The mode of permutation, together with \p sel, it
+/// further defines the behavior of data selection. \p mode and \p sel define 4
+/// selectors (s[i], i=0, 1, 2, 3), which are used as index to fetch 4 bytes from
+/// the 8 bytes value constructed by \p low32 and \p high32; the byte selected by s[i]
+/// is used to fill the i-th byte of the result.
+/// The available mode values are:
+/// mode value 0: s[i] = sel[i * 4 + 3 : i * 4]
+/// mode value 1: s[i] = sel[1 : 0] + i
+/// mode value 2: s[i] = (sel[1 : 0] + 7) % 8
+/// mode value 3: s[i] = sel[1 : 0]
+/// mode value 4: s[i] = max(sel[1 : 0], i)
+/// mode value 5: s[i] = min(sel[1 : 0], i)
+/// mode value 6: s[0] = sel[0 : 0] * 2
+///               s[1] = sel[0 : 0] * 2 + 1
+///               s[2] = s[0]
+///               s[3] = s[1]
+/// other value: illegal, undefined behavior, return 0.
+inline uint32_t byte_level_permute_custom(uint32_t low32, uint32_t high32,
+                                          uint32_t sel, int mode = 0) {
+  constexpr uint16_t lookup[6][4] = {
+      {0x3210, 0x4321, 0x5432, 0x6543}, // Forward 4-byte extract
+      {0x5670, 0x6701, 0x7012, 0x0123}, // Backward 4-byte extract
+      {0x0000, 0x1111, 0x2222, 0x3333}, // Replicate 8-bit values
+      {0x3210, 0x3211, 0x3222, 0x3333}, // Edge clamp left
+      {0x0000, 0x1110, 0x2210, 0x3210}, // Edge clamp right
+      {0x1010, 0x3232, 0x1010, 0x3232}  // Replicate 16-bit values
+  };
+
+  if (mode >= 1 && mode <= 6) {
+    return byte_level_permute(low32, high32, lookup[mode - 1][sel & 0x3]);
+  } else if (!mode) {
+    return byte_level_permute(low32, high32, sel);
+  }
+  return 0;
+}
+
 /// Find position of first least significant set bit in an integer.
 /// ffs(0) returns 0.
 ///

clang/test/dpct/asm/prmt.cu

@@ -0,0 +1,38 @@
+// 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/prmt %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck %s --match-full-lines --input-file %T/prmt/prmt.dp.cpp
+// RUN: %if build_lit %{icpx -c -fsycl %T/prmt/prmt.dp.cpp -o %T/prmt/prmt.dp.o %}
+
+// clang-format off
+#include <cstdint>
+#include <cuda_runtime.h>
+
+__global__ void testKernel1(uint32_t *d_result, uint32_t a) {
+  static constexpr uint32_t sel = 0x3210;
+  static constexpr uint32_t b = 0;
+  uint32_t d;
+
+  // CHECK: d = dpct::byte_level_permute_custom(a, b, sel, 0);
+  asm volatile("prmt.b32 %0, %1, %2, %3;\n" : "=r"(d) : "r"(a), "n"(b), "n"(sel));
+
+  // CHECK: d = dpct::byte_level_permute_custom(a, b, sel, 1);
+  asm volatile("prmt.b32.f4e %0, %1, %2, %3;\n" : "=r"(d) : "r"(a), "n"(b), "n"(sel));
+
+  // CHECK: d = dpct::byte_level_permute_custom(a, b, sel, 2);
+  asm volatile("prmt.b32.b4e %0, %1, %2, %3;\n" : "=r"(d) : "r"(a), "n"(b), "n"(sel));
+
+  // CHECK: d = dpct::byte_level_permute_custom(a, b, sel, 3);
+  asm volatile("prmt.b32.rc8 %0, %1, %2, %3;\n" : "=r"(d) : "r"(a), "n"(b), "n"(sel));
+
+  // CHECK: d = dpct::byte_level_permute_custom(a, b, sel, 4);
+  asm volatile("prmt.b32.ecl  %0, %1, %2, %3;\n" : "=r"(d) : "r"(a), "n"(b), "n"(sel));
+
+  // CHECK: d = dpct::byte_level_permute_custom(a, b, sel, 5);
+  asm volatile("prmt.b32.ecr %0, %1, %2, %3;\n" : "=r"(d) : "r"(a), "n"(b), "n"(sel));
+
+  // CHECK: d = dpct::byte_level_permute_custom(a, b, sel, 6);
+  asm volatile("prmt.b32.rc16 %0, %1, %2, %3;\n" : "=r"(d) : "r"(a), "n"(b), "n"(sel));
+}
+
+// clang-format on

docs/dev_guide/api-mapping-status/ASM_API_migration_status.csv

@@ -67,7 +67,7 @@ or,YES,
 pmevent,NO,
 popc,YES,
 prefetch,YES, Partial
-prmt,NO,
+prmt,YES,
 rcp,YES,
 red,NO,
 redux,NO,