[SYCLomatic] Ensure wrappers are only generated under C language link…

…age and for __global__ functions (#338) * [SYCLomatic] Ensure wrappers are only generated under C language linkage and for __global__ functions Signed-off-by: Cai, Justin <justin.cai@intel.com> * Fix other module lit test * Move check to isModuleFunction * Fix windows lit test * Fix windows lit test 2 * Add more cases in lit test Signed-off-by: Cai, Justin <justin.cai@intel.com>
oneapi-src · Nov 17, 2022 · 57a1c25 · 57a1c25
1 parent af325b6
commit 57a1c25
Show file tree

Hide file tree

Showing 4 changed files with 41 additions and 10 deletions.
diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp
@@ -2488,12 +2488,12 @@ void DeviceFunctionDeclInModule::buildCallInfo(const FunctionDecl *FD) {
 
 bool isModuleFunction(const FunctionDecl *FD) {
   auto &SM = DpctGlobalInfo::getSourceManager();
-  if (DpctGlobalInfo::getModuleFiles().find(
+  return
+    FD->getLanguageLinkage() == CLanguageLinkage
+    && FD->hasAttr<CUDAGlobalAttr>()
+    && DpctGlobalInfo::getModuleFiles().find(
           DpctGlobalInfo::getLocInfo(SM.getExpansionLoc(FD->getBeginLoc()))
-              .first) != DpctGlobalInfo::getModuleFiles().end()) {
-    return true;
-  }
-  return false;
+              .first) != DpctGlobalInfo::getModuleFiles().end();
 }
 
 DeviceFunctionDecl::DeviceFunctionDecl(unsigned Offset,

diff --git a/clang/test/dpct/module_kernel.cu b/clang/test/dpct/module_kernel.cu
@@ -5,13 +5,13 @@
 //CHECK: dpct::image_wrapper_base_p tex;
 CUtexref tex;
 
-//CHECK: void foo(float* k, float* y, sycl::nd_item<3> item_ct1, uint8_t *dpct_local,
-//CHECK-NEXT:     unsigned int *const_data);
+//CHECK: extern "C" void foo(float* k, float* y, sycl::nd_item<3> item_ct1,
+//CHECK-NEXT:     uint8_t *dpct_local, unsigned int *const_data);
 
 //CHECK:extern "C" {
 //CHECK-NEXT:void foo_wrapper(sycl::queue &queue, const sycl::nd_range<3> &nr, unsigned int localMemSize, void **kernelParams, void **extra);
 //CHECK-NEXT:}
-__global__ void foo(float* k, float* y);
+extern "C" __global__ void foo(float* k, float* y);
 
 
 __constant__ unsigned int const_data[3] = {1, 2, 3};

diff --git a/clang/test/dpct/module_kernel_win.cu b/clang/test/dpct/module_kernel_win.cu
@@ -5,12 +5,13 @@
 //CHECK: dpct::image_wrapper_base_p tex;
 CUtexref tex;
 
-//CHECK: void foo(float* k, float* y, sycl::nd_item<3> item_ct1, uint8_t *dpct_local);
+//CHECK: extern "C" void foo(float* k, float* y, sycl::nd_item<3> item_ct1,
+//CHECK-NEXT: uint8_t *dpct_local);
 
 //CHECK: extern "C" {
 //CHECK-NEXT: __declspec(dllexport) void foo_wrapper(sycl::queue &queue, const sycl::nd_range<3> &nr, unsigned int localMemSize, void **kernelParams, void **extra);
 //CHECK-NEXT: }
-__global__ void foo(float* k, float* y);
+extern "C" __global__ void foo(float* k, float* y);
 
 
 //CHECK: void foo(float* k, float* y, sycl::nd_item<3> item_ct1, uint8_t *dpct_local){

diff --git a/clang/test/dpct/module_wrapper_gen.cu b/clang/test/dpct/module_wrapper_gen.cu
@@ -0,0 +1,30 @@
+// RUN: dpct --extra-arg="--ptx" \
+// RUN:      --out-root=%T/module_wrapper_gen \
+// RUN:      --cuda-include-path="%cuda-path/include" %s
+// RUN: FileCheck %s --input-file=%T/module_wrapper_gen/module_wrapper_gen.dp.cpp
+
+// START
+
+__device__ float2 operator+(float2 a, float2 b) {
+  return float2{a.x - b.x, a.y - b.y};
+}
+
+extern "C" __device__ void externCNonKernel() {}
+
+__device__ void deviceFun() {}
+
+__global__ void nonExternCKernel() {}
+
+extern "C" __global__ void exampleKernel() {
+  float2 x;
+  float2 y;
+  float2 z = x + y;
+}
+
+// END
+
+// Only one _wrapper should be generated.
+// CHECK:     // START
+// CHECK:     _wrapper
+// CHECK-NOT: _wrapper
+// CHECK:     // END