Merge pull request #196 from KernelTuner/add-pmt-observer

Add pmt observer

Author: benvanwerkhoven

examples/cuda/going_green_performance_model.py

@@ -33,7 +33,7 @@
     raise e
 
 from kernel_tuner.energy import energy
-from kernel_tuner.nvml import get_nvml_gr_clocks
+from kernel_tuner.observers.nvml import get_nvml_gr_clocks
 
 def get_default_parser():
     parser = argparse.ArgumentParser(
@@ -60,7 +60,7 @@ def get_default_parser():
     args = parser.parse_args()
 
     ridge_frequency, freqs, nvml_power, fitted_params, scaling = energy.create_power_frequency_model(device=args.device,
-                                                                                               n_samples=args.samples,
+                                                                                               n_samples=int(args.samples),
                                                                                                verbose=True,
                                                                                                nvidia_smi_fallback=args.nvidia_smi_fallback,
                                                                                                use_locked_clocks=args.locked_clocks)

examples/cuda/vector_add_observers_pmt.py

@@ -0,0 +1,49 @@
+#!/usr/bin/env python
+"""This is the minimal example from the README"""
+
+import json
+from collections import OrderedDict
+
+import numpy
+from kernel_tuner import tune_kernel
+from kernel_tuner.observers.pmt import PMTObserver
+
+def tune():
+
+    kernel_string = """
+    __global__ void vector_add(float *c, float *a, float *b, int n) {
+        int i = blockIdx.x * block_size_x + threadIdx.x;
+        if (i<n) {
+            c[i] = a[i] + b[i];
+        }
+    }
+    """
+
+    size = 80000000
+
+    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 = [c, a, b, n]
+
+    tune_params = dict()
+    tune_params["block_size_x"] = [128+64*i for i in range(15)]
+
+    pmtobserver = PMTObserver(["nvml", "rapl"])
+
+    metrics = OrderedDict()
+    metrics["GPU W"] = lambda p: p["nvml_power"]
+    metrics["CPU W"] = lambda p: p["rapl_power"]
+
+    results, env = tune_kernel("vector_add", kernel_string, size, args, tune_params, observers=[pmtobserver], metrics=metrics, iterations=32)
+
+    with open("vector_add.json", 'w') as fp:
+        json.dump(results, fp)
+
+    return results
+
+
+if __name__ == "__main__":
+    tune()

kernel_tuner/observers/pmt.py

@@ -0,0 +1,80 @@
+import numpy as np
+
+from kernel_tuner.observers.observer import BenchmarkObserver
+
+# check if pmt is installed
+try:
+    import pmt
+except ImportError:
+    pmt = None
+
+
+class PMTObserver(BenchmarkObserver):
+    """Observer that uses the PMT library to measure power
+
+    :param observables:  One of:
+        - A string specifying a single power meter to use
+        - A list of string, specifying one or more power meters to use
+        - A dictionary, specifying one or more power meters to use,
+          including the device identifier. For arduino this should be for
+          instance "/dev/ttyACM0". For nvml, it should correspond to the GPU
+          id (e.g. '0', or '1'). For some sensors (such as rapl) the device
+          id is not used, it should be 'None' in those cases.
+        This observer will report "<platform>_energy>" and "<platform>_power" for
+        all specified platforms.
+    :type observables: string,list/dictionary
+
+    """
+
+    def __init__(self, observable=None):
+        if not pmt:
+            raise ImportError("could not import pmt")
+
+        # User specifices a dictonary of platforms and corresponding device
+        if type(observable) is dict:
+            pass
+        elif type(observable) is list:
+            # user specifies a list of platforms as observable
+            observable = dict([(obs, 0) for obs in observable])
+        else:
+            # User specifices a string (single platform) as observable
+            observable = {observable: None}
+        supported = ["arduino", "jetson", "likwid", "nvml", "rapl", "rocm", "xilinx"]
+        for obs in observable.keys():
+            if not obs in supported:
+                raise ValueError(f"Observable {obs} not in supported: {supported}")
+
+        self.pms = [pmt.get_pmt(obs[0], obs[1]) for obs in observable.items()]
+        self.pm_names = list(observable.keys())
+
+        self.begin_states = [None] * len(self.pms)
+        self.initialize_results(self.pm_names)
+
+    def initialize_results(self, pm_names):
+        self.results = dict()
+        for pm_name in pm_names:
+            energy_result_name = f"{pm_name}_energy"
+            power_result_name = f"{pm_name}_power"
+            self.results[energy_result_name] = []
+            self.results[power_result_name] = []
+
+    def after_start(self):
+        self.begin_states = [pm.read() for pm in self.pms]
+
+    def after_finish(self):
+        end_states = [pm.read() for pm in self.pms]
+        for i in range(len(self.pms)):
+            begin_state = self.begin_states[i]
+            end_state = end_states[i]
+            measured_energy = pmt.joules(begin_state, end_state)
+            measured_power = pmt.watts(begin_state, end_state)
+            pm_name = self.pm_names[i]
+            energy_result_name = f"{pm_name}_energy"
+            power_result_name = f"{pm_name}_power"
+            self.results[energy_result_name].append(measured_energy)
+            self.results[power_result_name].append(measured_power)
+
+    def get_results(self):
+        averages = {key: np.average(values) for key, values in self.results.items()}
+        self.initialize_results(self.pm_names)
+        return averages