[SYCLomatic] Fix the issue that SYCL_EXTERNAL not inserted into right position in macro scenario(#2734)

Signed-off-by: intwanghao <[email protected]>

Author: intwanghao

clang/lib/DPCT/AnalysisInfo.cpp

@@ -4823,6 +4823,8 @@ DeviceFunctionDecl::DeviceFunctionDecl(
       FuncInfo(getFuncInfo(FD)) {
   if (FD->isFunctionTemplateSpecialization()) {
     SourceRange ReturnTypeRange = FD->getReturnTypeSourceRange();
+    ReturnTypeRange = clang::dpct::getDefinitionRange(
+        ReturnTypeRange.getBegin(), ReturnTypeRange.getEnd());
     OffsetForAttr =
         DpctGlobalInfo::getLocInfo(ReturnTypeRange.getBegin()).second;
   }
@@ -4859,6 +4861,11 @@ DeviceFunctionDecl::DeviceFunctionDecl(
       FuncInfo(getFuncInfo(Specialization)) {
   IsDefFilePathNeeded = false;
 
+  auto ReturnLoc = FTL.getReturnLoc();
+  auto ReturnRange = clang::dpct::getDefinitionRange(ReturnLoc.getBeginLoc(),
+                                                     ReturnLoc.getEndLoc());
+  OffsetForAttr = DpctGlobalInfo::getLocInfo(ReturnRange.getBegin()).second;
+
   buildReplaceLocInfo(FTL, Attrs);
   buildTextureObjectParamsInfo(FTL.getParams());
   if (Specialization->hasAttr<CUDAGlobalAttr>()) {

clang/test/dpct/mf-kernel.cu

@@ -10,6 +10,15 @@ __global__ void Reset_kernel_parameters(void)
     g_mutex=0;
 }
 
+// CHECK: template <typename T> SYCL_EXTERNAL void template_test(T v) {};
+// CHECK: #define FUNC(T) template SYCL_EXTERNAL void template_test<T>(T v);
+// CHECK: FUNC(int)
+template <typename T> __global__ void template_test(T v) {};
+
+#define FUNC(T) template __global__ void template_test<T>(T v);
+
+FUNC(int)
+
 // CHECK: SYCL_EXTERNAL void kernel_extern(const sycl::nd_item<3> &item_ct1, int *a) {
 __global__ void kernel_extern() {
   __shared__ int a[360];

clang/test/dpct/mf-kernel.cuh

@@ -5,6 +5,9 @@
 // CHECK: SYCL_EXTERNAL void Reset_kernel_parameters(volatile int &g_mutex);
 __global__ void Reset_kernel_parameters(void);
 
+// CHECK: template <typename T> SYCL_EXTERNAL void template_test(T v);
+template <typename T> __global__ void template_test(T v);
+
 // CHECK: SYCL_EXTERNAL void test_foo();
 __device__ void test_foo(void);
 

clang/test/dpct/mf-test.cu

@@ -77,6 +77,9 @@ void test() {
   // CHECK-NEXT:       });
   // CHECK-NEXT:   }
   test_fp<<<1, 1>>>();
+
+  template_test<int><<<1, 1>>>(0);
+
 }