modified copy_constant_memory_args, implemented test_copy_constant_memory_args --> passed

Author: MiloLurati

kernel_tuner/backends/hip.py

@@ -68,9 +68,11 @@
 _libhip.hipEventQuery.restype = ctypes.c_int
 _libhip.hipEventQuery.argtypes = [ctypes.c_void_p]
 _libhip.hipModuleGetGlobal.restype = ctypes.c_int
-_libhip.hipModuleGetGlobal.argtypes = [ctypes.c_void_p, ctypes.c_size_t, ctypes.c_void_p, ctypes.c_char_p]
+_libhip.hipModuleGetGlobal.argtypes = [ctypes.POINTER(ctypes.c_void_p), ctypes.POINTER(ctypes.c_size_t), ctypes.c_void_p, ctypes.c_char_p]
 _libhip.hipMemset.restype = ctypes.c_int
 _libhip.hipMemset.argtypes = [ctypes.c_void_p, ctypes.c_int, ctypes.c_size_t]
+_libhip.hipMemcpyToSymbol.restype = ctypes.c_int
+_libhip.hipMemcpyToSymbol.argtypes = [ctypes.c_void_p, ctypes.c_void_p, ctypes.c_size_t, ctypes.c_size_t, ctypes.c_int]
 
 
 hipSuccess = 0
@@ -313,16 +315,23 @@ 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():
-            symbol = ctypes.c_void_p
-            size_kernel = ctypes.c_size_t
-            status = _libhip.hipModuleGetGlobal(symbol, size_kernel, self.current_module, str.encode(k))
+            #Format arguments, call hipModuleGetGlobal, and check return status
+            symbol_string = ctypes.c_char_p(k.encode('utf-8'))
+            symbol = ctypes.c_void_p()
+            symbol_ptr = ctypes.POINTER(ctypes.c_void_p)(symbol)
+            size_kernel = ctypes.c_size_t(0)
+
+            size_kernel_ptr = ctypes.POINTER(ctypes.c_size_t)(size_kernel)
+            status = _libhip.hipModuleGetGlobal(symbol_ptr, size_kernel_ptr, self.current_module, symbol_string)
             hip.hipCheckStatus(status)
+
+            #Format arguments and call hipMemcpy_htod
             dtype_str = str(v.dtype)
-            hip.hipMemcpy_htod(symbol, ctypes.byref(v.ctypes), ctypes.sizeof(dtype_map[dtype_str]) * v.size)
+            v_c = v.ctypes.data_as(ctypes.POINTER(dtype_map[dtype_str]))
+
+            hip.hipMemcpy_htod(symbol_ptr.contents, v_c, v.nbytes)
 
     def copy_shared_memory_args(self, smem_args):
         """add shared memory arguments to the kernel"""

test/test_hip_functions.py

@@ -56,7 +56,7 @@ def test_compile():
     """
 
     kernel_name = "vector_add"
-    kernel_sources = KernelSource(kernel_name, kernel_string, "cuda")
+    kernel_sources = KernelSource(kernel_name, kernel_string, "HIP")
     kernel_instance = KernelInstance(kernel_name, kernel_sources, kernel_string, [], None, None, dict(), [])
     dev = kt_hip.HipFunctions(0)
     try:
@@ -71,11 +71,11 @@ def test_memset_and_memcpy_dtoh():
     x = np.array(a).astype(np.int8)
     x_d = hip.hipMalloc(x.nbytes)
 
-    Hipfunc = kt_hip.HipFunctions()
-    Hipfunc.memset(x_d, 4, x.nbytes)
+    dev = kt_hip.HipFunctions()
+    dev.memset(x_d, 4, x.nbytes)
 
     output = np.empty(4, dtype=np.int8)
-    Hipfunc.memcpy_dtoh(output, x_d)
+    dev.memcpy_dtoh(output, x_d)
 
     assert all(output == np.full(4, 4))
 
@@ -86,12 +86,45 @@ def test_memcpy_htod():
     x_d = hip.hipMalloc(x.nbytes)
     output = np.empty(4, dtype=np.float32)
 
-    Hipfunc = kt_hip.HipFunctions()
-    Hipfunc.memcpy_htod(x_d, x)
-    Hipfunc.memcpy_dtoh(output, x_d)
+    dev = kt_hip.HipFunctions()
+    dev.memcpy_htod(x_d, x)
+    dev.memcpy_dtoh(output, x_d)
 
     assert all(output == x)
 
+@skip_if_no_pyhip
+def test_copy_constant_memory_args():
+    kernel_string = """
+    __constant__ float my_constant_data[100];
+    __global__ void copy_data_kernel(float* output) {
+        int idx = threadIdx.x + blockIdx.x * blockDim.x;
+        if (idx < 100) {
+            output[idx] = my_constant_data[idx];
+        }
+    }
+    """
+
+    kernel_name = "copy_data_kernel"
+    kernel_sources = KernelSource(kernel_name, kernel_string, "HIP")
+    kernel_instance = KernelInstance(kernel_name, kernel_sources, kernel_string, [], None, None, dict(), [])
+    dev = kt_hip.HipFunctions(0)
+    kernel = dev.compile(kernel_instance)
+
+    my_constant_data = np.full(100, 23).astype(np.float32)
+    cmem_args = {'my_constant_data': my_constant_data}
+    dev.copy_constant_memory_args(cmem_args)
+
+    output = np.full(100, 0).astype(np.float32)
+    gpu_args = dev.ready_argument_list([output])
+    
+    threads = (100, 1, 1)
+    grid = (1, 1, 1)
+    dev.run_kernel(kernel, gpu_args, threads, grid)
+
+    dev.memcpy_dtoh(output, gpu_args.field0)
+
+    assert(my_constant_data == output).all()
+
 def dummy_func(a, b, block=0, grid=0, stream=None, shared=0, texrefs=None):
     pass