[SYCLomatic] Make the kernel function wrapper keep origin inline or static specifier (#2742)

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

Author: intwanghao

clang/lib/DPCT/AnalysisInfo.cpp

@@ -5249,7 +5249,9 @@ void DeviceFunctionDecl::insertWrapper() {
         Printer << ">";
         Printer.newLine();
       }
-      Printer << "void " << FuncName << "_wrapper(";
+      Printer << (IsStaticSpecified ? "static " : "")
+              << (IsInlineSpecified ? "inline " : "") << "void " << FuncName
+              << "_wrapper(";
       for (size_t i = 0; i < ParamsInfo.size(); i++) {
         Printer << (i == 0 ? "" : " ,") << ParamsInfo[i].first << " "
                 << ParamsInfo[i].second << ParameterDefaultValueMap[i];
@@ -5329,7 +5331,8 @@ void DeviceFunctionDecl::collectInfoForWrapper(const FunctionDecl *FD) {
   if (HasBody && FD != Def) {
     HasBody = false;
   }
-
+  IsInlineSpecified = FD->isInlineSpecified();
+  IsStaticSpecified = FD->isStatic();
   if (auto FTD = FD->getDescribedFunctionTemplate()) {
     if (auto TemplateParmsList = FTD->getTemplateParameters()) {
       for (size_t i = 0; i < TemplateParmsList->size(); ++i) {

clang/lib/DPCT/AnalysisInfo.h

@@ -2617,6 +2617,8 @@ class DeviceFunctionDecl {
   std::vector<std::shared_ptr<TextureObjectInfo>> TextureObjectList;
   FormatInfo FormatInformation;
   bool HasBody = false;
+  bool IsInlineSpecified = false;
+  bool IsStaticSpecified = false;
   size_t DeclEnd = 0;
   std::map<int, std::string> TemplateParameterDefaultValueMap;
   std::map<int, std::string> ParameterDefaultValueMap;

clang/test/dpct/function_pointer.cu

@@ -5,14 +5,14 @@
 #include <cuda_runtime.h>
 #include <iostream>
 
-__global__ void vectorAdd(const int *A, int *B, int *C, int N) {
+__global__ static inline void vectorAdd(const int *A, int *B, int *C, int N) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i < N) {
         C[i] = A[i] + B[i];
     }
 }
 
-// CHECK:  void vectorAdd_wrapper(const int * A ,int * B ,int * C ,int N) {
+// CHECK:  static inline void vectorAdd_wrapper(const int * A ,int * B ,int * C ,int N) {
 // CHECK:        sycl::queue queue = *dpct::kernel_launcher::_que;
 // CHECK:        unsigned int localMemSize = dpct::kernel_launcher::_local_mem_size;
 // CHECK:        sycl::nd_range<3> nr = dpct::kernel_launcher::_nr;
@@ -24,15 +24,15 @@ __global__ void vectorAdd(const int *A, int *B, int *C, int N) {
 // CHECK:  }
 
 template<typename T>
-__global__ void vectorTemplateAdd(const T *A, T *B, T *C, int N) {
+__global__ static inline void vectorTemplateAdd(const T *A, T *B, T *C, int N) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i < N) {
         C[i] = A[i] + B[i];
     }
 }
 
 // CHECK:  template<typename T>
-// CHECK:  void vectorTemplateAdd_wrapper(const T * A ,T * B ,T * C ,int N) {
+// CHECK:  static inline void vectorTemplateAdd_wrapper(const T * A ,T * B ,T * C ,int N) {
 // CHECK:      sycl::queue queue = *dpct::kernel_launcher::_que;
 // CHECK:      unsigned int localMemSize = dpct::kernel_launcher::_local_mem_size;
 // CHECK:      sycl::nd_range<3> nr = dpct::kernel_launcher::_nr;