[SYCLomatic] Fix a bug that kernel arg name may be incorrect when defined in macro (#2425)

Signed-off-by: Huang, Andy <[email protected]>

Author: AndyCHHuang

clang/lib/DPCT/AnalysisInfo.cpp

@@ -5351,7 +5351,7 @@ KernelCallExpr::ArgInfo::ArgInfo(const ParmVarDecl *PVD, KernelCallExpr *Kernel)
   IsRedeclareRequired = false;
 }
 KernelCallExpr::ArgInfo::ArgInfo(std::shared_ptr<TextureObjectInfo> Obj,
-                                 KernelCallExpr *BASE)
+                                 KernelCallExpr *BASE, std::string ArgStr)
     : IsUsedAsLvalueAfterMalloc(false), Texture(Obj) {
   IsPointer = false;
   IsRedeclareRequired = false;
@@ -5360,7 +5360,7 @@ KernelCallExpr::ArgInfo::ArgInfo(std::shared_ptr<TextureObjectInfo> Obj,
   if (auto S = std::dynamic_pointer_cast<StructureTextureObjectInfo>(Obj)) {
     IsDoublePointer = S->containsVirtualPointer();
   }
-  ArgString = Obj->getName();
+  ArgString = ArgStr;
   IdString = ArgString + "_";
   ArgSize = MapNames::KernelArgTypeSizeMap.at(KernelArgType::KAT_Texture);
 }
@@ -5869,14 +5869,16 @@ void KernelCallExpr::buildArgsInfo(const CallExpr *CE) {
       getTheLastCompleteImmediateRange(CE->getBeginLoc(), CE->getEndLoc());
   Analysis.setCallSpelling(KCallSpellingRange.first, KCallSpellingRange.second);
   auto &TexList = getTextureObjectList();
-
   const auto *FD = CE->getDirectCallee();
   const auto *FTD = FD ? FD->getPrimaryTemplate() : nullptr;
   for (unsigned Idx = 0; Idx < CE->getNumArgs(); ++Idx) {
+    auto Arg = CE->getArg(Idx);
+    auto CallDefRange = getDefinitionRange(CE->getBeginLoc(), CE->getEndLoc());
+    auto ArgString = getStringInRange(
+        Arg->getSourceRange(), CallDefRange.getBegin(), CallDefRange.getEnd());
     if (auto Obj = TexList[Idx]) {
-      ArgsInfo.emplace_back(Obj, this);
+      ArgsInfo.emplace_back(Obj, this, ArgString);
     } else {
-      auto Arg = CE->getArg(Idx);
       bool Used = true;
       if (auto *ArgDRE = dyn_cast<DeclRefExpr>(Arg->IgnoreImpCasts()))
         Used = isArgUsedAsLvalueUntil(ArgDRE, CE);

clang/lib/DPCT/AnalysisInfo.h

@@ -2741,7 +2741,8 @@ class KernelCallExpr : public CallFunctionExpr {
     ArgInfo(const ParmVarDecl *PVD, const std::string &ArgsArrayName,
             KernelCallExpr *Kernel);
     ArgInfo(const ParmVarDecl *PVD, KernelCallExpr *Kernel);
-    ArgInfo(std::shared_ptr<TextureObjectInfo> Obj, KernelCallExpr *BASE);
+    ArgInfo(std::shared_ptr<TextureObjectInfo> Obj, KernelCallExpr *BASE,
+            std::string ArgStr);
     inline const std::string &getArgString() const;
     inline const std::string &getTypeString() const;
     inline std::string getIdStringWithIndex() const {

clang/lib/DPCT/CallExprRewriterTexture.cpp

@@ -8,6 +8,7 @@
 
 #include "CallExprRewriter.h"
 #include "CallExprRewriterCommon.h"
+#include "Utility.h"
 
 namespace clang {
 namespace dpct {
@@ -198,6 +199,11 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
             return createRewriter(Call, RetAssign, SourceName);
           }
         } else if (auto DRE = dyn_cast<DeclRefExpr>(SourceExpr)) {
+          auto CallDefRange =
+              getDefinitionRange(Call->getBeginLoc(), Call->getEndLoc());
+          auto DREString =
+              getStringInRange(DRE->getSourceRange(), CallDefRange.getBegin(),
+                               CallDefRange.getEnd());
           auto TexInfo = CallInfo->addTextureObjectArg(SourceIdx, DRE, false);
           if (TexInfo) {
             TexInfo->setType(DpctGlobalInfo::getUnqualifiedTypeName(TargetType),

clang/lib/DPCT/Utility.cpp

@@ -2385,6 +2385,17 @@ getRangeInRange(SourceRange Range, SourceLocation SearchRangeBegin,
                                                    Range.getEnd());
 }
 
+std::string getStringInRange(clang::SourceRange Range,
+                             clang::SourceLocation RangeBegin,
+                             clang::SourceLocation RangeEnd,
+                             bool IncludeLastToken) {
+  auto ResultRange =
+      getRangeInRange(Range, RangeBegin, RangeEnd, IncludeLastToken);
+  auto &SM = dpct::DpctGlobalInfo::getSourceManager();
+  return std::string(SM.getCharacterData(ResultRange.first),
+              SM.getCharacterData(ResultRange.second));
+}
+
 unsigned int calculateIndentWidth(const CUDAKernelCallExpr *Node,
                                   clang::SourceLocation SL, bool &Flag) {
   Flag = true;

clang/lib/DPCT/Utility.h

@@ -446,6 +446,10 @@ getRangeInRange(const clang::Stmt *E, clang::SourceLocation RangeBegin,
 std::pair<clang::SourceLocation, clang::SourceLocation>
 getRangeInRange(clang::SourceRange Range, clang::SourceLocation RangeBegin,
                 clang::SourceLocation RangeEnd, bool IncludeLastToken = true);
+std::string getStringInRange(clang::SourceRange Range,
+                             clang::SourceLocation RangeBegin,
+                             clang::SourceLocation RangeEnd,
+                             bool IncludeLastToken = true);
 unsigned int calculateIndentWidth(const clang::CUDAKernelCallExpr *Node,
                                   clang::SourceLocation SL, bool &Flag);
 bool isIncludedFile(const clang::tooling::UnifiedPath &CurrentFile,

clang/test/dpct/macro_test.cu

@@ -1394,4 +1394,24 @@ __global__ void foo40_kernel() {
 void foo40() {
   foo40_kernel<<<1, 1>>>();
 }
+
+
+template <class T> class MyClass {};
+
+__global__ void foo41(MyClass<float> m) {}
+
+void foo42(MyClass<float> &vecs) {
+#define RUN_APPEND2(DATA) foo41<<<1, 1, 0>>>(DATA);
+//CHECK: #define RUN_APPEND2(DATA)                                                      \
+//CHECK-NEXT:   dpct::get_in_order_queue().submit([&](sycl::handler &cgh) {                  \
+//CHECK-NEXT:     auto DATA_ct0 = DATA;                                                      \
+//CHECK-NEXT:                                                                                \
+//CHECK-NEXT:     cgh.parallel_for(                                                          \
+//CHECK-NEXT:         sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),   \
+//CHECK-NEXT:         [=](sycl::nd_item<3> item_ct1) { foo41(DATA_ct0); });                  \
+//CHECK-NEXT:   });
+  RUN_APPEND2(vecs);
+}
+
+
 #endif