[SYCLomatic][PTX] Refine migration of asm PTX instrution "lop3.b32"

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

Author: tomflinda

clang/lib/DPCT/RulesAsm/AsmMigration.cpp

@@ -942,82 +942,14 @@ class SYCLGen : public SYCLGenBase {
       return SYCLGenError();
     OS() << " = ";
 
-    std::string Op[3];
-    for (auto Idx : llvm::seq(0, 3)) {
+    std::string Op[4];
+    for (auto Idx : llvm::seq(0, 4)) {
       if (tryEmitStmt(Op[Idx], I->getInputOperand(Idx)))
         return SYCLGenError();
     }
 
-    if (!isa<InlineAsmIntegerLiteral>(I->getInputOperand(3)))
-      return SYCLGenError();
-    unsigned Imm = dyn_cast<InlineAsmIntegerLiteral>(I->getInputOperand(3))
-                       ->getValue()
-                       .getZExtValue();
-
-#define EMPTY nullptr
-#define EMPTY4 EMPTY, EMPTY, EMPTY, EMPTY
-#define EMPTY16 EMPTY4, EMPTY4, EMPTY4, EMPTY4
-    constexpr const char *FastMap[256] = {
-        /*0x00*/ "0",
-        // clang-format off
-      EMPTY16, EMPTY4, EMPTY4, EMPTY,
-      /*0x1a*/ "({0} & {1} | {2}) ^ {0}",
-      EMPTY, EMPTY, EMPTY,
-      /*0x1e*/ "{0} ^ ({1} | {2})",
-      EMPTY4, EMPTY4, EMPTY4, EMPTY, EMPTY,
-      /*0x2d*/ "~{0} ^ (~{1} & {2})",
-      EMPTY16, EMPTY, EMPTY,
-      /*0x40*/ "{0} & {1} & ~{2}",
-      EMPTY16, EMPTY16, EMPTY16, EMPTY4, EMPTY, EMPTY, EMPTY,
-      /*0x78*/ "{0} ^ ({1} & {2})",
-      EMPTY4, EMPTY, EMPTY, EMPTY,
-      /*0x80*/ "{0} & {1} & {2}",
-      EMPTY16, EMPTY4, EMPTY,
-      /*0x96*/ "{0} ^ {1} ^ {2}",
-      EMPTY16, EMPTY4, EMPTY4, EMPTY4, EMPTY,
-      /*0xb4*/ "{0} ^ ({1} & ~{2})",
-      EMPTY, EMPTY, EMPTY,
-      /*0xb8*/ "({0} ^ ({1} & ({2} ^ {0})))",
-      EMPTY16, EMPTY4, EMPTY4, EMPTY,
-      /*0xd2*/ "{0} ^ (~{1} & {2})",
-      EMPTY16, EMPTY4, EMPTY,
-      /*0xe8*/ "(({0} & ({1} | {2})) | ({1} & {2}))",
-      EMPTY,
-      /*0xea*/ "({0} & {1}) | {2}",
-      EMPTY16, EMPTY, EMPTY, EMPTY,
-      /*0xfe*/ "{0} | {1} | {2}",
-      /*0xff*/ "uint32_t(-1)"};
-    // clang-format on
-
-#undef EMPTY16
-#undef EMPTY4
-#undef EMPTY
-    // clang-format off
-    constexpr const char *SlowMap[8] = {
-      /* 0x01*/ "(~{0} & ~{1} & ~{2})",
-      /* 0x02*/ "(~{0} & ~{1} & {2})",
-      /* 0x04*/ "(~{0} & {1} & ~{2})",
-      /* 0x08*/ "(~{0} & {1} & {2})",
-      /* 0x10*/ "({0} & ~{1} & ~{2})",
-      /* 0x20*/ "({0} & ~{1} & {2})",
-      /* 0x40*/ "({0} & {1} & ~{2})",
-      /* 0x80*/ "({0} & {1} & {2})"
-    };
-    // clang-format on
-
-    if (FastMap[Imm]) {
-      OS() << llvm::formatv(FastMap[Imm], Op[0], Op[1], Op[2]);
-    } else {
-      SmallVector<std::string, 8> Templates;
-      for (auto Bit : llvm::seq(0, 8)) {
-        if (Imm & (1U << Bit)) {
-          Templates.push_back(
-              llvm::formatv(SlowMap[Bit], Op[0], Op[1], Op[2]).str());
-        }
-      }
-
-      OS() << llvm::join(Templates, " | ");
-    }
+    OS() << MapNames::getDpctNamespace() << "lop3(" << Op[0] << ", " << Op[1]
+         << ", " << Op[2] << ", " << Op[3] << ")";
 
     endstmt();
     return SYCLGenSuccess();

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

@@ -1205,6 +1205,36 @@ template <typename Func, std::size_t N> struct nth_argument_type {
   using type = decltype(helper(std::declval<Func>()));
 };
 
+/// The functions performs bitwise logical operations on three input values of \p a, \p b and
+/// \p c based on the specified 8-bit truth table \p lut and return the result
+/// \param [in] a Input value
+/// \param [in] b Input value
+/// \param [in] c Input value
+/// \param [in] lut truth table for looking up
+/// \returns The result
+uint32_t lop3(uint32_t a, uint32_t b, uint32_t c, uint8_t lut) {
+  uint32_t result = 0;
+
+  // Iterate through all 32 bits
+  for (int i = 0; i < 32; i++) {
+    // Extract the i-th bit from each input
+    uint8_t a_bit_val = (a >> i) & 1;
+    uint8_t b_bit_val = (b >> i) & 1;
+    uint8_t c_bit_val = (c >> i) & 1;
+
+    // Compute the index for the truth table using the three bits
+    uint8_t index = a_bit_val | (b_bit_val << 1) | (c_bit_val << 2);
+
+    // Extract the corresponding bit from the mask
+    uint8_t output_bit = (lut >> index) & 1;
+
+    // Set the output bit in the result
+    result |= (output_bit << i);
+  }
+
+  return result;
+}
+
 #ifdef _WIN32
 #define DPCT_EXPORT __declspec(dllexport)
 #else

clang/test/dpct/asm/lop3.cu

@@ -7,32 +7,45 @@
 // a^b^c
 static __device__ __forceinline__ uint32_t LOP3LUT_XOR(uint32_t a, uint32_t b, uint32_t c) {
   uint32_t d1;
-  // CHECK: d1 = a ^ b ^ c;
+  // CHECK: d1 = dpct::lop3(a, b, c, 0x96);
   asm("lop3.b32 %0, %1, %2, %3, 0x96;" : "=r"(d1) : "r"(a), "r"(b), "r"(c));
   return d1;
 }
 
 // (a ^ (c & (b ^ a)))
 static __device__ __forceinline__ uint32_t LOP3LUT_XORAND(uint32_t a, uint32_t b, uint32_t c) {
   uint32_t d2;
-  // CHECK: d2 = (a ^ (c & (b ^ a)));
+  // CHECK: d2 = dpct::lop3(a, c, b, 0xb8);
   asm("lop3.b32 %0, %1, %3, %2, 0xb8;" : "=r"(d2) : "r"(a), "r"(b), "r"(c));
   return d2;
 }
 
 // ((a & (b | b)) | (b & b))
 static __device__ __forceinline__ uint32_t LOP3LUT_ANDOR(uint32_t a, uint32_t b) {
   uint32_t d3;
-  // CHECK: d3 = ((a & (b | b)) | (b & b));
+  // CHECK: d3 = dpct::lop3(a, b, b, 0xe8);
   asm("lop3.b32 %0, %1, %2, %2, 0xe8;" : "=r"(d3) : "r"(a), "r"(b));
   return d3;
 }
 
 #define B 3
 __device__  int hard(int a) {
   int d4;
-  // CHECK: d4 = (~(a + B) & B & ~3) | (~(a + B) & B & 3) | ((a + B) & ~B & ~3);
+  // CHECK: d4 = dpct::lop3((a + B), B, 3, 0x1C);
   asm("lop3.b32 %0, %1, %2, 3, 0x1C;" : "=r"(d4) : "r"(a + B), "r"(B));
   return d4;
 }
+
+// CHECK: template <int lut, typename T> inline T lop3(T a, T b, T c) {
+// CHECK-NEXT:  T res;
+// CHECK-NEXT:  res = dpct::lop3(a, b, c, lut);
+// CHECK-NEXT:  return res;
+// CHECK-NEXT:}
+template <int lut, typename T> __device__ inline T lop3(T a, T b, T c) {
+  T res;
+  asm volatile("lop3.b32 %0, %1, %2, %3, %4;\n"
+               : "=r"(res)
+               : "r"(a), "r"(b), "r"(c), "n"(lut));
+  return res;
+}
 // clang-format on