add support for templates in cuda-python backend

Author: benvanwerkhoven

kernel_tuner/core.py

@@ -68,7 +68,7 @@ def __init__(self, kernel_name, kernel_sources, lang, defines=None):
             lang = util.detect_language(kernel_string)
 
         # The validity of lang is checked later, when creating the DeviceInterface
-        self.lang = lang
+        self.lang = lang.upper()
 
     def get_kernel_string(self, index=0, params=None):
         """ retrieve the kernel source with the given index and return as a string
@@ -529,7 +529,7 @@ def create_kernel_instance(self, kernel_source, kernel_options, params, verbose)
                                                                               kernel_options.block_size_names)
 
         #check for templated kernel
-        if kernel_source.lang == "CUDA" and "<" in name and ">" in name:
+        if kernel_source.lang in ["CUDA", "NVCUDA"] and "<" in name and ">" in name:
             kernel_string, name = wrap_templated_kernel(kernel_string, name)
 
         #collect everything we know about this instance and return it

kernel_tuner/nvcuda.py

@@ -199,21 +199,30 @@ def compile(self, kernel_instance):
             self.compiler_options.append(f"--gpu-architecture=compute_{self.cc}")
 
         err, program = nvrtc.nvrtcCreateProgram(str.encode(kernel_string), b"CUDAProgram", 0, [], [])
-        error_check(err)
-        err = nvrtc.nvrtcCompileProgram(program, len(compiler_options), compiler_options)
-        error_check(err)
-        err, size = nvrtc.nvrtcGetPTXSize(program)
-        error_check(err)
-        buff = b' ' * size
-        err = nvrtc.nvrtcGetPTX(program, buff)
-        error_check(err)
-        err, self.current_module = cuda.cuModuleLoadData(np.char.array(buff))
-        if err == cuda.CUresult.CUDA_ERROR_INVALID_PTX:
-            raise SkippableFailure("uses too much shared data")
-        else:
+        try:
             error_check(err)
-        err, self.func = cuda.cuModuleGetFunction(self.current_module, str.encode(kernel_name))
-        error_check(err)
+            err = nvrtc.nvrtcCompileProgram(program, len(compiler_options), compiler_options)
+            error_check(err)
+            err, size = nvrtc.nvrtcGetPTXSize(program)
+            error_check(err)
+            buff = b' ' * size
+            err = nvrtc.nvrtcGetPTX(program, buff)
+            error_check(err)
+            err, self.current_module = cuda.cuModuleLoadData(np.char.array(buff))
+            if err == cuda.CUresult.CUDA_ERROR_INVALID_PTX:
+                raise SkippableFailure("uses too much shared data")
+            else:
+                error_check(err)
+            err, self.func = cuda.cuModuleGetFunction(self.current_module, str.encode(kernel_name))
+            error_check(err)
+
+        except RuntimeError as re:
+            _, n = nvrtc.nvrtcGetProgramLogSize(program)
+            log = b' ' * n
+            nvrtc.nvrtcGetProgramLog(program, log)
+            print(log.decode('utf-8'))
+            raise re
+
         return self.func
 
     def start_event(self):