From 894917a4916cf47878a3f8e1310adeb284c05439 Mon Sep 17 00:00:00 2001 From: Ben van Werkhoven Date: Fri, 11 Sep 2026 11:44:36 +0200 Subject: [PATCH 1/4] trying to get this to work with cuModuleLoadData --- kernel_tuner/backends/nvcuda.py | 107 +++++++++++++++++++++++++++----- 1 file changed, 90 insertions(+), 17 deletions(-) diff --git a/kernel_tuner/backends/nvcuda.py b/kernel_tuner/backends/nvcuda.py index c5f9b577..a22787c1 100644 --- a/kernel_tuner/backends/nvcuda.py +++ b/kernel_tuner/backends/nvcuda.py @@ -18,7 +18,7 @@ def preload_python_nvrtc(): # search site-packages for nvidia-cuda-nvrtc wheels site_packages = sysconfig.get_paths()["purelib"] nvidia_path = os.path.join(site_packages, "nvidia", "nvrtc", "lib") - + # fall back to CUDA_HOME / CUDA_PATH if present cuda_home = os.environ.get("CUDA_HOME") or os.environ.get("CUDA_PATH") cuda_home_lib = os.path.join(cuda_home, "lib64") if cuda_home else None @@ -290,38 +290,70 @@ def compile(self, kernel_instance): err, program = nvrtc.nvrtcCreateProgram(str.encode(kernel_string), b"CUDAProgram", 0, [], []) try: + cuda_error_check(err) # Add the kernel as an expression. This is necessary for templated kernels to ensure that the # compiler actually instantiates the kernel that we want to compile. - cuda_error_check(err) err = nvrtc.nvrtcAddNameExpression(program, expression_name) + cuda_error_check(err) # Compile the program - cuda_error_check(err) err = nvrtc.nvrtcCompileProgram(program, len(compiler_options), compiler_options) - - # Get the PTX - cuda_error_check(err) - err, size = nvrtc.nvrtcGetPTXSize(program) - cuda_error_check(err) - buff = b" " * size - err = nvrtc.nvrtcGetPTX(program, buff) cuda_error_check(err) + if b"-enable-tile" in compiler_options: + # Get the Tile IR + err, size = nvrtc.nvrtcGetTileIRSize(program) + cuda_error_check(err) + buff = b" " * size + err = nvrtc.nvrtcGetTileIR(program, buff) + cuda_error_check(err) + else: + # Get the PTX + err, size = nvrtc.nvrtcGetPTXSize(program) + cuda_error_check(err) + buff = b" " * size + err = nvrtc.nvrtcGetPTX(program, buff) + cuda_error_check(err) + # Load the module - err, self.current_module = driver.cuModuleLoadData(np.char.array(buff)) + err, self.current_module = _load_module_with_logs(np.char.array(buff)) + + # It the compile succeeded, but loading the PTX failed it is most likely + # that the kernel uses too much shared memory if err == driver.CUresult.CUDA_ERROR_INVALID_PTX: raise SkippableFailure("uses too much shared data") else: cuda_error_check(err) - # First, get the "lowered" name of the kernel (i.e., the name inside the PTX). - # After, we can use the lowered name to lookup the kernel in the module. + # Get the lowered name (resolves C++ mangling for both regular and Tile kernels). err, lowered_name = nvrtc.nvrtcGetLoweredName(program, expression_name) cuda_error_check(err) - err, self.func = driver.cuModuleGetFunction( - self.current_module, lowered_name - ) - cuda_error_check(err) + + if b"-enable-tile" in compiler_options: + # For Tile kernels the entry point name in the module is the lowered name + # plus encoded Tile template parameters. Enumerate all functions in the + # loaded module and find the one whose name contains the kernel name. + err, func_count = driver.cuModuleGetFunctionCount(self.current_module) + cuda_error_check(err) + err, functions = driver.cuModuleEnumerateFunctions(func_count, self.current_module) + cuda_error_check(err) + self.func = None + for func in functions: + err, func_name = driver.cuFuncGetName(func) + if err != driver.CUresult.CUDA_SUCCESS: + continue + if isinstance(func_name, bytes): + func_name = func_name.decode() + if kernel_name in func_name: + self.func = func + break + if self.func is None: + raise RuntimeError(f"Could not find Tile kernel '{kernel_name}' in compiled module") + else: + err, self.func = driver.cuModuleGetFunction(self.current_module, lowered_name) + if err == driver.CUresult.CUDA_ERROR_NOT_FOUND: + err, self.func = driver.cuModuleGetFunction(self.current_module, expression_name) + cuda_error_check(err) # get the number of registers per thread used in this kernel num_regs = driver.cuFuncGetAttribute(driver.CUfunction_attribute.CU_FUNC_ATTRIBUTE_NUM_REGS, self.func) @@ -479,3 +511,44 @@ def memcpy_htod(dest, src): units = {"time": "ms"} last_selected_device = None + + + +def _load_module_with_logs(image: bytes, log_size: int = 8192): + """Load a PTX/cubin/Tile IR image via cuModuleLoadDataEx with + error/info log capture. Raises RuntimeError with the compiler's + own diagnostic text on failure, instead of a bare CUresult.""" + + error_log = bytearray(log_size) + info_log = bytearray(log_size) + + options = [ + driver.CUjit_option.CU_JIT_ERROR_LOG_BUFFER, + driver.CUjit_option.CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES, + driver.CUjit_option.CU_JIT_INFO_LOG_BUFFER, + driver.CUjit_option.CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES, + driver.CUjit_option.CU_JIT_LOG_VERBOSE, + ] + option_values = [ + error_log, + log_size, + info_log, + log_size, + 1, # verbose logging on + ] + + result = driver.cuModuleLoadDataEx( + image, len(options), options, option_values + ) + + if result[0] != driver.CUresult.CUDA_SUCCESS: + err_text = error_log.split(b"\x00", 1)[0].decode(errors="replace") + info_text = info_log.split(b"\x00", 1)[0].decode(errors="replace") + raise RuntimeError( + f"cuModuleLoadDataEx failed: {result[0]}\n" + f"--- error log ---\n{err_text}\n" + f"--- info log ---\n{info_text}" + ) + + (module,) = result[1:] + return result[0], module From 9b483f5c54f6d8444530f88274caaf3ce4a237c3 Mon Sep 17 00:00:00 2001 From: Ben van Werkhoven Date: Fri, 11 Sep 2026 14:17:05 +0200 Subject: [PATCH 2/4] ensure template wrapper kernel is only used for PYCUDA backend --- kernel_tuner/core.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/kernel_tuner/core.py b/kernel_tuner/core.py index 4a715696..f51e7ddd 100644 --- a/kernel_tuner/core.py +++ b/kernel_tuner/core.py @@ -336,6 +336,7 @@ def __init__( from kernel_tuner.backends.pycuda import PyCudaFunctions backend = PyCudaFunctions + lang = "PYCUDA" elif lang.upper() == "CUPY": from kernel_tuner.backends.cupy import CupyFunctions @@ -782,7 +783,7 @@ def create_kernel_instance(self, kernel_source, kernel_options, params, verbose) ) # check for templated kernel - if kernel_source.lang in ["CUDA"] and "<" in name and ">" in name: + if kernel_source.lang == "PYCUDA" and "<" in name and ">" in name: kernel_string, name = wrap_templated_kernel(kernel_string, name) # Preprocess GPU arguments. Require for handling `Tunable` arguments From 20e5742673f28155e3185160d2aaf6f97671f5cf Mon Sep 17 00:00:00 2001 From: Ben van Werkhoven Date: Fri, 11 Sep 2026 14:17:58 +0200 Subject: [PATCH 3/4] add test to ensure failed compilation of SIMT kernels are detected correctly --- test/test_cuda_functions.py | 27 +++++++++++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/test/test_cuda_functions.py b/test/test_cuda_functions.py index 99989cc7..95208014 100644 --- a/test/test_cuda_functions.py +++ b/test/test_cuda_functions.py @@ -6,6 +6,7 @@ from kernel_tuner.core import KernelInstance, KernelSource from kernel_tuner.utils.nvcuda import cuda_error_check +from kernel_tuner import util from .context import skip_if_no_cuda from .test_runners import env # noqa: F401 @@ -152,3 +153,29 @@ def test_copy_constant_memory_args(): dev.memcpy_dtoh(output, gpu_args[0]) assert (my_constant_data == output).all() + + +@skip_if_no_cuda +def test_detect_kernel_too_much_shared_memory(env): + + # The kernel below uses a crazy amount of shared memory to trigger + # an error at compile time + kernel_string = """ + extern "C" __global__ void vector_add(float *c, float *a, float *b, int n) { + __shared__ float sh_mem_buf[block_size_x * 1000]; + int i = blockIdx.x * block_size_x + threadIdx.x; + if (i 0 + + for res in result: + assert '__error__' in res + assert isinstance(res['__error__'], util.CompilationFailedConfig) From e78a11a8fc1ff17f4e833c70670ddc21b28c93b7 Mon Sep 17 00:00:00 2001 From: Ben van Werkhoven Date: Fri, 11 Sep 2026 14:24:57 +0200 Subject: [PATCH 4/4] initial support for CUDA Tile C++ kernels --- examples/cuda-c++/vector_add_tile.cu | 26 +++ examples/cuda-c++/vector_add_tile.py | 35 ++++ kernel_tuner/backends/nvcuda.py | 266 ++++++++++++++++++--------- 3 files changed, 240 insertions(+), 87 deletions(-) create mode 100644 examples/cuda-c++/vector_add_tile.cu create mode 100644 examples/cuda-c++/vector_add_tile.py diff --git a/examples/cuda-c++/vector_add_tile.cu b/examples/cuda-c++/vector_add_tile.cu new file mode 100644 index 00000000..6742cd75 --- /dev/null +++ b/examples/cuda-c++/vector_add_tile.cu @@ -0,0 +1,26 @@ +#include "cuda_tile.h" + +__tile_global__ void vector_add_tile(float* a, float* b, float* out, int n) { + namespace ct = cuda::tiles; + using namespace ct::literals; + + a = ct::assume_aligned(a, 16_ic); + b = ct::assume_aligned(b, 16_ic); + out = ct::assume_aligned(out, 16_ic); + + // Step 1: attach a shape to each raw pointer. n is a runtime value (dynamic extent). + auto aSpan = ct::tensor_span{a, ct::extents{n}}; + auto bSpan = ct::tensor_span{b, ct::extents{n}}; + auto oSpan = ct::tensor_span{out, ct::extents{n}}; + + // Step 2: partition each span into tiles of TILE_SIZE elements (tuned by kernel_tuner). + constexpr auto tile = ct::integral_constant{}; + auto aView = ct::partition_view{aSpan, ct::shape{tile}}; + auto bView = ct::partition_view{bSpan, ct::shape{tile}}; + auto oView = ct::partition_view{oSpan, ct::shape{tile}}; + + int bx = ct::bid().x; // this block's tile-space index along .x + auto aTile = aView.load(bx); // pick the bx-th tile of a + auto bTile = bView.load(bx); + oView.store(aTile + bTile, bx); // write the tile back at the bx-th position of out +} diff --git a/examples/cuda-c++/vector_add_tile.py b/examples/cuda-c++/vector_add_tile.py new file mode 100644 index 00000000..0301156e --- /dev/null +++ b/examples/cuda-c++/vector_add_tile.py @@ -0,0 +1,35 @@ +#!/usr/bin/env python +""" This is a minimal example to tune a CUDA Tile vector add kernel """ + +import numpy +from kernel_tuner import tune_kernel + +def tune(): + + size = 3*2**24 + + a = numpy.random.randn(size).astype(numpy.float32) + b = numpy.random.randn(size).astype(numpy.float32) + c = numpy.zeros_like(b) + n = numpy.int32(size) + + args = [a, b, c, n] + + # TILE_SIZE controls how many elements each tile processes. kernel_tuner injects + # it as a #define so the kernel can use it as a compile-time constant. + # grid_x = size / TILE_SIZE (number of tiles), computed automatically by kernel_tuner. + tune_params = {"TILE_SIZE": [8, 16, 32, 64, 128, 256]} + + answer = [None, None, a+b, None] + + compiler_options = ["-enable-tile", "-std=c++20"] + + results, env = tune_kernel("vector_add_tile", "vector_add_tile.cu", size, args, + tune_params, lang="NVCUDA", compiler_options=compiler_options, + answer=answer, verbose=True, grid_div_x=["TILE_SIZE"]) + + return results + + +if __name__ == "__main__": + tune() diff --git a/kernel_tuner/backends/nvcuda.py b/kernel_tuner/backends/nvcuda.py index a22787c1..e535cab0 100644 --- a/kernel_tuner/backends/nvcuda.py +++ b/kernel_tuner/backends/nvcuda.py @@ -3,6 +3,8 @@ import numpy as np import uuid import os +import subprocess +import tempfile from kernel_tuner.backends.backend import GPUBackend from kernel_tuner.observers.nvcuda import CudaRuntimeObserver @@ -49,6 +51,12 @@ def preload_python_nvrtc(): except ImportError: driver = None +try: + from cuda.core import Kernel as CudaCoreKernel, launch as cuda_core_launch, LaunchConfig as CudaCoreConfig + _cuda_core_available = True +except ImportError: + _cuda_core_available = False + class CudaFunctions(GPUBackend): """Class that groups the Cuda functions and it maintains state about the device.""" @@ -102,6 +110,7 @@ def __init__(self, device=0, iterations=7, compiler_options=None, observers=None self.cc = f"{major}{minor}" self.iterations = iterations self.current_module = None + self.current_library = None self.func = None self.compiler_options = compiler_options or [] @@ -267,47 +276,50 @@ def compile(self, kernel_instance): """ kernel_string = kernel_instance.kernel_string kernel_name = kernel_instance.name - expression_name = str.encode(kernel_name) compiler_options = list(self.compiler_options) - # Add -std=c++11 - if not any(opt.startswith(("-std=", "--std=")) for opt in self.compiler_options): - compiler_options.append("--std=c++11") + # Detect Tile kernels: user passes "-enable-tile" (NVRTC flag) as a signal. + # Tile kernels are compiled with nvcc -tilecubin rather than NVRTC. + is_tile_kernel = any(str(opt).strip() == "-enable-tile" for opt in compiler_options) - # Add -arch - if not any(opt.startswith(("-arch", "--arch", "--gpu-architecture=")) for opt in self.compiler_options): - arch_val = to_valid_nvrtc_gpu_arch_cc(self.cc) - compiler_options.append(f"--gpu-architecture=compute_{arch_val}") + if is_tile_kernel: + if not _cuda_core_available: + raise RuntimeError("Tile kernels require 'cuda-core' (pip install cuda-core)") + self.func = self._compile_tile_kernel_nvcc(kernel_string, kernel_name, compiler_options) + self.num_regs = 0 + else: + expression_name = str.encode(kernel_name) - # Add CUDA home to include path - cuda_home = find_cuda_home() - if cuda_home: - cuda_include = os.path.join(cuda_home, "include") - compiler_options.append(f"-I{cuda_include}") + # Add -std=c++11 + if not any(opt.startswith(("-std=", "--std=")) for opt in self.compiler_options): + compiler_options.append("--std=c++11") - # nvrtcCompileProgram requires bytes instead of str - compiler_options = [str(opt).encode("UTF-8") for opt in compiler_options] + # Add -arch + if not any(opt.startswith(("-arch", "--arch", "--gpu-architecture=")) for opt in self.compiler_options): + arch_val = to_valid_nvrtc_gpu_arch_cc(self.cc) + compiler_options.append(f"--gpu-architecture=compute_{arch_val}") - err, program = nvrtc.nvrtcCreateProgram(str.encode(kernel_string), b"CUDAProgram", 0, [], []) - try: - cuda_error_check(err) - # Add the kernel as an expression. This is necessary for templated kernels to ensure that the - # compiler actually instantiates the kernel that we want to compile. - err = nvrtc.nvrtcAddNameExpression(program, expression_name) - cuda_error_check(err) + # Add CUDA home to include path + cuda_home = find_cuda_home() + if cuda_home: + cuda_include = os.path.join(cuda_home, "include") + compiler_options.append(f"-I{cuda_include}") - # Compile the program - err = nvrtc.nvrtcCompileProgram(program, len(compiler_options), compiler_options) - cuda_error_check(err) + # nvrtcCompileProgram requires bytes instead of str + compiler_options = [str(opt).encode("UTF-8") for opt in compiler_options] - if b"-enable-tile" in compiler_options: - # Get the Tile IR - err, size = nvrtc.nvrtcGetTileIRSize(program) + err, program = nvrtc.nvrtcCreateProgram(str.encode(kernel_string), b"CUDAProgram", 0, [], []) + try: cuda_error_check(err) - buff = b" " * size - err = nvrtc.nvrtcGetTileIR(program, buff) + # Add the kernel as an expression. This is necessary for templated kernels to ensure that the + # compiler actually instantiates the kernel that we want to compile. + err = nvrtc.nvrtcAddNameExpression(program, expression_name) cuda_error_check(err) - else: + + # Compile the program + err = nvrtc.nvrtcCompileProgram(program, len(compiler_options), compiler_options) + cuda_error_check(err) + # Get the PTX err, size = nvrtc.nvrtcGetPTXSize(program) cuda_error_check(err) @@ -315,60 +327,124 @@ def compile(self, kernel_instance): err = nvrtc.nvrtcGetPTX(program, buff) cuda_error_check(err) - # Load the module - err, self.current_module = _load_module_with_logs(np.char.array(buff)) - - # It the compile succeeded, but loading the PTX failed it is most likely - # that the kernel uses too much shared memory - if err == driver.CUresult.CUDA_ERROR_INVALID_PTX: - raise SkippableFailure("uses too much shared data") - else: - cuda_error_check(err) + # Load the module + err, self.current_module = _load_module_with_logs(buff) - # Get the lowered name (resolves C++ mangling for both regular and Tile kernels). - err, lowered_name = nvrtc.nvrtcGetLoweredName(program, expression_name) - cuda_error_check(err) + # If the compile succeeded but loading the PTX failed it is most likely + # that the kernel uses too much shared memory + if err == driver.CUresult.CUDA_ERROR_INVALID_PTX: + raise SkippableFailure("uses too much shared data") + else: + cuda_error_check(err) - if b"-enable-tile" in compiler_options: - # For Tile kernels the entry point name in the module is the lowered name - # plus encoded Tile template parameters. Enumerate all functions in the - # loaded module and find the one whose name contains the kernel name. - err, func_count = driver.cuModuleGetFunctionCount(self.current_module) - cuda_error_check(err) - err, functions = driver.cuModuleEnumerateFunctions(func_count, self.current_module) + # Get the lowered name (resolves C++ mangling) and look up the function + err, lowered_name = nvrtc.nvrtcGetLoweredName(program, expression_name) cuda_error_check(err) - self.func = None - for func in functions: - err, func_name = driver.cuFuncGetName(func) - if err != driver.CUresult.CUDA_SUCCESS: - continue - if isinstance(func_name, bytes): - func_name = func_name.decode() - if kernel_name in func_name: - self.func = func - break - if self.func is None: - raise RuntimeError(f"Could not find Tile kernel '{kernel_name}' in compiled module") - else: err, self.func = driver.cuModuleGetFunction(self.current_module, lowered_name) if err == driver.CUresult.CUDA_ERROR_NOT_FOUND: err, self.func = driver.cuModuleGetFunction(self.current_module, expression_name) cuda_error_check(err) - # get the number of registers per thread used in this kernel - num_regs = driver.cuFuncGetAttribute(driver.CUfunction_attribute.CU_FUNC_ATTRIBUTE_NUM_REGS, self.func) - assert num_regs[0] == 0, f"Retrieving number of registers per thread unsuccesful: code {num_regs[0]}" - self.num_regs = num_regs[1] + # get the number of registers per thread used in this kernel + num_regs = driver.cuFuncGetAttribute(driver.CUfunction_attribute.CU_FUNC_ATTRIBUTE_NUM_REGS, self.func) + assert num_regs[0] == 0, f"Retrieving number of registers per thread unsuccesful: code {num_regs[0]}" + self.num_regs = num_regs[1] - except RuntimeError as re: - _, n = nvrtc.nvrtcGetProgramLogSize(program) - log = b" " * n - nvrtc.nvrtcGetProgramLog(program, log) - print(log.decode("utf-8")) - raise re + 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 _compile_tile_kernel_nvcc(self, kernel_string, kernel_name, compiler_options): + """Compile a CUDA Tile kernel using nvcc -tilecubin and load it via cuda.core. + + Tile kernels cannot be compiled to PTX via NVRTC and loaded as regular + modules. Instead, nvcc produces a 'tile cubin' that can be loaded by + cuda.core.ObjectCode.from_cubin and launched with cuda.core.launch. + + Returns a cuda.core.Kernel that must be launched with block=1. + """ + from cuda.core import ObjectCode + + nvcc = os.environ.get("KERNEL_TUNER_NVCC", "nvcc") + + # Determine arch: translate --gpu-architecture=compute_XX or -arch=compute_XX + # to -arch=sm_XX (required for cubin output). + arch_flag = f"-arch=sm_{self.cc}" + for opt in compiler_options: + s = str(opt).strip() + for prefix in ("--gpu-architecture=compute_", "--gpu-architecture=sm_", + "-arch=compute_", "-arch=sm_"): + if s.startswith(prefix): + suffix = s[len(prefix):] + arch_flag = f"-arch=sm_{suffix}" + break + + # Build the nvcc command, forwarding safe options and skipping NVRTC-only ones. + _skip = {"-enable-tile"} + _skip_prefixes = ("--gpu-architecture=", "-arch=") + nvcc_opts = [] + for opt in compiler_options: + s = str(opt).strip() + if s in _skip or any(s.startswith(p) for p in _skip_prefixes): + continue + nvcc_opts.append(s) + + with tempfile.TemporaryDirectory() as tmpdir: + cu_file = os.path.join(tmpdir, "kernel.cu") + cubin_file = os.path.join(tmpdir, "kernel.cubin") + + with open(cu_file, "w") as f: + f.write(kernel_string) + + cmd = [nvcc, "-tilecubin", "--tile-only", arch_flag] + nvcc_opts + ["-o", cubin_file, cu_file] + try: + result = subprocess.run(cmd, capture_output=True, text=True, check=True) + except subprocess.CalledProcessError as e: + raise RuntimeError( + f"nvcc Tile kernel compilation failed:\n{e.stderr}" + ) from e + + cubin_bytes = open(cubin_file, "rb").read() + + # Find the mangled kernel name by loading the cubin as a module and + # enumerating its functions (same approach as TileGym). + err, mod = driver.cuModuleLoadData(cubin_bytes) + cuda_error_check(err) + err, func_count = driver.cuModuleGetFunctionCount(mod) + cuda_error_check(err) + err, functions = driver.cuModuleEnumerateFunctions(func_count, mod) + cuda_error_check(err) + # Strip template args (e.g. "vector_add_tile<8>" -> "vector_add_tile") before + # searching mangled names, since template args are encoded in mangled form. + base_name = kernel_name.split('<')[0].strip() + mangled_name = None + for func in functions: + err2, name_bytes = driver.cuFuncGetName(func) + if err2 != driver.CUresult.CUDA_SUCCESS: + continue + name = name_bytes.decode() if isinstance(name_bytes, bytes) else name_bytes + if base_name in name: + mangled_name = name + break + if mangled_name is None and func_count > 0: + _, name_bytes = driver.cuFuncGetName(functions[0]) + mangled_name = name_bytes.decode() if isinstance(name_bytes, bytes) else name_bytes + driver.cuModuleUnload(mod) + if mangled_name is None: + raise RuntimeError(f"Could not find Tile kernel '{kernel_name}' in compiled cubin") + + # Load the cubin via cuda.core and retrieve the Kernel object. + # Store the ObjectCode as self.current_library to keep it alive for the + # duration of the kernel's use (the Kernel handle is only valid while the + # ObjectCode/library is alive). + self.current_library = ObjectCode.from_cubin(cubin_bytes) + return self.current_library.get_kernel(mangled_name) + def start_event(self): """Records the event that marks the start of a measurement.""" err = runtime.cudaEventRecord(self.start, self.stream) @@ -450,20 +526,36 @@ def run_kernel(self, func, gpu_args, threads, grid, stream=None): else: arg_types.append(np.ctypeslib.as_ctypes_type(arg.dtype)) kernel_args = (tuple(gpu_args), tuple(arg_types)) - err = driver.cuLaunchKernel( - func, - grid[0], - grid[1], - grid[2], - threads[0], - threads[1], - threads[2], - self.smem_size, - stream, - kernel_args, - 0, - ) - cuda_error_check(err) + if _cuda_core_available and isinstance(func, CudaCoreKernel): + # Tile kernels are launched via cuda.core.launch with block=1; + # the Tile runtime handles the thread-to-tile mapping internally. + from cuda.core import Stream as CudaCoreStream + tile_args = [] + for arg in gpu_args: + if isinstance(arg, driver.CUdeviceptr): + tile_args.append(np.uint64(int(arg))) + else: + tile_args.append(arg) + # Stream.from_handle takes the raw integer CUstream handle value, + # the same convention as torch_stream.cuda_stream in PyTorch. + core_stream = CudaCoreStream.from_handle(int(stream)) + config = CudaCoreConfig(grid=(grid[0], grid[1], grid[2]), block=1, shmem_size=self.smem_size) + cuda_core_launch(core_stream, config, func, *tile_args) + else: + err = driver.cuLaunchKernel( + func, + grid[0], + grid[1], + grid[2], + threads[0], + threads[1], + threads[2], + self.smem_size, + stream, + kernel_args, + 0, + ) + cuda_error_check(err) @staticmethod def memset(allocation, value, size):