diff --git a/INSTALL.rst b/INSTALL.rst
index cfa6275fb..53fef4926 100644
--- a/INSTALL.rst
+++ b/INSTALL.rst
@@ -111,6 +111,31 @@ Or you could install Kernel Tuner and PyOpenCL together if you haven't done so a
 
 If this fails, please see the PyOpenCL installation guide (https://wiki.tiker.net/PyOpenCL/Installation)
 
+HIP and PyHIP
+-------------
+
+Before we can install PyHIP, you'll need to have the HIP runtime and compiler installed on your system. 
+The HIP compiler is included as part of the ROCm software stack. Here is AMD's installation guide:
+
+* `ROCm Documentation: HIP Installation Guide <https://docs.amd.com/bundle/HIP-Installation-Guide-v5.3/page/Introduction_to_HIP_Installation_Guide.html>`__
+
+After you've installed HIP, you will need to install PyHIP. Run the following command in your terminal to install:
+
+.. code-block:: bash
+
+    pip install pyhip-interface
+
+Alternatively, you can install PyHIP from the source code. First, clone the repository from GitHub:
+
+.. code-block:: bash
+
+    git clone https://github.com/jatinx/PyHIP
+
+Then, navigate to the repository directory and run the following command to install:
+
+.. code-block:: bash
+    
+    python setup.py install
 
 Installing the git version
 --------------------------
@@ -128,6 +153,7 @@ You can install Kernel Tuner with several optional dependencies, the full list i
 
 - `cuda`: install pycuda along with kernel_tuner
 - `opencl`: install pycuda along with kernel_tuner
+- `hip`: install pyhip along with kernel_tuner
 - `doc`: installs packages required to build the documentation
 - `tutorial`: install packages required to run the guides
 - `dev`: install everything you need to start development on Kernel Tuner
diff --git a/README.rst b/README.rst
index 08e641061..2f8e8bcc0 100644
--- a/README.rst
+++ b/README.rst
@@ -28,9 +28,14 @@ To tune OpenCL kernels:
 - First, make sure you have an OpenCL compiler for your intended OpenCL platform
 - Then type: ``pip install kernel_tuner[opencl]``
 
-Or both:
+To tune HIP kernels:
 
-- ``pip install kernel_tuner[cuda,opencl]``
+- First, make sure you have an HIP runtime and compiler installed
+- Then type: ``pip install kernel_tuner[hip]``
+
+Or all:
+
+- ``pip install kernel_tuner[cuda,opencl,hip]``
 
 More information about how to install Kernel Tuner and its
 dependencies can be found in the `installation guide 
diff --git a/doc/source/architecture.png b/doc/source/architecture.png
new file mode 100644
index 000000000..bedeefa5b
Binary files /dev/null and b/doc/source/architecture.png differ
diff --git a/doc/source/design.rst b/doc/source/design.rst
index 614ac608d..4ca515e26 100644
--- a/doc/source/design.rst
+++ b/doc/source/design.rst
@@ -12,7 +12,7 @@ The Kernel Tuner is designed to be extensible and support
 different search and execution strategies. The current architecture of 
 the Kernel Tuner can be seen as:
 
-.. image:: architecture_0.4.3.png
+.. image:: architecture.png
    :width: 500pt
 
 At the top we have the kernel code and the Python script that tunes it, 
@@ -48,7 +48,7 @@ building blocks for implementing runners.
 The observers are explained in :ref:`observers`.
 
 At the bottom, the backends are shown. 
-PyCUDA, CuPy, cuda-python and PyOpenCL are for tuning either CUDA or OpenCL kernels.
+PyCUDA, CuPy, cuda-python, PyOpenCL and PyHIP are for tuning either CUDA, OpenCL, or HIP kernels.
 The C 
 Functions implementation can actually call any compiler, typically NVCC 
 or GCC is used. There is limited support for tuning Fortran kernels. 
@@ -128,6 +128,12 @@ kernel_tuner.backends.c.CFunctions
     :special-members: __init__
     :members:
 
+kernel_tuner.backends.hip.HipFunctions
+~~~~~~~~~~~~~~~~~~~~~~~~~
+.. autoclass:: kernel_tuner.backends.hip.HipFunctions
+    :special-members: __init__
+    :members:
+
 
 Util Functions
 --------------
diff --git a/doc/source/index.rst b/doc/source/index.rst
index 4a66d893b..c8af5e745 100644
--- a/doc/source/index.rst
+++ b/doc/source/index.rst
@@ -27,9 +27,14 @@ To tune OpenCL kernels:
 - First, make sure you have an OpenCL compiler for your intended OpenCL platform
 - Then type: ``pip install kernel_tuner[opencl]``
 
-Or both:
+To tune HIP kernels:
 
-- ``pip install kernel_tuner[cuda,opencl]``
+- First, make sure you have an HIP runtime and compiler installed
+- Then type: ``pip install kernel_tuner[hip]``
+
+Or all:
+
+- ``pip install kernel_tuner[cuda,opencl,hip]``
 
 More information about how to install Kernel Tuner and its
 dependencies can be found under :ref:`install`. 
diff --git a/examples/hip/test_vector_add.py b/examples/hip/test_vector_add.py
new file mode 100644
index 000000000..6c342632e
--- /dev/null
+++ b/examples/hip/test_vector_add.py
@@ -0,0 +1,40 @@
+#!/usr/bin/env python
+"""Minimal example for a HIP Kernel unit test with the Kernel Tuner"""
+
+import numpy
+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():
+
+    kernel_string = """
+    __global__ void vector_add(float *c, float *a, float *b, int n) {
+        int i = blockIdx.x * block_size_x + threadIdx.x;
+        if (i<n) {
+            c[i] = a[i] + b[i];
+        }
+    }
+    """
+
+    size = 10000000
+    problem_size = (size, 1)
+
+    a = numpy.random.randn(size).astype(numpy.float32)
+    b = numpy.random.randn(size).astype(numpy.float32)
+    c = numpy.zeros_like(b)
+    n = numpy.int32(size)
+
+    args = [c, a, b, n]
+    params = {"block_size_x": 512}
+
+    answer = run_kernel("vector_add", kernel_string, problem_size, args, params, lang="HIP")
+
+    assert numpy.allclose(answer[0], a+b, atol=1e-8)
\ No newline at end of file
diff --git a/examples/hip/vector_add.py b/examples/hip/vector_add.py
new file mode 100644
index 000000000..7e2810711
--- /dev/null
+++ b/examples/hip/vector_add.py
@@ -0,0 +1,43 @@
+#!/usr/bin/env python
+"""This is the minimal example from the README"""
+
+import numpy
+from kernel_tuner import tune_kernel
+from kernel_tuner.file_utils import store_output_file, store_metadata_file
+import logging
+from collections import OrderedDict
+
+def tune():
+
+    kernel_string = """
+    __global__ void vector_add(float *c, float *a, float *b, int n) {
+        int i = blockIdx.x * block_size_x + threadIdx.x;
+        if (i<n) {
+            c[i] = a[i] + b[i];
+        }
+    }
+    """
+
+    size = 10000000
+
+    a = numpy.random.randn(size).astype(numpy.float32)
+    b = numpy.random.randn(size).astype(numpy.float32)
+    c = numpy.zeros_like(b)
+    n = numpy.int32(size)
+
+    args = [c, a, b, n]
+
+    tune_params = OrderedDict()
+    tune_params["block_size_x"] = [128+64*i for i in range(15)]
+
+    results, env = tune_kernel("vector_add", kernel_string, size, args, tune_params, lang="HIP", 
+                               cache="vector_add_cache.json", log=logging.DEBUG)
+
+    # Store the metadata of this run
+    store_metadata_file("vector_add-metadata.json")
+
+    return results
+
+
+if __name__ == "__main__":
+    tune()
\ No newline at end of file
diff --git a/examples/hip/vector_add_simulation_mode.py b/examples/hip/vector_add_simulation_mode.py
new file mode 100644
index 000000000..a3368bca8
--- /dev/null
+++ b/examples/hip/vector_add_simulation_mode.py
@@ -0,0 +1,45 @@
+#!/usr/bin/env python
+"""This is the minimal example from the README"""
+
+import numpy
+from kernel_tuner import tune_kernel
+from kernel_tuner.file_utils import store_output_file, store_metadata_file
+import logging
+from collections import OrderedDict
+import os
+
+def tune():
+
+    kernel_string = """
+    __global__ void vector_add(float *c, float *a, float *b, int n) {
+        int i = blockIdx.x * block_size_x + threadIdx.x;
+        if (i<n) {
+            c[i] = a[i] + b[i];
+        }
+    }
+    """
+
+    size = 10000000
+
+    a = numpy.random.randn(size).astype(numpy.float32)
+    b = numpy.random.randn(size).astype(numpy.float32)
+    c = numpy.zeros_like(b)
+    n = numpy.int32(size)
+
+    args = [c, a, b, n]
+
+    tune_params = OrderedDict()
+    tune_params["block_size_x"] = [128+64*i for i in range(15)]
+
+    filename = "vector_add_cache.json"
+    if os.path.isfile(filename):
+        results, env = tune_kernel("vector_add", kernel_string, size, args, tune_params, 
+                                strategy="random_sample",  strategy_options=dict(max_fevals=10),
+                                lang="HIP", simulation_mode=True, cache="vector_add_cache.json")
+
+    else:
+        print(f"{filename} does not exist in the directory, run vector_add.py first.")
+
+
+if __name__ == "__main__":
+    tune()
\ No newline at end of file
diff --git a/kernel_tuner/backends/hip.py b/kernel_tuner/backends/hip.py
new file mode 100644
index 000000000..682f5e6a9
--- /dev/null
+++ b/kernel_tuner/backends/hip.py
@@ -0,0 +1,311 @@
+"""This module contains all HIP specific kernel_tuner functions"""
+
+import numpy as np
+import ctypes
+import ctypes.util
+import sys
+import logging
+
+from kernel_tuner.backends.backend import GPUBackend
+from kernel_tuner.observers.hip import HipRuntimeObserver
+
+# embedded in try block to be able to generate documentation
+# and run tests without pyhip installed
+try:
+    from pyhip import hip, hiprtc
+except ImportError:
+    print("Not able to import pyhip, check if PYTHONPATH includes PyHIP")
+    hip = None
+    hiprtc = None
+
+dtype_map = {
+    "bool": ctypes.c_bool,
+    "int8": ctypes.c_int8,
+    "int16": ctypes.c_int16,
+    "int32": ctypes.c_int32,
+    "int64": ctypes.c_int64,
+    "uint8": ctypes.c_uint8,
+    "uint16": ctypes.c_uint16,
+    "uint32": ctypes.c_uint32,
+    "uint64": ctypes.c_uint64,
+    "float32": ctypes.c_float,
+    "float64": ctypes.c_double,
+}
+
+hipSuccess = 0
+
+class HipFunctions(GPUBackend):
+    """Class that groups the HIP functions on maintains state about the device"""
+
+    def __init__(self, device=0, iterations=7, compiler_options=None, observers=None):
+        """instantiate HipFunctions object used for interacting with the HIP device
+
+        Instantiating this object will inspect and store certain device properties at
+        runtime, which are used during compilation and/or execution of kernels by the
+        kernel tuner. It also maintains a reference to the most recently compiled
+        source module for copying data to constant memory before kernel launch.
+
+        :param device: Number of HIP device to use for this context
+        :type device: int
+
+        :param iterations: Number of iterations used while benchmarking a kernel, 7 by default.
+        :type iterations: int
+        """
+        logging.debug("HipFunction instantiated")
+        
+        self.hipProps = hip.hipGetDeviceProperties(device)
+
+        self.name = self.hipProps._name.decode('utf-8')
+        self.max_threads = self.hipProps.maxThreadsPerBlock
+        self.device = device
+        self.compiler_options = compiler_options or []
+        self.iterations = iterations
+
+        env = dict()
+        env["device_name"] = self.name
+        env["iterations"] = self.iterations
+        env["compiler_options"] = compiler_options
+        self.env = env
+
+        # create a stream and events
+        self.stream = hip.hipStreamCreate()
+        self.start = hip.hipEventCreate()
+        self.end = hip.hipEventCreate()
+
+        # default dynamically allocated shared memory size, can be overwritten using smem_args
+        self.smem_size = 0
+
+        self.current_module = None
+
+        # setup observers
+        self.observers = observers or []
+        self.observers.append(HipRuntimeObserver(self))
+        for obs in self.observers:
+            obs.register_device(self)
+
+
+    def ready_argument_list(self, arguments):
+        """ready argument list to be passed to the HIP function
+
+        :param arguments: List of arguments to be passed to the HIP function.
+            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: Ctypes structure of arguments to be passed to the HIP function.
+        :rtype: ctypes structure
+        """
+        logging.debug("HipFunction ready_argument_list called")
+
+        ctype_args = []
+        data_ctypes = None
+        for arg in arguments:
+            dtype_str = str(arg.dtype)
+            # Allocate space on device for array and convert to ctypes
+            if isinstance(arg, np.ndarray):
+                if dtype_str in dtype_map.keys():
+                    device_ptr = hip.hipMalloc(arg.nbytes)
+                    data_ctypes = arg.ctypes.data_as(ctypes.POINTER(dtype_map[dtype_str]))
+                    hip.hipMemcpy_htod(device_ptr, data_ctypes, arg.nbytes)
+                    ctype_args.append(device_ptr)
+                else:
+                    raise TypeError("unknown dtype for ndarray")  
+            # Convert valid non-array arguments to ctypes      
+            elif isinstance(arg, np.generic):
+                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)]
+            def __getitem__(self, key):
+                return getattr(self, self._fields_[key][0])
+        
+        return ArgListStructure(*ctype_args)
+            
+    
+    def compile(self, kernel_instance):
+        """call the HIP compiler to compile the kernel, return the function
+        
+        :param kernel_instance: An object representing the specific instance of the tunable kernel
+            in the parameter space.
+        :type kernel_instance: kernel_tuner.core.KernelInstance
+        
+        :returns: An ctypes function that can be called directly.
+        :rtype: ctypes._FuncPtr
+        """
+        logging.debug("HipFunction compile called")
+
+        #Format and create program
+        kernel_string = kernel_instance.kernel_string
+        kernel_name = kernel_instance.name
+        if 'extern "C"' not in kernel_string:
+            kernel_string = 'extern "C" {\n' + kernel_string + "\n}"
+        kernel_ptr = hiprtc.hiprtcCreateProgram(kernel_string, kernel_name, [], [])
+        
+        try:
+            #Compile based on device (Not yet tested for non-AMD devices)
+            plat = hip.hipGetPlatformName()
+            if plat == "amd":
+                options_list = [f'--offload-arch={self.hipProps.gcnArchName}']
+                options_list.extend(self.compiler_options)
+                hiprtc.hiprtcCompileProgram(kernel_ptr, options_list)
+            else:
+                options_list = []
+                options_list.extend(self.compiler_options)
+                hiprtc.hiprtcCompileProgram(kernel_ptr, options_list)
+            
+            #Get module and kernel from compiled kernel string
+            code = hiprtc.hiprtcGetCode(kernel_ptr)
+            module = hip.hipModuleLoadData(code)
+            self.current_module = module
+            kernel = hip.hipModuleGetFunction(module, kernel_name)
+
+        except Exception as e:
+            log = hiprtc.hiprtcGetProgramLog(kernel_ptr)
+            print(log)
+            raise e
+        
+        return kernel
+    
+    def start_event(self):
+        """Records the event that marks the start of a measurement"""
+        logging.debug("HipFunction start_event called")
+
+        hip.hipEventRecord(self.start, self.stream)
+
+    def stop_event(self):
+        """Records the event that marks the end of a measurement"""
+        logging.debug("HipFunction stop_event called")
+
+        hip.hipEventRecord(self.end, self.stream)
+
+    def kernel_finished(self):
+        """Returns True if the kernel has finished, False otherwise"""
+        logging.debug("HipFunction kernel_finished called")
+        
+        # Query the status of the event
+        return hip.hipEventQuery(self.end)
+
+    def synchronize(self):
+        """Halts execution until device has finished its tasks"""
+        logging.debug("HipFunction synchronize called")
+
+        hip.hipDeviceSynchronize()
+
+    def run_kernel(self, func, gpu_args, threads, grid, stream=None):
+        """runs the HIP kernel passed as 'func'
+
+        :param func: A HIP kernel compiled for this specific kernel configuration
+        :type func: ctypes pionter
+
+        :param gpu_args: A ctypes structure 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
+
+        :param threads: A tuple listing the number of threads in each dimension of
+            the thread block
+        :type threads: tuple(int, int, int)
+
+        :param grid: A tuple listing the number of thread blocks in each dimension
+            of the grid
+        :type grid: tuple(int, int, int)
+        """
+        logging.debug("HipFunction run_kernel called")
+
+        if stream is None:
+            stream = self.stream
+
+        hip.hipModuleLaunchKernel(func, 
+                                  grid[0], grid[1], grid[2], 
+                                  threads[0], threads[1], threads[2],
+                                  self.smem_size,
+                                  stream,
+                                  gpu_args)
+
+    def memset(self, allocation, value, size):
+        """set the memory in allocation to the value in value
+
+        :param allocation: A GPU memory allocation unit
+        :type allocation: ctypes ptr
+
+        :param value: The value to set the memory to
+        :type value: a single 8-bit unsigned int
+
+        :param size: The size of to the allocation unit in bytes
+        :type size: int
+
+        """
+        logging.debug("HipFunction memset called")
+        
+        hip.hipMemset(allocation, value, size)
+
+    def memcpy_dtoh(self, dest, src):
+        """perform a device to host memory copy
+
+        :param dest: A numpy array in host memory to store the data
+        :type dest: numpy.ndarray
+
+        :param src: A GPU memory allocation unit
+        :type src: ctypes ptr
+        """
+        logging.debug("HipFunction memcpy_dtoh called")
+
+        # Format arguments to correct type and perform memory copy
+        dtype_str = str(dest.dtype)
+        dest_c = dest.ctypes.data_as(ctypes.POINTER(dtype_map[dtype_str]))
+        hip.hipMemcpy_dtoh(dest_c, src, dest.nbytes)
+
+    def memcpy_htod(self, dest, src):
+        """perform a host to device memory copy
+
+        :param dest: A GPU memory allocation unit
+        :type dest: ctypes ptr
+
+        :param src: A numpy array in host memory to store the data
+        :type src: numpy.ndarray
+        """
+        logging.debug("HipFunction memcpy_htod called")
+
+        # Format arguments to correct type and perform memory copy
+        dtype_str = str(src.dtype)
+        src_c = src.ctypes.data_as(ctypes.POINTER(dtype_map[dtype_str]))
+        hip.hipMemcpy_htod(dest, src_c, src.nbytes)
+
+    def copy_constant_memory_args(self, cmem_args):
+        """adds constant memory arguments to the most recently compiled module
+
+        :param cmem_args: A dictionary containing the data to be passed to the
+            device constant memory. The format to be used is as follows: A
+            string key is used to name the constant memory symbol to which the
+            value needs to be copied. Similar to regular arguments, these need
+            to be numpy objects, such as numpy.ndarray or numpy.int32, and so on.
+        :type cmem_args: dict( string: numpy.ndarray, ... )
+        """
+        logging.debug("HipFunction copy_constant_memory_args called")
+
+        # Iterate over dictionary
+        for k, v in cmem_args.items():
+            #Get symbol pointer
+            symbol_ptr, _ = hip.hipModuleGetGlobal(self.current_module, k)
+
+            #Format arguments and perform memory copy
+            dtype_str = str(v.dtype)
+            v_c = v.ctypes.data_as(ctypes.POINTER(dtype_map[dtype_str]))
+            hip.hipMemcpy_htod(symbol_ptr, v_c, v.nbytes)
+
+    def copy_shared_memory_args(self, smem_args):
+        """add shared memory arguments to the kernel"""
+        logging.debug("HipFunction copy_shared_memory_args called")
+
+        self.smem_size = smem_args["size"]
+
+    def copy_texture_memory_args(self, texmem_args):
+        logging.debug("HipFunction copy_texture_memory_args called")
+
+        raise NotImplementedError("HIP backend does not support texture memory")
+
+    units = {"time": "ms", "power": "s,mW", "energy": "J"}
diff --git a/kernel_tuner/core.py b/kernel_tuner/core.py
index accb58cc1..1eeeb0d8d 100644
--- a/kernel_tuner/core.py
+++ b/kernel_tuner/core.py
@@ -19,6 +19,7 @@
 from kernel_tuner.backends.opencl import OpenCLFunctions
 from kernel_tuner.backends.c import CFunctions
 from kernel_tuner.backends.opencl import OpenCLFunctions
+from kernel_tuner.backends.hip import HipFunctions
 import kernel_tuner.util as util
 
 try:
@@ -212,7 +213,7 @@ def __init__(self, kernel_source, device=0, platform=0, quiet=False, compiler=No
         :type device: int
 
         :param lang: Specifies the language used for GPU kernels.
-            Currently supported: "CUDA", "OpenCL", or "C"
+            Currently supported: "CUDA", "OpenCL", "HIP" or "C"
         :type lang: string
 
         :param compiler_options: The compiler options to use when compiling kernels for this device.
@@ -239,6 +240,8 @@ def __init__(self, kernel_source, device=0, platform=0, quiet=False, compiler=No
             dev = OpenCLFunctions(device, platform, compiler_options=compiler_options, iterations=iterations, observers=observers)
         elif lang.upper() in ["C", "FORTRAN"]:
             dev = CFunctions(compiler=compiler, compiler_options=compiler_options, iterations=iterations)
+        elif lang.upper() == "HIP":
+            dev = HipFunctions(device, compiler_options=compiler_options, iterations=iterations, observers=observers)
         else:
             raise ValueError("Sorry, support for languages other than CUDA, OpenCL, or C is not implemented yet")
 
@@ -422,48 +425,48 @@ def compile_and_benchmark(self, kernel_source, gpu_args, params, kernel_options,
 
         instance = self.create_kernel_instance(kernel_source, kernel_options, params, verbose)
         if isinstance(instance, util.ErrorConfig):
-            return instance
-
-        try:
-            # compile the kernel
-            start_compilation = time.perf_counter()
-            func = self.compile_kernel(instance, verbose)
-            if not func:
-                result[to.objective] = util.CompilationFailedConfig()
-            else:
-                # add shared memory arguments to compiled module
-                if kernel_options.smem_args is not None:
-                    self.dev.copy_shared_memory_args(util.get_smem_args(kernel_options.smem_args, params))
-                # add constant memory arguments to compiled module
-                if kernel_options.cmem_args is not None:
-                    self.dev.copy_constant_memory_args(kernel_options.cmem_args)
-                # add texture memory arguments to compiled module
-                if kernel_options.texmem_args is not None:
-                    self.dev.copy_texture_memory_args(kernel_options.texmem_args)
-
-            # stop compilation stopwatch and convert to miliseconds
-            last_compilation_time = 1000 * (time.perf_counter() - start_compilation)
-
-            # test kernel for correctness
-            if func and (to.answer or to.verify):
-                start_verification = time.perf_counter()
-                self.check_kernel_output(func, gpu_args, instance, to.answer, to.atol, to.verify, verbose)
-                last_verification_time = 1000 * (time.perf_counter() - start_verification)
-
-            # benchmark
-            if func:
-                start_benchmark = time.perf_counter()
-                result.update(self.benchmark(func, gpu_args, instance, verbose, to.objective))
-                last_benchmark_time = 1000 * (time.perf_counter() - start_benchmark)
-
-        except Exception as e:
-            # dump kernel sources to temp file
-            temp_filenames = instance.prepare_temp_files_for_error_msg()
-            print("Error while compiling or benchmarking, see source files: " + " ".join(temp_filenames))
-            raise e
+            result[to.objective] = util.InvalidConfig()
+        else:
+            try:
+                # compile the kernel
+                start_compilation = time.perf_counter()
+                func = self.compile_kernel(instance, verbose)
+                if not func:
+                    result[to.objective] = util.CompilationFailedConfig()
+                else:
+                    # add shared memory arguments to compiled module
+                    if kernel_options.smem_args is not None:
+                        self.dev.copy_shared_memory_args(util.get_smem_args(kernel_options.smem_args, params))
+                    # add constant memory arguments to compiled module
+                    if kernel_options.cmem_args is not None:
+                        self.dev.copy_constant_memory_args(kernel_options.cmem_args)
+                    # add texture memory arguments to compiled module
+                    if kernel_options.texmem_args is not None:
+                        self.dev.copy_texture_memory_args(kernel_options.texmem_args)
+
+                # stop compilation stopwatch and convert to miliseconds
+                last_compilation_time = 1000 * (time.perf_counter() - start_compilation)
+
+                # test kernel for correctness
+                if func and (to.answer or to.verify):
+                    start_verification = time.perf_counter()
+                    self.check_kernel_output(func, gpu_args, instance, to.answer, to.atol, to.verify, verbose)
+                    last_verification_time = 1000 * (time.perf_counter() - start_verification)
+
+                # benchmark
+                if func:
+                    start_benchmark = time.perf_counter()
+                    result.update(self.benchmark(func, gpu_args, instance, verbose, to.objective))
+                    last_benchmark_time = 1000 * (time.perf_counter() - start_benchmark)
+
+            except Exception as e:
+                # dump kernel sources to temp file
+                temp_filenames = instance.prepare_temp_files_for_error_msg()
+                print("Error while compiling or benchmarking, see source files: " + " ".join(temp_filenames))
+                raise e
 
-        #clean up any temporary files, if no error occured
-        instance.delete_temp_files()
+            #clean up any temporary files, if no error occured
+            instance.delete_temp_files()
 
         result['compile_time'] = last_compilation_time or 0
         result['verification_time'] = last_verification_time or 0
diff --git a/kernel_tuner/file_utils.py b/kernel_tuner/file_utils.py
index 321d0db80..0d5024187 100644
--- a/kernel_tuner/file_utils.py
+++ b/kernel_tuner/file_utils.py
@@ -105,6 +105,8 @@ def store_output_file(output_filename: str, results, tune_params, objective="tim
         timings["validation"] = result["verification_time"]
         if "times" in result:
             timings["runtimes"] = result["times"]
+        else:
+            timings["runtimes"] = []
         out["times"] = timings
 
         # encode the validity of the configuration
@@ -225,7 +227,7 @@ def store_metadata_file(metadata_filename: str):
             hardware_description_string = "[" + hardware_description_string + "]"
         metadata["operating_system"] = os_string
     except:
-        hardware_description_string = "[error retrieving hardware description]"
+        hardware_description_string = '["error retrieving hardware description"]'
         metadata["operating_system"] = "unidentified OS"
     metadata["hardware"] = dict(hardware_description=json.loads(hardware_description_string))
 
@@ -233,14 +235,14 @@ def store_metadata_file(metadata_filename: str):
     device_query = {}
     try:
         device_query["nvidia-smi"] = get_device_query("nvidia")
-    except FileNotFoundError:
-        # ignore if nvidia-smi is not found
+    except Exception:
+        # ignore if nvidia-smi is not found, or parse error occurs
         pass
 
     try:
         device_query["rocm-smi"] = get_device_query("amd")
-    except FileNotFoundError:
-        # ignore if rocm-smi is not found
+    except Exception:
+        # ignore if rocm-smi is not found, or parse error occurs
         pass
 
     metadata["environment"] = dict(device_query=device_query, requirements=get_dependencies())
diff --git a/kernel_tuner/interface.py b/kernel_tuner/interface.py
index 9abfa1de9..15fe04ec1 100644
--- a/kernel_tuner/interface.py
+++ b/kernel_tuner/interface.py
@@ -98,7 +98,7 @@ def __deepcopy__(self, _):
     (
         "kernel_source",
         (
-            """The CUDA, OpenCL, or C kernel code.
+            """The CUDA, OpenCL, HIP, or C kernel code.
             It is allowed for the code to be passed as a string, a filename, a function
             that returns a string of code, or a list when the code needs auxilliary files.
 
@@ -124,7 +124,7 @@ def __deepcopy__(self, _):
             """Specifies the language used for GPU kernels. The kernel_tuner
         automatically detects the language, but if it fails, you may specify
         the language using this argument, currently supported: "CUDA", "Cupy",
-        "OpenCL", or "C".""",
+        "OpenCL", "HIP", or "C".""",
             "string",
         ),
     ),
diff --git a/kernel_tuner/observers/hip.py b/kernel_tuner/observers/hip.py
new file mode 100644
index 000000000..72a3cb4fe
--- /dev/null
+++ b/kernel_tuner/observers/hip.py
@@ -0,0 +1,31 @@
+import numpy as np
+
+from kernel_tuner.observers.observer import BenchmarkObserver
+
+try:
+    from pyhip import hip, hiprtc
+except ImportError:
+    print("Not able to import pyhip, check if PYTHONPATH includes PyHIP")
+    hip = None
+    hiprtc = None
+
+
+class HipRuntimeObserver(BenchmarkObserver):
+    """Observer that measures time using CUDA events during benchmarking"""
+
+    def __init__(self, dev):
+        self.dev = dev
+        self.stream = dev.stream
+        self.start = dev.start
+        self.end = dev.end
+        self.times = []
+
+    def after_finish(self):
+        # Time is measured in milliseconds
+        EventElapsedTime = hip.hipEventElapsedTime(self.start, self.end)
+        self.times.append(EventElapsedTime)
+
+    def get_results(self):
+        results = {"time": np.average(self.times), "times": self.times.copy()}
+        self.times = []
+        return results
diff --git a/kernel_tuner/util.py b/kernel_tuner/util.py
index fd7ee3c37..0e3dee4d6 100644
--- a/kernel_tuner/util.py
+++ b/kernel_tuner/util.py
@@ -618,8 +618,8 @@ def prepare_kernel_string(kernel_name, kernel_string, params, grid, threads, blo
         v = str(v)
         v = v.replace("\n", "\\\n")
 
-        if "loop_unroll_factor" in k and lang == "CUDA":
-            # this handles the special case that in CUDA
+        if "loop_unroll_factor" in k and lang in ("CUDA", "HIP"):
+            # this handles the special case that in CUDA/HIP
             # pragma unroll loop_unroll_factor, loop_unroll_factor should be a constant integer expression
             # in OpenCL this isn't the case and we can just insert "#define loop_unroll_factor N"
             # using 0 to disable specifying a loop unrolling factor for this loop
diff --git a/setup.py b/setup.py
index 040a1c931..cc4a64928 100644
--- a/setup.py
+++ b/setup.py
@@ -77,6 +77,7 @@ def readme():
         "cuda": ["pycuda", "nvidia-ml-py", "pynvml>=11.4.1"],
         "opencl": ["pyopencl"],
         "cuda_opencl": ["pycuda", "pyopencl"],
+        "hip": ["pyhip-interface"],
         "tutorial": ["jupyter", "matplotlib", "pandas"],
         "dev": [
             "numpy>=1.13.3",
diff --git a/test/context.py b/test/context.py
index 08d2a190a..763b3b7c9 100644
--- a/test/context.py
+++ b/test/context.py
@@ -1,6 +1,7 @@
 import sys
 import subprocess
 import shutil
+import os
 
 import pytest
 
@@ -44,6 +45,15 @@
 except Exception:
     cuda_present = False
 
+PYHIP_PATH = os.environ.get('PYHIP_PATH')  # get the PYHIP_PATH environment variable
+try: 
+    if PYHIP_PATH is not None:
+        sys.path.insert(0, PYHIP_PATH)
+    from pyhip import hip, hiprtc
+    pyhip_present = True
+except ImportError:
+    pyhip_present = False
+
 skip_if_no_pycuda = pytest.mark.skipif(not pycuda_present, reason="PyCuda not installed or no CUDA device detected")
 skip_if_no_pynvml = pytest.mark.skipif(not pynvml_present, reason="NVML not installed")
 skip_if_no_cupy = pytest.mark.skipif(not cupy_present, reason="CuPy not installed or no CUDA device detected")
@@ -52,6 +62,7 @@
 skip_if_no_gcc = pytest.mark.skipif(not gcc_present, reason="No gcc on PATH")
 skip_if_no_gfortran = pytest.mark.skipif(not gfortran_present, reason="No gfortran on PATH")
 skip_if_no_openmp = pytest.mark.skipif(not openmp_present, reason="No OpenMP found")
+skip_if_no_pyhip = pytest.mark.skipif(not pyhip_present, reason="No PyHIP found")
 
 
 def skip_backend(backend: str):
diff --git a/test/test_hip_functions.py b/test/test_hip_functions.py
new file mode 100644
index 000000000..ce3eb0642
--- /dev/null
+++ b/test/test_hip_functions.py
@@ -0,0 +1,167 @@
+import numpy as np
+import ctypes
+from .context import skip_if_no_pyhip
+from collections import OrderedDict
+
+import pytest
+import kernel_tuner
+from kernel_tuner import tune_kernel
+from kernel_tuner.backends import hip as kt_hip
+from kernel_tuner.core import KernelSource, KernelInstance
+
+try: 
+    from pyhip import hip, hiprtc
+    hip_present = True
+except ImportError:
+    pass
+
+@pytest.fixture
+def env():
+    kernel_string = """
+    extern "C" __global__ void vector_add(float *c, float *a, float *b, int n) {
+        int i = blockIdx.x * block_size_x + threadIdx.x;
+        if (i<n) {
+            c[i] = a[i] + b[i];
+        }
+    }
+    """
+
+    size = 100
+    a = np.random.randn(size).astype(np.float32)
+    b = np.random.randn(size).astype(np.float32)
+    c = np.zeros_like(b)
+    n = np.int32(size)
+
+    args = [c, a, b, n]
+    tune_params = OrderedDict()
+    tune_params["block_size_x"] = [128 + 64 * i for i in range(15)]
+
+    return ["vector_add", kernel_string, size, args, tune_params]
+
+@skip_if_no_pyhip
+def test_ready_argument_list():
+
+    size = 1000
+    a = np.int32(75)
+    b = np.random.randn(size).astype(np.float32)
+    c = np.bool_(True)
+    d = np.zeros_like(b)
+
+    arguments = [d, a, b, c]
+
+    class ArgListStructure(ctypes.Structure):
+        _fields_ = [("field0", ctypes.POINTER(ctypes.c_float)),
+                    ("field1", ctypes.c_int),
+                    ("field2", ctypes.POINTER(ctypes.c_float)),
+                    ("field3", ctypes.c_bool)]
+        def __getitem__(self, key):
+                return getattr(self, self._fields_[key][0])
+
+    dev = kt_hip.HipFunctions(0)
+    gpu_args = dev.ready_argument_list(arguments)
+
+    argListStructure = ArgListStructure(d.ctypes.data_as(ctypes.POINTER(ctypes.c_float)),
+                                        ctypes.c_int(a),
+                                        b.ctypes.data_as(ctypes.POINTER(ctypes.c_float)),
+                                        ctypes.c_bool(c))
+    
+    assert(gpu_args[1] == argListStructure[1])
+    assert(gpu_args[3] == argListStructure[3])
+
+@skip_if_no_pyhip
+def test_compile():
+
+    kernel_string = """
+    __global__ void vector_add(float *c, float *a, float *b, int n) {
+        int i = blockIdx.x * blockDim.x + threadIdx.x;
+        if (i<n) {
+            c[i] = a[i] + b[i];
+        }
+    }
+    """
+
+    kernel_name = "vector_add"
+    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:
+        dev.compile(kernel_instance)
+    except Exception as e:
+        pytest.fail("Did not expect any exception:" + str(e))
+
+
+@skip_if_no_pyhip
+def test_memset_and_memcpy_dtoh():
+    a = [1, 2, 3, 4]
+    x = np.array(a).astype(np.int8)
+    x_d = hip.hipMalloc(x.nbytes)
+
+    dev = kt_hip.HipFunctions()
+    dev.memset(x_d, 4, x.nbytes)
+
+    output = np.empty(4, dtype=np.int8)
+    dev.memcpy_dtoh(output, x_d)
+
+    assert all(output == np.full(4, 4))
+
+@skip_if_no_pyhip
+def test_memcpy_htod():
+    a = [1, 2, 3, 4]
+    x = np.array(a).astype(np.float32)
+    x_d = hip.hipMalloc(x.nbytes)
+    output = np.empty(4, dtype=np.float32)
+
+    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()
+
+@skip_if_no_pyhip
+def test_smem_args(env):
+    result, _ = tune_kernel(*env,
+                            smem_args=dict(size="block_size_x*4"),
+                            verbose=True, lang="HIP")
+    tune_params = env[-1]
+    assert len(result) == len(tune_params["block_size_x"])
+    result, _ = tune_kernel(
+        *env,
+        smem_args=dict(size=lambda p: p['block_size_x'] * 4),
+        verbose=True, lang="HIP")
+    tune_params = env[-1]
+    assert len(result) == len(tune_params["block_size_x"])
+
+