Split backends and runtime observers.

Author: isazi

kernel_tuner/backends/c.py

@@ -13,7 +13,7 @@
 import numpy.ctypeslib
 
 from kernel_tuner.backends.backend import CompilerBackend
-from kernel_tuner.observers.observer import BenchmarkObserver
+from kernel_tuner.observers.c import CRuntimeObserver
 from kernel_tuner.util import get_temp_filename, delete_temp_file, write_file, SkippableFailure
 
 dtype_map = {"int8": C.c_int8,
@@ -33,26 +33,6 @@
 Argument = namedtuple("Argument", ["numpy", "ctypes"])
 
 
-class CRuntimeObserver(BenchmarkObserver):
-    """ Observer that collects results returned by benchmarking function in the C backend """
-
-    def __init__(self, dev):
-        self.dev = dev
-        self.objective = "time"
-        self.times = []
-
-    def after_finish(self):
-        self.times.append(self.dev.last_result)
-
-    def get_results(self):
-        results = {
-            self.objective: np.average(self.times),
-            self.objective + "s": self.times.copy()
-        }
-        self.times = []
-        return results
-
-
 class CFunctions(CompilerBackend):
     """Class that groups the code for running and compiling C functions"""
 

kernel_tuner/backends/cupy.py

@@ -7,7 +7,8 @@
 import numpy as np
 
 from kernel_tuner.backends.backend import GPUBackend
-from kernel_tuner.observers.observer import BenchmarkObserver
+from kernel_tuner.observers.cupy import CupyRuntimeObserver
+
 
 #embedded in try block to be able to generate documentation
 #and run tests without cupy installed
@@ -17,24 +18,6 @@
     cp = None
 
 
-class CupyRuntimeObserver(BenchmarkObserver):
-    """ Observer that measures time using CUDA events during benchmarking in the CuPy backend """
-    def __init__(self, dev):
-        self.dev = dev
-        self.stream = dev.stream
-        self.start = dev.start
-        self.end = dev.end
-        self.times = []
-
-    def after_finish(self):
-        self.times.append(cp.cuda.get_elapsed_time(self.start, self.end)) #ms
-
-    def get_results(self):
-        results = {"time": np.average(self.times), "times": self.times.copy()}
-        self.times = []
-        return results
-
-
 class CupyFunctions(GPUBackend):
     """Class that groups the Cupy functions on maintains state about the device"""
 

kernel_tuner/backends/nvcuda.py

@@ -2,8 +2,8 @@
 import numpy as np
 
 from kernel_tuner.backends.backend import GPUBackend
-from kernel_tuner.observers.observer import BenchmarkObserver
-from kernel_tuner.util import SkippableFailure
+from kernel_tuner.observers.nvcuda import CudaRuntimeObserver
+from kernel_tuner.util import SkippableFailure, cuda_error_check
 
 #embedded in try block to be able to generate documentation
 #and run tests without cuda-python installed
@@ -13,46 +13,6 @@
     cuda = None
 
 
-def error_check(error):
-    """ Checking the status of CUDA calls """
-    if isinstance(error, cuda.CUresult):
-        if error != cuda.CUresult.CUDA_SUCCESS:
-            _, name = cuda.cuGetErrorName(error)
-            raise RuntimeError(f"CUDA error: {name.decode()}")
-    elif isinstance(error, cudart.cudaError_t):
-        if error != cudart.cudaError_t.cudaSuccess:
-            _, name = cudart.getErrorName(error)
-            raise RuntimeError(f"CUDART error: {name.decode()}")
-    elif isinstance(error, nvrtc.nvrtcResult):
-        if error != nvrtc.nvrtcResult.NVRTC_SUCCESS:
-            _, desc = nvrtc.nvrtcGetErrorString(error)
-            raise RuntimeError(f"NVRTC error: {desc.decode()}")
-
-
-class CudaRuntimeObserver(BenchmarkObserver):
-    """ Observer that measures time using CUDA events during benchmarking """
-    def __init__(self, dev):
-        self.dev = dev
-        self.stream = dev.stream
-        self.start = dev.start
-        self.end = dev.end
-        self.times = []
-
-    def after_finish(self):
-        # time in ms
-        err, time = cudart.cudaEventElapsedTime(self.start, self.end)
-        error_check(err)
-        self.times.append(time)
-
-    def get_results(self):
-        results = {
-            "time": np.average(self.times),
-            "times": self.times.copy()
-        }
-        self.times = []
-        return results
-
-
 class CudaFunctions(GPUBackend):
     """Class that groups the Cuda functions on maintains state about the device"""
 
@@ -82,23 +42,23 @@ def __init__(self, device=0, iterations=7, compiler_options=None, observers=None
 
         # initialize and select device
         err = cuda.cuInit(0)
-        error_check(err)
+        cuda_error_check(err)
         err, self.device = cuda.cuDeviceGet(device)
-        error_check(err)
+        cuda_error_check(err)
         err, self.context = cuda.cuDevicePrimaryCtxRetain(device)
-        error_check(err)
+        cuda_error_check(err)
         if CudaFunctions.last_selected_device != device:
             err = cuda.cuCtxSetCurrent(self.context)
-            error_check(err)
+            cuda_error_check(err)
             CudaFunctions.last_selected_device = device
 
         # compute capabilities and device properties
         err, major = cudart.cudaDeviceGetAttribute(cudart.cudaDeviceAttr.cudaDevAttrComputeCapabilityMajor, device)
-        error_check(err)
+        cuda_error_check(err)
         err, minor = cudart.cudaDeviceGetAttribute(cudart.cudaDeviceAttr.cudaDevAttrComputeCapabilityMinor, device)
-        error_check(err)
+        cuda_error_check(err)
         err, self.max_threads = cudart.cudaDeviceGetAttribute(cudart.cudaDeviceAttr.cudaDevAttrMaxThreadsPerBlock, device)
-        error_check(err)
+        cuda_error_check(err)
         self.cc = f"{major}{minor}"
         self.iterations = iterations
         self.current_module = None
@@ -110,11 +70,11 @@ def __init__(self, device=0, iterations=7, compiler_options=None, observers=None
 
         # create a stream and events
         err, self.stream = cuda.cuStreamCreate(0)
-        error_check(err)
+        cuda_error_check(err)
         err, self.start = cuda.cuEventCreate(0)
-        error_check(err)
+        cuda_error_check(err)
         err, self.end = cuda.cuEventCreate(0)
-        error_check(err)
+        cuda_error_check(err)
 
         # default dynamically allocated shared memory size, can be overwritten using smem_args
         self.smem_size = 0
@@ -127,7 +87,7 @@ def __init__(self, device=0, iterations=7, compiler_options=None, observers=None
 
         # collect environment information
         err, device_properties = cudart.cudaGetDeviceProperties(device)
-        error_check(err)
+        cuda_error_check(err)
         env = dict()
         env["device_name"] = device_properties.name.decode()
         env["cuda_version"] = cuda.CUDA_VERSION
@@ -142,7 +102,7 @@ def __del__(self):
         for device_memory in self.allocations:
             if isinstance(device_memory, cuda.CUdeviceptr):
                 err = cuda.cuMemFree(device_memory)
-                error_check(err)
+                cuda_error_check(err)
 
     def ready_argument_list(self, arguments):
         """ready argument list to be passed to the kernel, allocates gpu mem
@@ -160,7 +120,7 @@ def ready_argument_list(self, arguments):
             # if arg is a numpy array copy it to device
             if isinstance(arg, np.ndarray):
                 err, device_memory = cuda.cuMemAlloc(arg.nbytes)
-                error_check(err)
+                cuda_error_check(err)
                 self.allocations.append(device_memory)
                 gpu_args.append(device_memory)
                 self.memcpy_htod(device_memory, arg)
@@ -201,21 +161,21 @@ def compile(self, kernel_instance):
 
         err, program = nvrtc.nvrtcCreateProgram(str.encode(kernel_string), b"CUDAProgram", 0, [], [])
         try:
-            error_check(err)
+            cuda_error_check(err)
             err = nvrtc.nvrtcCompileProgram(program, len(compiler_options), compiler_options)
-            error_check(err)
+            cuda_error_check(err)
             err, size = nvrtc.nvrtcGetPTXSize(program)
-            error_check(err)
+            cuda_error_check(err)
             buff = b' ' * size
             err = nvrtc.nvrtcGetPTX(program, buff)
-            error_check(err)
+            cuda_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)
+                cuda_error_check(err)
             err, self.func = cuda.cuModuleGetFunction(self.current_module, str.encode(kernel_name))
-            error_check(err)
+            cuda_error_check(err)
 
         except RuntimeError as re:
             _, n = nvrtc.nvrtcGetProgramLogSize(program)
@@ -229,12 +189,12 @@ def compile(self, kernel_instance):
     def start_event(self):
         """ Records the event that marks the start of a measurement """
         err = cudart.cudaEventRecord(self.start, self.stream)
-        error_check(err)
+        cuda_error_check(err)
 
     def stop_event(self):
         """ Records the event that marks the end of a measurement """
         err = cudart.cudaEventRecord(self.end, self.stream)
-        error_check(err)
+        cuda_error_check(err)
 
     def kernel_finished(self):
         """ Returns True if the kernel has finished, False otherwise """
@@ -248,7 +208,7 @@ def kernel_finished(self):
     def synchronize():
         """ Halts execution until device has finished its tasks """
         err = cudart.cudaDeviceSynchronize()
-        error_check(err)
+        cuda_error_check(err)
 
 
     def copy_constant_memory_args(self, cmem_args):
@@ -263,9 +223,9 @@ def copy_constant_memory_args(self, cmem_args):
         """
         for k, v in cmem_args.items():
             err, symbol, _ = cuda.cuModuleGetGlobal(self.current_module, str.encode(k))
-            error_check(err)
+            cuda_error_check(err)
             err = cuda.cuMemcpyHtoD(symbol, v, v.nbytes)
-            error_check(err)
+            cuda_error_check(err)
 
     def copy_shared_memory_args(self, smem_args):
         """add shared memory arguments to the kernel"""
@@ -307,7 +267,7 @@ def run_kernel(self, func, gpu_args, threads, grid, stream=None):
                 arg_types.append(np.ctypeslib.as_ctypes_type(arg.dtype))
         kernel_args  = (tuple(gpu_args), tuple(arg_types))
         err = cuda.cuLaunchKernel(func, grid[0], grid[1], grid[2], threads[0], threads[1], threads[2], self.smem_size, stream, kernel_args, 0)
-        error_check(err)
+        cuda_error_check(err)
 
     @staticmethod
     def memset(allocation, value, size):
@@ -324,7 +284,7 @@ def memset(allocation, value, size):
 
         """
         err = cudart.cudaMemset(allocation, value, size)
-        error_check(err)
+        cuda_error_check(err)
 
     @staticmethod
     def memcpy_dtoh(dest, src):
@@ -337,7 +297,7 @@ def memcpy_dtoh(dest, src):
         :type src: cuda.CUdeviceptr
         """
         err = cuda.cuMemcpyDtoH(dest, src, dest.nbytes)
-        error_check(err)
+        cuda_error_check(err)
 
     @staticmethod
     def memcpy_htod(dest, src):
@@ -350,7 +310,7 @@ def memcpy_htod(dest, src):
         :type src: numpy.ndarray
         """
         err = cuda.cuMemcpyHtoD(dest, src, src.nbytes)
-        error_check(err)
+        cuda_error_check(err)
 
     units = {'time': 'ms'}
 

kernel_tuner/backends/opencl.py

@@ -4,7 +4,7 @@
 import numpy as np
 
 from kernel_tuner.backends.backend import GPUBackend
-from kernel_tuner.observers.observer import BenchmarkObserver
+from kernel_tuner.observers.opencl import OpenCLObserver
 
 #embedded in try block to be able to generate documentation
 try:
@@ -13,22 +13,6 @@
     cl = None
 
 
-class OpenCLObserver(BenchmarkObserver):
-    """ Observer that measures time using CUDA events during benchmarking """
-    def __init__(self, dev):
-        self.dev = dev
-        self.times = []
-
-    def after_finish(self):
-        event = self.dev.event
-        self.times.append((event.profile.end - event.profile.start)*1e-6) #ms
-
-    def get_results(self):
-        results = {"time": np.average(self.times), "times": self.times.copy()}
-        self.times = []
-        return results
-
-
 class OpenCLFunctions(GPUBackend):
     """Class that groups the OpenCL functions on maintains some state about the device"""
 

kernel_tuner/backends/pycuda.py

@@ -6,7 +6,7 @@
 import numpy as np
 
 from kernel_tuner.backends.backend import GPUBackend
-from kernel_tuner.observers.observer import BenchmarkObserver
+from kernel_tuner.observers.pycuda import PyCudaRuntimeObserver
 from kernel_tuner.observers.nvml import nvml
 from kernel_tuner.util import TorchPlaceHolder, SkippableFailure
 
@@ -52,28 +52,6 @@ def get_pointer(self):
         return self.t.data_ptr()
 
 
-class PyCudaRuntimeObserver(BenchmarkObserver):
-    """ Observer that measures time using CUDA events during benchmarking """
-
-    def __init__(self, dev):
-        self.dev = dev
-        self.stream = dev.stream
-        self.start = dev.start
-        self.end = dev.end
-        self.times = []
-
-    def after_finish(self):
-        self.times.append(self.end.time_since(self.start))    #ms
-
-    def get_results(self):
-        results = {
-            "time": np.average(self.times),
-            "times": self.times.copy()
-        }
-        self.times = []
-        return results
-
-
 class PyCudaFunctions(GPUBackend):
     """Class that groups the CUDA functions on maintains state about the device"""
 

kernel_tuner/observers/c.py

@@ -0,0 +1,23 @@
+import numpy as np
+
+from kernel_tuner.observers.observer import BenchmarkObserver
+
+
+class CRuntimeObserver(BenchmarkObserver):
+    """Observer that collects results returned by benchmarking function in the C backend"""
+
+    def __init__(self, dev):
+        self.dev = dev
+        self.objective = "time"
+        self.times = []
+
+    def after_finish(self):
+        self.times.append(self.dev.last_result)
+
+    def get_results(self):
+        results = {
+            self.objective: np.average(self.times),
+            self.objective + "s": self.times.copy(),
+        }
+        self.times = []
+        return results