Merge branch 'master' into searchspace_experiments

Author: fjwillemsen

doc/requirements.txt

@@ -84,4 +84,4 @@ tzdata==2025.2 ; python_version >= "3.10" and python_version < "4"
 urllib3==2.4.0 ; python_version >= "3.10" and python_version < "4"
 wcwidth==0.2.13 ; python_version >= "3.10" and python_version < "4"
 webencodings==0.5.1 ; python_version >= "3.10" and python_version < "4"
-xmltodict==0.14.2 ; python_version >= "3.10" and python_version < "4"
+xmltodict==0.14.2 ; python_version >= "3.10" and python_version < "4"

kernel_tuner/backends/hip.py

@@ -19,7 +19,6 @@
     "bool": ctypes.c_bool,
     "int8": ctypes.c_int8,
     "int16": ctypes.c_int16,
-    "float16": ctypes.c_int16,
     "int32": ctypes.c_int32,
     "int64": ctypes.c_int64,
     "uint8": ctypes.c_uint8,
@@ -40,7 +39,9 @@ def hip_check(call_result):
     if len(result) == 1:
         result = result[0]
     if isinstance(err, hip.hipError_t) and err != hip.hipError_t.hipSuccess:
-        raise RuntimeError(str(err))
+        _, error_name = hip.hipGetErrorName(err)
+        _, error_str = hip.hipGetErrorString(err)
+        raise RuntimeError(f"{error_name}: {error_str}")
     return result
 
 
@@ -120,25 +121,29 @@ def ready_argument_list(self, arguments):
 
             # Handle numpy arrays
             if isinstance(arg, np.ndarray):
-                if dtype_str in dtype_map.keys():
-                    # Allocate device memory
-                    device_ptr = hip_check(hip.hipMalloc(arg.nbytes))
+                # Allocate device memory
+                device_ptr = hip_check(hip.hipMalloc(arg.nbytes))
 
-                    # Copy data to device using hipMemcpy
-                    hip_check(hip.hipMemcpy(device_ptr, arg, arg.nbytes, hip.hipMemcpyKind.hipMemcpyHostToDevice))
+                # Copy data to device using hipMemcpy
+                hip_check(hip.hipMemcpy(device_ptr, arg, arg.nbytes, hip.hipMemcpyKind.hipMemcpyHostToDevice))
 
-                    prepared_args.append(device_ptr)
-                else:
-                    raise TypeError(f"Unknown dtype {dtype_str} for ndarray")
+                prepared_args.append(device_ptr)
 
             # Handle numpy scalar types
             elif isinstance(arg, np.generic):
                 # Convert numpy scalar to corresponding ctypes
-                ctype_arg = dtype_map[dtype_str](arg)
-                prepared_args.append(ctype_arg)
+                if dtype_str in dtype_map:
+                    ctype_arg = dtype_map[dtype_str](arg)
+                    prepared_args.append(ctype_arg)
+                # 16-bit float is not supported, view it as uint16
+                elif dtype_str in ("float16", "bfloat16"):
+                    ctype_arg = ctypes.c_uint16(arg.view(np.uint16))
+                    prepared_args.append(ctype_arg)
+                else:
+                    raise ValueError(f"Invalid argument type {dtype_str}: {arg}")
 
             else:
-                raise ValueError(f"Invalid argument type {type(arg)}, {arg}")
+                raise ValueError(f"Invalid argument type {type(arg)}: {arg}")
 
         return prepared_args
 

kernel_tuner/core.py

@@ -506,7 +506,7 @@ def check_kernel_output(
         # run the kernel
         check = self.run_kernel(func, gpu_args, instance)
         if not check:
-            # runtime failure occured that should be ignored, skip correctness check
+            # runtime failure occurred that should be ignored, skip correctness check
             return
 
         # retrieve gpu results to host memory
@@ -905,7 +905,7 @@ def split_argument_list(argument_list):
         match = re.match(regex, arg, re.S)
         if not match:
             raise ValueError("error parsing templated kernel argument list")
-        type_list.append(re.sub(r"\s+", " ", match.group(1).strip(), re.S))
+        type_list.append(re.sub(r"\s+", " ", match.group(1).strip(), flags=re.S))
         name_list.append(match.group(2).strip())
     return type_list, name_list
 

kernel_tuner/util.py

@@ -289,7 +289,7 @@ def check_restriction(restrict, params: dict) -> bool:
             return restrict(**selected_params)
     # otherwise, raise an error
     else:
-        raise ValueError(f"Unkown restriction type {type(restrict)} ({restrict})")
+        raise ValueError(f"Unknown restriction type {type(restrict)} ({restrict})")
 
 
 def check_restrictions(restrictions, params: dict, verbose: bool) -> bool:
@@ -357,7 +357,7 @@ def f_restrict(p):
 
     elif isinstance(restrict, (InSetConstraint, NotInSetConstraint, SomeInSetConstraint, SomeNotInSetConstraint)):
         raise NotImplementedError(
-            f"Restriction of the type {type(restrict)} is explicitely not supported in backwards compatibility mode, because the behaviour is too complex. Please rewrite this constraint to a function to use it with this algorithm."
+            f"Restriction of the type {type(restrict)} is explicitly not supported in backwards compatibility mode, because the behaviour is too complex. Please rewrite this constraint to a function to use it with this algorithm."
         )
     else:
         raise TypeError(f"Unrecognized restriction {restrict}")
@@ -602,7 +602,7 @@ def get_total_timings(results, env, overhead_time):
             total_verification_time += result["verification_time"]
             total_benchmark_time += result["benchmark_time"]
 
-    # add the seperate times to the environment dict
+    # add the separate time values to the environment dict
     env["total_framework_time"] = total_framework_time
     env["total_strategy_time"] = total_strategy_time
     env["total_compile_time"] = total_compile_time
@@ -778,7 +778,7 @@ def prepare_kernel_string(kernel_name, kernel_string, params, grid, threads, blo
             v = replace_param_occurrences(v, params)
 
         if not k.isidentifier():
-            raise ValueError("name is not a valid identifier: {k}")
+            raise ValueError(f"name is not a valid identifier: {k}")
 
         # Escape newline characters
         v = str(v)