diff --git a/extern/hipew/include/hipew.h b/extern/hipew/include/hipew.h
index f82654ffe93..df7498aaeae 100644
--- a/extern/hipew/include/hipew.h
+++ b/extern/hipew/include/hipew.h
@@ -185,13 +185,26 @@ typedef struct textureReference {
 
 typedef textureReference* hipTexRef;
 
+/**
+ * ROCm 6 and ROCm 5 memory types are different.
+ * For now, we include both in the enum and then use the get_hip_memory_type
+ * Function to convert. When removing ROCm 5 compatibility this can be simplified.
+*/
 typedef enum hipMemoryType {
-  hipMemoryTypeHost = 0x00,
-  hipMemoryTypeDevice = 0x01,
-  hipMemoryTypeArray = 0x02,
-  hipMemoryTypeUnified = 0x03,
+  hipMemoryTypeHost_v5 = 0x00,
+  hipMemoryTypeDevice_v5 = 0x01,
+  hipMemoryTypeArray_v5 = 0x02,
+  hipMemoryTypeUnified_v5 = 0x03,  
+  hipMemoryTypeUnregistered = 0,
+  hipMemoryTypeHost = 1,
+  hipMemoryTypeDevice = 2,
+  hipMemoryTypeManaged = 3,
+  hipMemoryTypeArray = 10,
+  hipMemoryTypeUnified = 11,
 } hipMemoryType;
 
+hipMemoryType get_hip_memory_type(hipMemoryType mem_type, int runtime_version);
+
 /**
  * Pointer attributes
  */
@@ -316,7 +329,7 @@ typedef enum hipDeviceAttribute_t {
   hipDeviceAttributeConcurrentManagedAccess,          ///< Device can coherently access managed memory concurrently with the CPU
   hipDeviceAttributeCooperativeLaunch,                ///< Support cooperative launch
   hipDeviceAttributeCooperativeMultiDeviceLaunch,     ///< Support cooperative launch on multiple devices
-  hipDeviceAttributeDeviceOverlap,                    ///< Cuda only. Device can concurrently copy memory and execute a kernel.
+  hipDeviceAttributeDeviceOverlap,               ///< Cuda only. Device can concurrently copy memory and execute a kernel.  
                                                       ///< Deprecated. Use instead asyncEngineCount.
   hipDeviceAttributeDirectManagedMemAccessFromHost,   ///< Host can directly access managed memory on
                                                       ///< the device without migration
@@ -420,6 +433,7 @@ typedef enum hipDeviceAttribute_t {
                                                               ///< hipStreamWaitValue64() , '0' otherwise.
   hipDeviceAttributeAmdSpecificEnd = 19999,
   hipDeviceAttributeVendorSpecificBegin = 20000,
+  hipDeviceAttribute
   // Extended attributes for vendors
 } hipDeviceAttribute_t;
 
@@ -1160,6 +1174,7 @@ typedef const char* HIPAPI thipGetErrorString(hipError_t error);
 typedef hipError_t HIPAPI thipGetLastError(hipError_t error);
 typedef hipError_t HIPAPI thipInit(unsigned int Flags);
 typedef hipError_t HIPAPI thipDriverGetVersion(int* driverVersion);
+typedef hipError_t HIPAPI thipRuntimeGetVersion(int* runtimeVersion);
 typedef hipError_t HIPAPI thipGetDevice(int* device);
 typedef hipError_t HIPAPI thipGetDeviceCount(int* count);
 typedef hipError_t HIPAPI thipGetDeviceProperties(hipDeviceProp_t* props, int deviceId);
@@ -1310,6 +1325,7 @@ extern thipGetErrorString* hipGetErrorString;
 extern thipGetLastError* hipGetLastError;
 extern thipInit *hipInit;
 extern thipDriverGetVersion *hipDriverGetVersion;
+extern thipRuntimeGetVersion *hipRuntimeGetVersion;
 extern thipGetDevice *hipGetDevice;
 extern thipGetDeviceCount *hipGetDeviceCount;
 extern thipGetDeviceProperties *hipGetDeviceProperties;
diff --git a/extern/hipew/src/hipew.c b/extern/hipew/src/hipew.c
index 4927f86b804..74b2ef111d6 100644
--- a/extern/hipew/src/hipew.c
+++ b/extern/hipew/src/hipew.c
@@ -35,6 +35,7 @@ thipGetErrorString *hipGetErrorString;
 thipGetLastError *hipGetLastError;
 thipInit *hipInit;
 thipDriverGetVersion *hipDriverGetVersion;
+thipRuntimeGetVersion *hipRuntimeGetVersion;
 thipGetDevice *hipGetDevice;
 thipGetDeviceCount *hipGetDeviceCount;
 thipGetDeviceProperties *hipGetDeviceProperties;
@@ -285,6 +286,7 @@ static int hipewHipInit(void) {
   HIP_LIBRARY_FIND_CHECKED(hipGetLastError);
   HIP_LIBRARY_FIND_CHECKED(hipInit);
   HIP_LIBRARY_FIND_CHECKED(hipDriverGetVersion);
+  HIP_LIBRARY_FIND_CHECKED(hipRuntimeGetVersion);
   HIP_LIBRARY_FIND_CHECKED(hipGetDevice);
   HIP_LIBRARY_FIND_CHECKED(hipGetDeviceCount);
   HIP_LIBRARY_FIND_CHECKED(hipGetDeviceProperties);
@@ -410,7 +412,28 @@ static int hipewHipInit(void) {
   return result;
 }
 
+hipMemoryType get_hip_memory_type(hipMemoryType mem_type, int runtime_version) {
+  /** Convert hipMemoryType for backwards compatibility with rocm5/6. 
+   * This can be removed when support for ROCm 5 is removed. */
 
+  /* If version is 5 we need to use the old enum vals (60000000 is start of ROCm 6) */
+  if (runtime_version > 60000000) {
+    return mem_type;
+  }
+
+  switch (mem_type) {
+    case hipMemoryTypeHost:
+      return hipMemoryTypeHost_v5;
+    case hipMemoryTypeDevice:
+      return hipMemoryTypeDevice_v5;
+    case hipMemoryTypeArray:
+      return hipMemoryTypeArray_v5;
+    case hipMemoryTypeUnified:
+      return hipMemoryTypeUnified_v5;
+    default:
+      return hipMemoryTypeUnregistered;  /* This should not happen. */
+  }
+}
 
 int hipewInit(hipuint32_t flags) {
   int result = HIPEW_SUCCESS;
diff --git a/intern/cycles/device/hip/device_impl.cpp b/intern/cycles/device/hip/device_impl.cpp
index 2739cb255ad..bbbe357bf76 100644
--- a/intern/cycles/device/hip/device_impl.cpp
+++ b/intern/cycles/device/hip/device_impl.cpp
@@ -116,6 +116,9 @@ HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
   hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
   hipDevArchitecture = major * 100 + minor * 10;
 
+  /* Get hip runtime Version needed for memory types. */
+  hip_assert(hipRuntimeGetVersion(&hipRuntimeVersion));
+
   /* Pop context set by hipCtxCreate. */
   hipCtxPopCurrent(NULL);
 }
@@ -745,9 +748,9 @@ void HIPDevice::tex_alloc(device_texture &mem)
 
     HIP_MEMCPY3D param;
     memset(&param, 0, sizeof(HIP_MEMCPY3D));
-    param.dstMemoryType = hipMemoryTypeArray;
+    param.dstMemoryType = get_memory_type(hipMemoryTypeArray);
     param.dstArray = array_3d;
-    param.srcMemoryType = hipMemoryTypeHost;
+    param.srcMemoryType = get_memory_type(hipMemoryTypeHost);
     param.srcHost = mem.host_pointer;
     param.srcPitch = src_pitch;
     param.WidthInBytes = param.srcPitch;
@@ -777,10 +780,10 @@ void HIPDevice::tex_alloc(device_texture &mem)
 
     hip_Memcpy2D param;
     memset(&param, 0, sizeof(param));
-    param.dstMemoryType = hipMemoryTypeDevice;
+    param.dstMemoryType = get_memory_type(hipMemoryTypeDevice);
     param.dstDevice = mem.device_pointer;
     param.dstPitch = dst_pitch;
-    param.srcMemoryType = hipMemoryTypeHost;
+    param.srcMemoryType = get_memory_type(hipMemoryTypeHost);
     param.srcHost = mem.host_pointer;
     param.srcPitch = src_pitch;
     param.WidthInBytes = param.srcPitch;
@@ -958,6 +961,11 @@ int HIPDevice::get_device_default_attribute(hipDeviceAttribute_t attribute, int
   return value;
 }
 
+hipMemoryType HIPDevice::get_memory_type(hipMemoryType mem_type)
+{
+  return get_hip_memory_type(mem_type, hipRuntimeVersion);
+}
+
 CCL_NAMESPACE_END
 
 #endif
diff --git a/intern/cycles/device/hip/device_impl.h b/intern/cycles/device/hip/device_impl.h
index bded023bd4a..fdbec2e3dbf 100644
--- a/intern/cycles/device/hip/device_impl.h
+++ b/intern/cycles/device/hip/device_impl.h
@@ -32,6 +32,7 @@ class HIPDevice : public GPUDevice {
   int pitch_alignment;
   int hipDevId;
   int hipDevArchitecture;
+  int hipRuntimeVersion;
   bool first_error;
 
   HIPDeviceKernels kernels;
@@ -102,6 +103,7 @@ class HIPDevice : public GPUDevice {
  protected:
   bool get_device_attribute(hipDeviceAttribute_t attribute, int *value);
   int get_device_default_attribute(hipDeviceAttribute_t attribute, int default_value);
+  hipMemoryType get_memory_type(hipMemoryType mem_type);
 };
 
 CCL_NAMESPACE_END