Made function for scaling the compute capability to a valid one, added tests for this function, removed setting --gpu-architecture for CuPy as it is already set internally

Author: fjwillemsen

kernel_tuner/backends/cupy.py

@@ -6,7 +6,6 @@
 
 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
@@ -127,11 +126,7 @@ 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 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")
+        # CuPy already sets the --gpu-architecture by itself, as per https://github.com/cupy/cupy/blob/20ccd63c0acc40969c851b1917dedeb032209e8b/cupy/cuda/compiler.py#L145
 
         options = tuple(compiler_options)
 

kernel_tuner/backends/nvcuda.py

@@ -5,7 +5,7 @@
 
 from kernel_tuner.backends.backend import GPUBackend
 from kernel_tuner.observers.nvcuda import CudaRuntimeObserver
-from kernel_tuner.util import SkippableFailure, cuda_error_check, is_valid_nvrtc_gpu_arch_cc
+from kernel_tuner.util import SkippableFailure, cuda_error_check, to_valid_nvrtc_gpu_arch_cc
 
 # embedded in try block to be able to generate documentation
 # and run tests without cuda-python installed
@@ -167,15 +167,12 @@ 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 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")
+        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_{to_valid_nvrtc_gpu_arch_cc(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_{to_valid_nvrtc_gpu_arch_cc(self.cc)}")
 
         err, program = nvrtc.nvrtcCreateProgram(
             str.encode(kernel_string), b"CUDAProgram", 0, [], []

kernel_tuner/util.py

@@ -221,7 +221,7 @@ def check_block_size_names(block_size_names):
         if not isinstance(block_size_names, list):
             raise ValueError("block_size_names should be a list of strings!")
         if len(block_size_names) > 3:
-            raise ValueError("block_size_names should not contain more than 3 names!")
+            raise ValueError(f"block_size_names should not contain more than 3 names! ({block_size_names=})")
         if not all([isinstance(name, "".__class__) for name in block_size_names]):
             raise ValueError("block_size_names should contain only strings!")
 
@@ -570,10 +570,22 @@ 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 to_valid_nvrtc_gpu_arch_cc(compute_capability: str) -> str:
+    """Returns a valid Compute Capability 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']    # must be in ascending order, when updating also update test_to_valid_nvrtc_gpu_arch_cc
+    compute_capability = str(compute_capability)
+    if len(compute_capability) < 2:
+        raise ValueError(f"Compute capability '{compute_capability}' must be at least of length 2, is {len(compute_capability)}")
+    if compute_capability in valid_cc:
+        return compute_capability
+    # if the compute capability does not match, scale down to the nearest matching
+    subset_cc = [cc for cc in valid_cc if compute_capability[0] == cc[0]]
+    if len(subset_cc) > 0:
+        # get the next-highest valid CC
+        highest_cc_index = max([i for i, cc in enumerate(subset_cc) if int(cc[1]) <= int(compute_capability[1])])
+        return subset_cc[highest_cc_index]
+    # if all else fails, return the default 52
+    return '52'
 
 
 def print_config(config, tuning_options, runner):

test/test_util_functions.py

@@ -146,6 +146,19 @@ def test_get_thread_block_dimensions():
     assert threads[2] == 1
 
 
+def test_to_valid_nvrtc_gpu_arch_cc():
+    assert to_valid_nvrtc_gpu_arch_cc("89") == "89"
+    assert to_valid_nvrtc_gpu_arch_cc("88") == "87"
+    assert to_valid_nvrtc_gpu_arch_cc("86") == "80"
+    assert to_valid_nvrtc_gpu_arch_cc("40") == "52"
+    assert to_valid_nvrtc_gpu_arch_cc("90b") == "90a"
+    assert to_valid_nvrtc_gpu_arch_cc("91c") == "90a"
+    assert to_valid_nvrtc_gpu_arch_cc("10123001") == "52"
+    with pytest.raises(ValueError):
+        assert to_valid_nvrtc_gpu_arch_cc("")
+        assert to_valid_nvrtc_gpu_arch_cc("1")
+
+
 def test_prepare_kernel_string():
     kernel = "this is a weird kernel"
     grid = (3, 7)