[SYCLomatic] Fix the issue that loop unroll rule insertion may conflict with nd_item declaration insertion (#2831)

Signed-off-by: intwanghao <hao3.wang@intel.com>

Author: intwanghao

clang/lib/DPCT/RulesLang/OptimizeMigration.cpp

@@ -85,7 +85,9 @@ void ForLoopUnrollRule::runRule(
     auto IndentStr = getIndent(ForLoopLoc, SM).str();
     Repl = Repl + "#pragma unroll" + getNL() + IndentStr;
     auto Begin = SM.getSpellingLoc(ForLoop->getBeginLoc());
-    emplaceTransformation(new InsertText(Begin, Repl));
+    emplaceTransformation(new InsertText(Begin, Repl, 0,
+                                         ReplacementType::RT_ForSYCLMigration,
+                                         InsertPosition::IP_Right));
   }
 }
 

clang/lib/DPCT/TextModification.h

@@ -231,9 +231,11 @@ class InsertText : public TextModification {
 
 public:
   InsertText(SourceLocation Loc, const std::string &S, unsigned PairID = 0,
-             ReplacementType IsForCodePin = RT_ForSYCLMigration)
+             ReplacementType IsForCodePin = RT_ForSYCLMigration,
+             InsertPosition IP = InsertPosition::IP_Left)
       : TextModification(TMID::InsertText), Begin(Loc), T(S), PairID(PairID) {
     this->IsForCodePin = IsForCodePin;
+    setInsertPosition(IP);
   }
   std::shared_ptr<ExtReplacement>
   getReplacement(const ASTContext &Context) const override;

clang/test/dpct/opt_loop_unroll.cu

@@ -110,9 +110,34 @@ __device__ void bar(float *a, float *b) {
   copy(a, b, 10)
 }
 
+
+// CHECK: __dpct_inline__ void foo2(float *a) {
+// CHECK:   auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
+// CHECK:   #pragma unroll
+// CHECK:   for(int i = 0; i < 10; i++) {
+// CHECK:     a[i] = item_ct1.get_local_id(2);
+// CHECK:   }
+// CHECK:   #pragma unroll
+// CHECK:   for(int i = 0; i < 10; i++) {
+// CHECK:     a[i] = item_ct1.get_local_id(1);
+// CHECK:   }
+// CHECK: }
+__global__ void foo2(float *a) {
+  for(int i = 0; i < 10; i++) {
+    a[i] = threadIdx.x;
+  }
+
+  for(int i = 0; i < 10; i++) {
+    a[i] = threadIdx.y;
+  }
+
+}
+
+
 int main(){
     float *a, *b, *c;
     kernel<<<10, 10>>>(a, b, c);
+    foo2<<<10, 10>>>(a);
     return 0;
 }
 #endif