[SYCLomatic] Cast the function name argument for get_kernel_function (#2806)

Author: the-slow-one

clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h

@@ -149,6 +149,8 @@ template <class SubExprT> class CastIfNotSameExprPrinter {
     clang::QualType ArgType = InputArg->getType().getCanonicalType();
     ArgType.removeLocalFastQualifiers(clang::Qualifiers::CVRMask);
     bool NeedParen = false;
+    std::cout << "Arg type: " << ArgType.getAsString() << "\n";
+    std::cout << "Given type " << TypeInfo << "\n";
     if (ArgType.getAsString() != TypeInfo) {
       NeedParen = needExtraParens(SubExpr);
       Stream << "(" << TypeInfo << ")";

clang/lib/DPCT/RulesLang/APINamesDriver.inc

@@ -47,7 +47,7 @@ FEATURE_REQUEST_FACTORY(
     ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY("cuModuleGetFunction", DEREF(0),
                                             CALL(MapNames::getDpctNamespace() +
                                                      "get_kernel_function",
-                                                 ARG(1), ARG(2)))))
+                                                 ARG(1), CAST_IF_NOT_SAME(makeLiteral("const char *"), ARG(2))))))
 
 FEATURE_REQUEST_FACTORY(
     HelperFeatureEnum::device_ext,

clang/test/dpct/kernel-function-typecast.cu

@@ -5,6 +5,8 @@
 #include <cstdint>
 #include <cuda.h>
 
+#include <string>
+
 typedef uint64_t u64;
 
 // CHECK: void exec_kernel(dpct::kernel_function cuFunc, dpct::kernel_library cuMod, dpct::queue_ptr stream) {
@@ -15,9 +17,10 @@ void exec_kernel(CUfunction cuFunc, CUmodule cuMod, CUstream stream) {
   // verify the conversion from dpct::kernel_library to uint64_t
   mod = (u64)cuMod;
 
+  std::string kernel_name{"kfoo"};
   // verify the conversion from uint64_t to dpct::kernel_library
-  // CHECK: cuFunc = dpct::get_kernel_function((dpct::kernel_library)mod, "kfoo");
-  cuModuleGetFunction(&cuFunc, (CUmodule)mod, "kfoo");
+  // CHECK: cuFunc = dpct::get_kernel_function((dpct::kernel_library)mod, kernel_name.c_str());
+  cuModuleGetFunction(&cuFunc, (CUmodule)mod, kernel_name.c_str());
 
   // verify the conversion from dpct::kernel_function to uint64_t
   function = (u64)cuFunc;
@@ -28,3 +31,18 @@ void exec_kernel(CUfunction cuFunc, CUmodule cuMod, CUstream stream) {
   // CHECK: dpct::invoke_kernel_function((dpct::kernel_function)function, *stream, sycl::range<3>(100, 100, 100), sycl::range<3>(100, 100, 100), 1024, NULL, config);
   cuLaunchKernel((CUfunction)function, 100, 100, 100, 100, 100, 100, 1024, stream, NULL, config);
 }
+
+class CString {
+  private:
+    char *str;
+  public:
+    CString(): str(NULL) {};
+    operator const char* () const { return str; }
+
+    operator char* () { return str; }
+};
+
+void test_casting(CUmodule mod, CUfunction func, const CString &name) {
+  // CHECK: func = dpct::get_kernel_function(mod, (const char *)name);
+  cuModuleGetFunction(&func, mod, name);
+}