Added RegisterObserver with common interface among backends

Author: fjwillemsen

kernel_tuner/backends/cupy.py

@@ -132,6 +132,7 @@ def compile(self, kernel_instance):
         )
 
         self.func = self.current_module.get_function(kernel_name)
+        self.num_regs = self.func.num_regs
         return self.func
 
     def start_event(self):

kernel_tuner/backends/nvcuda.py

@@ -192,6 +192,11 @@ def compile(self, kernel_instance):
             )
             cuda_error_check(err)
 
+            # get the number of registers per thread used in this kernel
+            num_regs = cuda.cuFuncGetAttribute(cuda.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

kernel_tuner/backends/pycuda.py

@@ -218,6 +218,7 @@ def compile(self, kernel_instance):
             )
 
             self.func = self.current_module.get_function(kernel_name)
+            self.num_regs = self.func.num_regs
             return self.func
         except drv.CompileError as e:
             if "uses too much shared data" in e.stderr:

kernel_tuner/observers/register.py

@@ -0,0 +1,12 @@
+from kernel_tuner.observers.observer import BenchmarkObserver
+
+class RegisterObserver(BenchmarkObserver):
+    """Observer for counting the number of registers."""
+
+    def __init__(self) -> None:
+        super().__init__()
+
+    def get_results(self):
+        return {
+            "num_regs": self.dev.num_regs
+        }