Added function to check for compute capability validity, improved check on gpu-architecture compiler option, added gpu-architecture auto-adding to CuPy

Author: fjwillemsen

kernel_tuner/backends/cupy.py

@@ -1,10 +1,12 @@
 """This module contains all Cupy specific kernel_tuner functions."""
 from __future__ import print_function
+from warnings import warn
 
 import numpy as np
 
 from kernel_tuner.backends.backend import GPUBackend
 from kernel_tuner.observers.cupy import CupyRuntimeObserver
+from kernel_tuner.util import is_valid_nvrtc_gpu_arch_cc
 
 # embedded in try block to be able to generate documentation
 # and run tests without cupy installed
@@ -125,10 +127,11 @@ def compile(self, kernel_instance):
         compiler_options = self.compiler_options
         if not any(["-std=" in opt for opt in self.compiler_options]):
             compiler_options = ["--std=c++11"] + self.compiler_options
-        if not any([b"--gpu-architecture=" in opt for opt in compiler_options]):
-            compiler_options.append(
-                f"--gpu-architecture=compute_{self.cc}".encode("UTF-8")
-            )
+        if is_valid_nvrtc_gpu_arch_cc(self.cc):
+            if not any(["--gpu-architecture=" in opt or "-arch" in opt for opt in compiler_options]):
+                compiler_options.append(f"--gpu-architecture=compute_{self.cc}")
+        else:
+            warn(f"Could not add compiler option '--gpu-architecture=compute_{self.cc}' as {self.cc} is an invalid target")
 
         options = tuple(compiler_options)
 

kernel_tuner/backends/nvcuda.py

@@ -1,9 +1,11 @@
 """This module contains all NVIDIA cuda-python specific kernel_tuner functions."""
+from warnings import warn
+
 import numpy as np
 
 from kernel_tuner.backends.backend import GPUBackend
 from kernel_tuner.observers.nvcuda import CudaRuntimeObserver
-from kernel_tuner.util import SkippableFailure, cuda_error_check
+from kernel_tuner.util import SkippableFailure, cuda_error_check, is_valid_nvrtc_gpu_arch_cc
 
 # embedded in try block to be able to generate documentation
 # and run tests without cuda-python installed
@@ -165,12 +167,15 @@ def compile(self, kernel_instance):
             compiler_options.append(b"--std=c++11")
         if not any(["--std=" in opt for opt in self.compiler_options]):
             self.compiler_options.append("--std=c++11")
-        if not any([b"--gpu-architecture=" in opt for opt in compiler_options]):
-            compiler_options.append(
-                f"--gpu-architecture=compute_{self.cc}".encode("UTF-8")
-            )
-        if not any(["--gpu-architecture=" in opt for opt in self.compiler_options]):
-            self.compiler_options.append(f"--gpu-architecture=compute_{self.cc}")
+        if is_valid_nvrtc_gpu_arch_cc(self.cc):
+            if not any([b"--gpu-architecture=" in opt or b"-arch" in opt for opt in compiler_options]):
+                compiler_options.append(
+                    f"--gpu-architecture=compute_{self.cc}".encode("UTF-8")
+                )
+            if not any(["--gpu-architecture=" in opt or "-arch" in opt for opt in self.compiler_options]):
+                self.compiler_options.append(f"--gpu-architecture=compute_{self.cc}")
+        else:
+            warn(f"Could not add compiler option '--gpu-architecture=compute_{self.cc}' as {self.cc} is an invalid target")
 
         err, program = nvrtc.nvrtcCreateProgram(
             str.encode(kernel_string), b"CUDAProgram", 0, [], []

kernel_tuner/util.py

@@ -570,6 +570,12 @@ def get_total_timings(results, env, overhead_time):
     return env
 
 
+def is_valid_nvrtc_gpu_arch_cc(compute_capability: str) -> bool:
+    """Returns whether the Compute Capability is a valid argument for NVRTC `--gpu-architecture=`, as per https://docs.nvidia.com/cuda/nvrtc/index.html#group__options."""
+    valid_cc = ['50', '52', '53', '60', '61', '62', '70', '72', '75', '80', '87', '89', '90', '90a']
+    return str(compute_capability) in valid_cc
+
+
 def print_config(config, tuning_options, runner):
     """Print the configuration string with tunable parameters and benchmark results."""
     print_config_output(tuning_options.tune_params, config, runner.quiet, tuning_options.metrics, runner.units)