[SYCLomatic] Enable module-kernel test

features/config/TEMPLATE_module.xml

@@ -3,8 +3,10 @@
 <test driverID="test_feature" name="TEMPLATE">
     <description>test</description>
     <files>
-      <file path="feature_case/module/module-kernel.cu" />
-      <file path="feature_case/module/module-main.cu" />
-      <file path="feature_case/module/module-helper.cpp" />
+      <file path="feature_case/module" />
     </files>
+    <rules>
+        <platformRule OSFamily="Linux" kit="CUDA12.0" kitRange="LATER_OR_EQUAL" runOnThisPlatform="false"/>
+        <platformRule OSFamily="Windows" kit="CUDA12.0" kitRange="LATER_OR_EQUAL" runOnThisPlatform="false"/>
+    </rules>
 </test>

features/feature_case/module/do_test.py

@@ -0,0 +1,42 @@
+# ====------ do_test.py---------- *- Python -* ----===##
+#
+# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#
+#
+# ===----------------------------------------------------------------------===#
+import subprocess
+import platform
+import os
+import sys
+from test_config import CT_TOOL
+
+from test_utils import *
+
+def setup_test():
+    change_dir(test_config.current_test)
+    return True
+
+def migrate_test():
+    # clean previous migration output
+    if (os.path.exists("dpct_output")):
+        shutil.rmtree("dpct_output")    
+    call_subprocess(test_config.CT_TOOL + " --cuda-include-path=" + test_config.include_path + " module-helper.cpp module-main.cu module-kernel.cu --extra-arg=--ptx")
+    return os.path.exists(os.path.join("dpct_output", "module-kernel.dp.cpp"))
+def build_test():
+    # make shared library
+    if (platform.system() == 'Windows'):
+        ret = call_subprocess("icpx -fsycl              dpct_output/module-kernel.dp.cpp       -shared -o module-kernel.dll")
+    else:
+        ret = call_subprocess(test_config.DPCXX_COM + " dpct_output/module-kernel.dp.cpp -fPIC -shared -o module-kernel.so")
+    if not ret:
+        print("Could not make module-kernel.* shared library.")
+        return False
+
+    srcs = []
+    srcs.append(os.path.join("dpct_output", "module-helper.cpp"))
+    srcs.append(os.path.join("dpct_output", "module-main.dp.cpp"))
+    return compile_and_link(srcs)
+def run_test():
+    return call_subprocess(os.path.join(os.path.curdir, test_config.current_test + '.run '))

features/feature_case/module/module-kernel.cu

@@ -8,9 +8,9 @@
 // ===----------------------------------------------------------------------===//
 static texture<int, 3> tex;
 
-__global__ void foo(float* k, float* y);
+extern "C" __global__ void foo(float* k, float* y);
 
-__global__ void foo(float* k, float* y){
+extern "C" __global__ void foo(float* k, float* y){
     extern __shared__ int s[];
     int a = threadIdx.x;
 }

features/feature_case/module/module-main.cu

@@ -10,10 +10,14 @@
 int main(){
     CUmodule M;
     CUfunction F;
-    std::string Path, FunctionName, Data;
+#ifdef _WIN32
+    std::string Path{"./module-kernel.dll"};
+#else
+    std::string Path{"./module-kernel.so"};
+#endif
+    std::string FunctionName{"foo"}, Data;
     FunctionName = "foo";
     cuModuleLoad(&M, Path.c_str());
-    cuModuleLoadData(&M, Data.c_str());
     cuModuleGetFunction(&F, M, FunctionName.c_str());
     float *param[2] = {0};
     cudaMalloc(&param[0], sizeof(float));

features/features.xml

@@ -102,7 +102,7 @@
     <test testName="cufft-func-ptr" configFile="config/TEMPLATE_cuFFT.xml" />
     <test testName="cufft-others" configFile="config/TEMPLATE_cuFFT.xml" />
     <test testName="cufft-external-workspace" configFile="config/cufft-external-workspace.xml" splitGroup="double"/>
-    <!-- <test testName="module-kernel" configFile="config/TEMPLATE_module.xml" /> -->
+    <test testName="module-kernel" configFile="config/TEMPLATE_module.xml" />
     <test testName="macro" configFile="config/TEMPLATE_macro.xml" />
     <test testName="driverMem" configFile="config/TEMPLATE_driverMem.xml" />
     <test testName="driverCtx" configFile="config/TEMPLATE_driverCtx.xml" />