got the code working again

Author: benvanwerkhoven

examples/cuda/going_green_performance_model.py

@@ -1,3 +1,4 @@
+#!/usr/bin/env python
 """
 This example demonstrates how to use the performance model presented in
 
@@ -58,27 +59,31 @@ def get_default_parser():
     parser = get_default_parser()
     args = parser.parse_args()
 
-    ridge_frequency, fitted_params, scaling = energy.create_performance_frequency_model(device=args.device,
-                                                                                        n_samples=args.samples,
-                                                                                        verbose=True,
-                                                                                        nvidia_smi_fallback=args.nvidia_smi_fallback,
-                                                                                        use_locked_clocks=args.locked_clocks)
+    ridge_frequency, freqs, nvml_power, fitted_params, scaling = energy.create_performance_frequency_model(device=args.device,
+                                                                                               n_samples=args.samples,
+                                                                                               verbose=True,
+                                                                                               nvidia_smi_fallback=args.nvidia_smi_fallback,
+                                                                                               use_locked_clocks=args.locked_clocks)
 
     all_frequencies = np.array(get_nvml_gr_clocks(args.device)['nvml_gr_clock'])
 
-    frequency_selection = energy.get_frequency_range_around_ridge(all_frequencies, args.range, args.number)
+    frequency_selection = energy.get_frequency_range_around_ridge(ridge_frequency, all_frequencies, args.range, args.number)
     print(f"Search space reduction: {np.round(100 - len(frequency_selection) / len(all_frequencies) * 100, 1)} %%")
 
     xs = np.linspace(all_frequencies[0], all_frequencies[-1], 100)
     # scale to start at 0
     xs -= scaling[0]
-    modelled_power = estimated_power(xs, *fitted_params)
+    modelled_power = energy.estimated_power(xs, *fitted_params)
     # undo scaling
     xs += scaling[0]
     modelled_power *= scaling[1]
 
     # Add point for ridge frequency
-    P_ridge = estimated_power([ridge_frequency - scaling[0]], *fitted_params) * scaling[1]
+    P_ridge = energy.estimated_power([ridge_frequency - scaling[0]], *fitted_params) * scaling[1]
+
+    # Add the frequency range
+    min_freq = 1e-2 * (100 - int(args.range)) * ridge_frequency
+    max_freq = 1e-2 * (100 + int(args.range)) * ridge_frequency
 
     # plot measurements with model
     try:
@@ -101,5 +106,6 @@ def get_default_parser():
     plt.xlabel('Core frequency (MHz)')
     plt.ylabel('Power consumption (W)')
     plt.legend()
+    plt.show()
 
     plt.savefig("GPU_power_consumption_model.pdf")

kernel_tuner/energy/energy.py

@@ -2,12 +2,18 @@
 This module contains a set of helper functions specifically for auto-tuning codes
 for energy efficiency.
 """
+from collections import OrderedDict
 import numpy as np
 import math
 from scipy import optimize
 
 from kernel_tuner import tune_kernel
-from kernel_tuner.nvml import nvml, NVMLObserver
+from kernel_tuner.nvml import NVMLObserver, get_nvml_gr_clocks, get_idle_power
+
+try:
+    import pycuda.driver as drv
+except ImportError:
+    pass
 
 fp32_kernel_string = """
 __device__ void fp32_n_8(
@@ -36,7 +42,7 @@
 }
 """
 
-def get_frequency_power_relation_fp32(device, n_samples=10, use_locked_clocks=False, nvidia_smi_fallback=None):
+def get_frequency_power_relation_fp32(device, n_samples=10, nvidia_smi_fallback=None, use_locked_clocks=False):
     """ Use NVML and PyCUDA with a synthetic kernel to obtain samples of frequency-power pairs """
 
     if drv is None:
@@ -46,17 +52,17 @@ def get_frequency_power_relation_fp32(device, n_samples=10, use_locked_clocks=Fa
     drv.init()
     dev = drv.Device(device)
     device_name = dev.name().replace(' ', '_')
-    multiprocessor_count = dev.get_attribute(
-        drv.device_attribute.MULTIPROCESSOR_COUNT)
-    max_block_dim_x = dev.get_attribute(drv.device_attribute.MAX_BLOCK_DIM_X)
+    multiprocessor_count = int(dev.get_attribute(
+        drv.device_attribute.MULTIPROCESSOR_COUNT))
+    max_block_dim_x = int(dev.get_attribute(drv.device_attribute.MAX_BLOCK_DIM_X))
 
     # kernel arguments
     data_size = (multiprocessor_count, max_block_dim_x)
     data = np.random.random(np.prod(data_size)).astype(float)
     arguments = [data]
 
     # setup clocks
-    nvml_gr_clocks = get_nvml_gr_clocks(device, n=n_samples)
+    nvml_gr_clocks = get_nvml_gr_clocks(device, n=n_samples, quiet=True)
 
     # idle power
     power_idle = get_idle_power(device)
@@ -68,17 +74,19 @@ def get_frequency_power_relation_fp32(device, n_samples=10, use_locked_clocks=Fa
     tune_params["nr_inner"] = [1024]
     tune_params.update(nvml_gr_clocks)
 
+    tune_params["nvml_gr_clock"] = [int(c) for c in tune_params["nvml_gr_clock"]]
+
     # metrics
     metrics = OrderedDict()
     metrics["f"] = lambda p: p["core_freq"]
 
     nvmlobserver = NVMLObserver(
-        ["core_freq", "nvml_power"], device=device, nvidia_smi_fallback=nvidia_smi_fallback)
+        ["core_freq", "nvml_power"], device=device, nvidia_smi_fallback=nvidia_smi_fallback, use_locked_clocks=use_locked_clocks)
 
     results, _ = tune_kernel("fp32_kernel", fp32_kernel_string, problem_size=(multiprocessor_count, 64),
                              arguments=arguments, tune_params=tune_params, observers=[nvmlobserver],
                              verbose=False, quiet=True, metrics=metrics, iterations=10,
-                             grid_div_x=[], grid_div_y=[])
+                             grid_div_x=[], grid_div_y=[], cache=f"synthetic_fp32_cache_{device_name}.json")
 
     freqs = np.array([res["core_freq"] for res in results])
     nvml_power = np.array([res["nvml_power"] for res in results])
@@ -172,27 +180,27 @@ def create_performance_frequency_model(device=0, n_samples=10, verbose=False, nv
     :rtype: float
 
     """
-    freqs, nvml_power = get_frequency_power_relation(device, n_samples, nvidia_smi_fallback, use_locked_clocks)
+    freqs, nvml_power = get_frequency_power_relation_fp32(device, n_samples, nvidia_smi_fallback, use_locked_clocks)
 
     if verbose:
         print("Clock frequencies:", freqs.tolist())
         print("Power consumption:", nvml_power.tolist())
 
-    ridge_frequency, fitted_params, scaling = fit_model(freqs, nvml_power)
+    ridge_frequency, fitted_params, scaling = fit_performance_frequency_model(freqs, nvml_power)
 
     if verbose:
         print(f"Modelled most energy efficient frequency: {ridge_frequency} MHz")
 
-    all_frequencies = np.array(get_nvml_gr_clocks(device)['nvml_gr_clock'])
+    all_frequencies = np.array(get_nvml_gr_clocks(device, quiet=True)['nvml_gr_clock'])
     ridge_frequency_final = all_frequencies[np.argmin(abs(all_frequencies - ridge_frequency))]
 
     if verbose:
-        print(f"Closest configurable most energy efficient frequency: {ridge_frequency2} MHz")
+        print(f"Closest configurable most energy efficient frequency: {ridge_frequency_final} MHz")
 
-    return ridge_frequency_final, fitted_params, scaling
+    return ridge_frequency_final, freqs, nvml_power, fitted_params, scaling
 
 
-def get_frequency_range_around_ridge(ridge_frequency, all_frequencies, freq_range, number_of_freqs, verbose=False)
+def get_frequency_range_around_ridge(ridge_frequency, all_frequencies, freq_range, number_of_freqs, verbose=False):
     """ Return number_of_freqs frequencies in a freq_range percentage around the ridge_frequency from among all_frequencies """
 
     min_freq = 1e-2 * (100 - int(freq_range)) * ridge_frequency

kernel_tuner/nvml.py

@@ -2,6 +2,7 @@
 import time
 import re
 import numpy as np
+from collections import OrderedDict
 
 from kernel_tuner.observers import BenchmarkObserver, ContinuousObserver
 
@@ -39,6 +40,7 @@ def __init__(self, device_id=0, nvidia_smi_fallback='nvidia-smi', use_locked_clo
             self._auto_boost = None
 
         #try to initialize application clocks
+        self.modified_clocks = False
         try:
             if not use_locked_clocks:
                 self.gr_clock_default = pynvml.nvmlDeviceGetDefaultApplicationsClock(self.dev, pynvml.NVML_CLOCK_GRAPHICS)
@@ -75,7 +77,8 @@ def __del__(self):
         #try to restore to defaults
         if self.pwr_limit_default is not None:
             self.pwr_limit = self.pwr_limit_default
-        self.reset_clocks()
+        if self.modified_clocks:
+            self.reset_clocks()
 
     @property
     def pwr_state(self):
@@ -115,6 +118,7 @@ def persistence_mode(self, new_mode):
 
     def set_clocks(self, mem_clock, gr_clock):
         """Set the memory and graphics clock for this device (may require permission)"""
+        self.modified_clocks = True
         if not mem_clock in self.supported_mem_clocks:
             raise ValueError("Illegal value for memory clock")
         if not gr_clock in self.supported_gr_clocks[mem_clock]: