modified run_kernel (modifies grpu_args from ctypes list to ctypes structure), ready_argument_list (returns ctypes list), and memcpy_dtoh; examples/hip/test_vector_add.py passes -> memcpy_dtoh works

Author: MiloLurati

examples/hip/test_vector_add.py

@@ -5,13 +5,15 @@
 from kernel_tuner import run_kernel
 import pytest
 
+#Check pyhip is installed and if a HIP capable device is present, if not skip the test
+try:
+    from pyhip import hip, hiprtc
+except ImportError:
+    pytest.skip("PyHIP not installed or PYTHONPATH does not includes PyHIP")
+    hip = None
+    hiprtc = None
+
 def test_vector_add():
-    #Check pyhip is installed and if a HIP capable device is present, if not skip the test
-    try:
-        import pyhip as hip
-        hip.hipGetDeviceProperties(0)
-    except (ImportError, Exception):
-        pytest.skip("PyHIP not installed or no HIP device detected")
 
     kernel_string = """
     __global__ void vector_add(float *c, float *a, float *b, int n) {

kernel_tuner/backends/hip.py

@@ -128,8 +128,8 @@ def ready_argument_list(self, arguments):
             The order should match the argument list on the HIP function.
             Allowed values are np.ndarray, and/or np.int32, np.float32, and so on.
         :type arguments: list(numpy objects)
-        :returns: A ctypes structure that can be passed to the HIP function.
-        :rtype: ctypes.Structure
+        :returns: List of ctypes arguments to be passed to the HIP function.
+        :rtype: list of ctypes
         """
         logging.debug("HipFunction ready_argument_list called")
 
@@ -151,13 +151,7 @@ def ready_argument_list(self, arguments):
                 data_ctypes = dtype_map[dtype_str](arg)
                 ctype_args.append(data_ctypes)  
 
-        # Determine the types of the fields in the structure
-        field_types = [type(x) for x in ctype_args]
-        # Define a new ctypes structure with the inferred layout
-        class ArgListStructure(ctypes.Structure):
-            _fields_ = [(f'field{i}', t) for i, t in enumerate(field_types)]
-        
-        return ArgListStructure(*ctype_args)
+        return ctype_args
 
 
     def compile(self, kernel_instance):
@@ -228,7 +222,7 @@ def run_kernel(self, func, gpu_args, threads, grid, stream=None):
         :param gpu_args: A list of arguments to the kernel, order should match the
             order in the code. Allowed values are either variables in global memory
             or single values passed by value.
-        :type gpu_args: ctypes.Structure
+        :type gpu_args: list of ctypes
 
         :param threads: A tuple listing the number of threads in each dimension of
             the thread block
@@ -241,6 +235,15 @@ def run_kernel(self, func, gpu_args, threads, grid, stream=None):
         logging.debug("HipFunction run_kernel called")
         if stream is None:
             stream = self.stream
+
+        # Determine the types of the fields in the structure
+        field_types = [type(x) for x in gpu_args]
+        # Define a new ctypes structure with the inferred layout
+        class ArgListStructure(ctypes.Structure):
+            _fields_ = [(f'field{i}', t) for i, t in enumerate(field_types)]
+        
+        gpu_args = ArgListStructure(*gpu_args)
+
         hip.hipModuleLaunchKernel(func, 
                                   grid[0], grid[1], grid[2], 
                                   threads[0], threads[1], threads[2],
@@ -261,6 +264,8 @@ def memset(self, allocation, value, size):
         :type size: int
 
         """
+        logging.debug("HipFunction memset called")
+        print("HipFunction memset called")
         ctypes_value = ctypes.c_int(value)
         ctypes_size = ctypes.c_size_t(size)
         status = _libhip.hipMemset(allocation, ctypes_value, ctypes_size)
@@ -276,8 +281,10 @@ def memcpy_dtoh(self, dest, src):
         :type src: ctypes ptr
         """
         logging.debug("HipFunction memcpy_dtoh called")
-        dtype_str = str(src.dtype)
-        hip.hipMemcpy_dtoh(ctypes.byref(dest.ctypes), src, ctypes.sizeof(dtype_map[dtype_str]) * src.size)
+        print("HipFunction memcpy_dtoh called")
+
+        address = dest.ctypes.data
+        hip.hipMemcpy_dtoh(ctypes.c_void_p(address), src, dest.nbytes)
 
     def memcpy_htod(self, dest, src):
         """perform a host to device memory copy
@@ -289,6 +296,7 @@ def memcpy_htod(self, dest, src):
         :type src: numpy.ndarray
         """
         logging.debug("HipFunction memcpy_htod called")
+        print("HipFunction memcpy_htod called")
         dtype_str = str(src.dtype)
         hip.hipMemcpy_htod(dest, ctypes.byref(src.ctypes), ctypes.sizeof(dtype_map[dtype_str]) * src.size)
 
@@ -303,6 +311,7 @@ def copy_constant_memory_args(self, cmem_args):
         :type cmem_args: dict( string: numpy.ndarray, ... )
         """
         logging.debug("HipFunction copy_constant_memory_args called")
+        print("HipFunction copy_constant_memory_args called")
         logging.debug("current module: " + str(self.current_module))
 
         for k, v in cmem_args.items():
@@ -316,11 +325,13 @@ def copy_constant_memory_args(self, cmem_args):
     def copy_shared_memory_args(self, smem_args):
         """add shared memory arguments to the kernel"""
         logging.debug("HipFunction copy_shared_memory_args called")
+        print("HipFunction copy_shared_memory_args called")
         self.smem_size = smem_args["size"]
 
     def copy_texture_memory_args(self, texmem_args):
         """This method must implement the allocation and copy of texture memory to the GPU."""
         logging.debug("HipFunction copy_texture_memory_args called")
+        print("HipFunction copy_texture_memory_args called")
         raise NotImplementedError("HIP backend does not support texture memory") # NOT SUPPORTED?
 
     units = {"time": "ms"}