[SYCLomatic] Fix __ldg used in macro (#2706)

Signed-off-by: Jiang, Zhiwei <zhiwei.jiang@intel.com>

Author: zhiweij1

clang/lib/DPCT/RuleInfra/CallExprRewriter.cpp

@@ -63,10 +63,24 @@ AddrOfExpr::AddrOfExpr(const Expr *E, const CallExpr *C) {
 }
 
 DerefExpr::DerefExpr(const Expr *E, const CallExpr *C) {
+  const auto &SM = DpctGlobalInfo::getSourceManager();
   this->C = C;
   // If E is UnaryOperator or CXXOperatorCallExpr D.E will has value
   this->E = getDereferencedExpr(E);
   if (this->E) {
+    if (C) {
+      // Check the addrof symbol (&) is in the parent range since only it
+      // will be merged with the deref symbol (*)
+      if (E->getBeginLoc().isMacroID()) {
+        auto Range = getDefinitionRange(C->getBeginLoc(), C->getEndLoc());
+        if (!isInRange(Range.getBegin(), Range.getEnd(),
+                       SM.getSpellingLoc(E->getBeginLoc()))) {
+          this->E = E;
+          this->NeedParens = true;
+          return;
+        }
+      }
+    }
     this->E = this->E->IgnoreParens();
     this->AddrOfRemoved = true;
   } else {

clang/lib/DPCT/RuleInfra/CallExprRewriter.h

@@ -800,6 +800,8 @@ class DerefExpr {
 
 public:
   DerefExpr(const Expr *E, const CallExpr *C = nullptr);
+  DerefExpr(std::pair<const CallExpr *, const Expr *> P)
+      : DerefExpr(P.second, P.first) {}
   template <class StreamT>
   void printArg(StreamT &Stream, ArgumentAnalysis &A) const {
     print(Stream);

clang/lib/DPCT/RulesLang/Math/RewriterHalfPrecisionConversionAndDataMovement.cpp

@@ -797,7 +797,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() {
                         EMPTY_FACTORY_ENTRY("__ldca"),
                         EMPTY_FACTORY_ENTRY("__ldca"),
                         WARNING_FACTORY_ENTRY(
-                            "__ldca", DEREF_FACTORY_ENTRY("__ldca", ARG(0)),
+                            "__ldca", DEREF_FACTORY_ENTRY("__ldca", ARG_WC(0)),
                             Diagnostics::MATH_EMULATION_EXPRESSION,
                             std::string("__ldca"), std::string("'*'"))))
       // __ldcg
@@ -807,7 +807,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() {
                         EMPTY_FACTORY_ENTRY("__ldcg"),
                         EMPTY_FACTORY_ENTRY("__ldcg"),
                         WARNING_FACTORY_ENTRY(
-                            "__ldcg", DEREF_FACTORY_ENTRY("__ldcg", ARG(0)),
+                            "__ldcg", DEREF_FACTORY_ENTRY("__ldcg", ARG_WC(0)),
                             Diagnostics::MATH_EMULATION_EXPRESSION,
                             std::string("__ldcg"), std::string("'*'"))))
       // __ldcs
@@ -817,7 +817,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() {
                         EMPTY_FACTORY_ENTRY("__ldcs"),
                         EMPTY_FACTORY_ENTRY("__ldcs"),
                         WARNING_FACTORY_ENTRY(
-                            "__ldcs", DEREF_FACTORY_ENTRY("__ldcs", ARG(0)),
+                            "__ldcs", DEREF_FACTORY_ENTRY("__ldcs", ARG_WC(0)),
                             Diagnostics::MATH_EMULATION_EXPRESSION,
                             std::string("__ldcs"), std::string("'*'"))))
       // __ldcv
@@ -827,7 +827,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() {
                         EMPTY_FACTORY_ENTRY("__ldcv"),
                         EMPTY_FACTORY_ENTRY("__ldcv"),
                         WARNING_FACTORY_ENTRY(
-                            "__ldcv", DEREF_FACTORY_ENTRY("__ldcv", ARG(0)),
+                            "__ldcv", DEREF_FACTORY_ENTRY("__ldcv", ARG_WC(0)),
                             Diagnostics::MATH_EMULATION_EXPRESSION,
                             std::string("__ldcv"), std::string("'*'"))))
       // __ldg
@@ -837,7 +837,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() {
                        EMPTY_FACTORY_ENTRY("__ldg"),
                        EMPTY_FACTORY_ENTRY("__ldg"),
                        WARNING_FACTORY_ENTRY(
-                           "__ldg", DEREF_FACTORY_ENTRY("__ldg", ARG(0)),
+                           "__ldg", DEREF_FACTORY_ENTRY("__ldg", ARG_WC(0)),
                            Diagnostics::MATH_EMULATION_EXPRESSION,
                            std::string("__ldg"), std::string("'*'"))))
       // __ldlu
@@ -847,7 +847,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() {
                         EMPTY_FACTORY_ENTRY("__ldlu"),
                         EMPTY_FACTORY_ENTRY("__ldlu"),
                         WARNING_FACTORY_ENTRY(
-                            "__ldlu", DEREF_FACTORY_ENTRY("__ldlu", ARG(0)),
+                            "__ldlu", DEREF_FACTORY_ENTRY("__ldlu", ARG_WC(0)),
                             Diagnostics::MATH_EMULATION_EXPRESSION,
                             std::string("__ldlu"), std::string("'*'"))))
       // __ll2half_rd

clang/test/dpct/macro_test.cu

@@ -1520,4 +1520,13 @@ extern EXPLICIT_DECL(half);
 #undef FROMDEV3
 #undef FROMDEV2
 
+// CHECK: #define VLLM_LDG(arg) *(arg)
+// CHECK-NEXT: void foo46(const float *__restrict__ input) {
+// CHECK-NEXT:   const float x = VLLM_LDG(&input[13]);
+// CHECK-NEXT: }
+#define VLLM_LDG(arg) __ldg(arg)
+__global__ void foo46(const float *__restrict__ input) {
+  const float x = VLLM_LDG(&input[13]);
+}
+
 #endif

clang/test/dpct/memory_management.cu

@@ -198,7 +198,7 @@ void checkError(cudaError_t err) {
 
 void cuCheckError(CUresult err) {
 }
-// CHECK: #define PITCH(a, b, c, d) a = (float *)dpct::dpct_malloc(b, c, d);
+// CHECK: #define PITCH(a, b, c, d) *(a) = (float *)dpct::dpct_malloc(*(b), c, d);
 #define PITCH(a, b, c, d) cudaMallocPitch(a, b, c, d);
 
 void testCommas() {