diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h
index 7ec70fdc3d..d72fd38d8d 100644
--- a/hipamd/include/hip/amd_detail/hip_prof_str.h
+++ b/hipamd/include/hip/amd_detail/hip_prof_str.h
@@ -370,10 +370,11 @@ enum hip_api_id_t {
   HIP_API_ID_hipGraphMemFreeNodeGetParams = 357,
   HIP_API_ID_hipModuleLaunchCooperativeKernel = 358,
   HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice = 359,
-  HIP_API_ID_LAST = 359,
+  HIP_API_ID_hipArray3DGetDescriptor = 360,
+  HIP_API_ID_hipArrayGetDescriptor = 361,
+  HIP_API_ID_hipArrayGetInfo = 362,
+  HIP_API_ID_LAST = 362,
 
-  HIP_API_ID_hipArray3DGetDescriptor = HIP_API_ID_NONE,
-  HIP_API_ID_hipArrayGetDescriptor = HIP_API_ID_NONE,
   HIP_API_ID_hipBindTexture = HIP_API_ID_NONE,
   HIP_API_ID_hipBindTexture2D = HIP_API_ID_NONE,
   HIP_API_ID_hipBindTextureToArray = HIP_API_ID_NONE,
@@ -420,8 +421,11 @@ static inline const char* hip_api_name(const uint32_t id) {
     case HIP_API_ID___hipPopCallConfiguration: return "__hipPopCallConfiguration";
     case HIP_API_ID___hipPushCallConfiguration: return "__hipPushCallConfiguration";
     case HIP_API_ID_hipArray3DCreate: return "hipArray3DCreate";
+    case HIP_API_ID_hipArray3DGetDescriptor: return "hipArray3DGetDescriptor";
     case HIP_API_ID_hipArrayCreate: return "hipArrayCreate";
     case HIP_API_ID_hipArrayDestroy: return "hipArrayDestroy";
+    case HIP_API_ID_hipArrayGetDescriptor: return "hipArrayGetDescriptor";
+    case HIP_API_ID_hipArrayGetInfo: return "hipArrayGetInfo";
     case HIP_API_ID_hipChooseDevice: return "hipChooseDevice";
     case HIP_API_ID_hipConfigureCall: return "hipConfigureCall";
     case HIP_API_ID_hipCreateSurfaceObject: return "hipCreateSurfaceObject";
@@ -782,8 +786,11 @@ static inline uint32_t hipApiIdByName(const char* name) {
   if (strcmp("__hipPopCallConfiguration", name) == 0) return HIP_API_ID___hipPopCallConfiguration;
   if (strcmp("__hipPushCallConfiguration", name) == 0) return HIP_API_ID___hipPushCallConfiguration;
   if (strcmp("hipArray3DCreate", name) == 0) return HIP_API_ID_hipArray3DCreate;
+  if (strcmp("hipArray3DGetDescriptor", name) == 0) return HIP_API_ID_hipArray3DGetDescriptor;
   if (strcmp("hipArrayCreate", name) == 0) return HIP_API_ID_hipArrayCreate;
   if (strcmp("hipArrayDestroy", name) == 0) return HIP_API_ID_hipArrayDestroy;
+  if (strcmp("hipArrayGetDescriptor", name) == 0) return HIP_API_ID_hipArrayGetDescriptor;
+  if (strcmp("hipArrayGetInfo", name) == 0) return HIP_API_ID_hipArrayGetInfo;
   if (strcmp("hipChooseDevice", name) == 0) return HIP_API_ID_hipChooseDevice;
   if (strcmp("hipConfigureCall", name) == 0) return HIP_API_ID_hipConfigureCall;
   if (strcmp("hipCreateSurfaceObject", name) == 0) return HIP_API_ID_hipCreateSurfaceObject;
@@ -1164,6 +1171,12 @@ typedef struct hip_api_data_s {
       const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray;
       HIP_ARRAY3D_DESCRIPTOR pAllocateArray__val;
     } hipArray3DCreate;
+    struct {
+      HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor;
+      HIP_ARRAY3D_DESCRIPTOR pArrayDescriptor__val;
+      hipArray* array;
+      hipArray array__val;
+    } hipArray3DGetDescriptor;
     struct {
       hipArray** pHandle;
       hipArray* pHandle__val;
@@ -1174,6 +1187,22 @@ typedef struct hip_api_data_s {
       hipArray* array;
       hipArray array__val;
     } hipArrayDestroy;
+    struct {
+      HIP_ARRAY_DESCRIPTOR* pArrayDescriptor;
+      HIP_ARRAY_DESCRIPTOR pArrayDescriptor__val;
+      hipArray* array;
+      hipArray array__val;
+    } hipArrayGetDescriptor;
+    struct {
+      hipChannelFormatDesc* desc;
+      hipChannelFormatDesc desc__val;
+      hipExtent* extent;
+      hipExtent extent__val;
+      unsigned int* flags;
+      unsigned int flags__val;
+      hipArray* array;
+      hipArray array__val;
+    } hipArrayGetInfo;
     struct {
       int* device;
       int device__val;
@@ -3252,6 +3281,11 @@ typedef struct hip_api_data_s {
   cb_data.args.hipArray3DCreate.array = (hipArray**)array; \
   cb_data.args.hipArray3DCreate.pAllocateArray = (const HIP_ARRAY3D_DESCRIPTOR*)pAllocateArray; \
 };
+// hipArray3DGetDescriptor[('HIP_ARRAY3D_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray*', 'array')]
+#define INIT_hipArray3DGetDescriptor_CB_ARGS_DATA(cb_data) { \
+  cb_data.args.hipArray3DGetDescriptor.pArrayDescriptor = (HIP_ARRAY3D_DESCRIPTOR*)pArrayDescriptor; \
+  cb_data.args.hipArray3DGetDescriptor.array = (hipArray*)array; \
+};
 // hipArrayCreate[('hipArray**', 'pHandle'), ('const HIP_ARRAY_DESCRIPTOR*', 'pAllocateArray')]
 #define INIT_hipArrayCreate_CB_ARGS_DATA(cb_data) { \
   cb_data.args.hipArrayCreate.pHandle = (hipArray**)array; \
@@ -3261,6 +3295,18 @@ typedef struct hip_api_data_s {
 #define INIT_hipArrayDestroy_CB_ARGS_DATA(cb_data) { \
   cb_data.args.hipArrayDestroy.array = (hipArray*)array; \
 };
+// hipArrayGetDescriptor[('HIP_ARRAY_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray*', 'array')]
+#define INIT_hipArrayGetDescriptor_CB_ARGS_DATA(cb_data) { \
+  cb_data.args.hipArrayGetDescriptor.pArrayDescriptor = (HIP_ARRAY_DESCRIPTOR*)pArrayDescriptor; \
+  cb_data.args.hipArrayGetDescriptor.array = (hipArray*)array; \
+};
+// hipArrayGetInfo[('hipChannelFormatDesc*', 'desc'), ('hipExtent*', 'extent'), ('unsigned int*', 'flags'), ('hipArray*', 'array')]
+#define INIT_hipArrayGetInfo_CB_ARGS_DATA(cb_data) { \
+  cb_data.args.hipArrayGetInfo.desc = (hipChannelFormatDesc*)desc; \
+  cb_data.args.hipArrayGetInfo.extent = (hipExtent*)extent; \
+  cb_data.args.hipArrayGetInfo.flags = (unsigned int*)flags; \
+  cb_data.args.hipArrayGetInfo.array = (hipArray*)array; \
+};
 // hipChooseDevice[('int*', 'device'), ('const hipDeviceProp_t*', 'prop')]
 #define INIT_hipChooseDevice_CB_ARGS_DATA(cb_data) { \
   cb_data.args.hipChooseDevice.device = (int*)device; \
@@ -5376,10 +5422,6 @@ typedef struct hip_api_data_s {
 #define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)
 
 // Macros for non-public API primitives
-// hipArray3DGetDescriptor()
-#define INIT_hipArray3DGetDescriptor_CB_ARGS_DATA(cb_data) {};
-// hipArrayGetDescriptor()
-#define INIT_hipArrayGetDescriptor_CB_ARGS_DATA(cb_data) {};
 // hipBindTexture()
 #define INIT_hipBindTexture_CB_ARGS_DATA(cb_data) {};
 // hipBindTexture2D()
@@ -5478,6 +5520,11 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
       if (data->args.hipArray3DCreate.array) data->args.hipArray3DCreate.array__val = *(data->args.hipArray3DCreate.array);
       if (data->args.hipArray3DCreate.pAllocateArray) data->args.hipArray3DCreate.pAllocateArray__val = *(data->args.hipArray3DCreate.pAllocateArray);
       break;
+// hipArray3DGetDescriptor[('HIP_ARRAY3D_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray*', 'array')]
+    case HIP_API_ID_hipArray3DGetDescriptor:
+      if (data->args.hipArray3DGetDescriptor.pArrayDescriptor) data->args.hipArray3DGetDescriptor.pArrayDescriptor__val = *(data->args.hipArray3DGetDescriptor.pArrayDescriptor);
+      if (data->args.hipArray3DGetDescriptor.array) data->args.hipArray3DGetDescriptor.array__val = *(data->args.hipArray3DGetDescriptor.array);
+      break;
 // hipArrayCreate[('hipArray**', 'pHandle'), ('const HIP_ARRAY_DESCRIPTOR*', 'pAllocateArray')]
     case HIP_API_ID_hipArrayCreate:
       if (data->args.hipArrayCreate.pHandle) data->args.hipArrayCreate.pHandle__val = *(data->args.hipArrayCreate.pHandle);
@@ -5487,6 +5534,18 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
     case HIP_API_ID_hipArrayDestroy:
       if (data->args.hipArrayDestroy.array) data->args.hipArrayDestroy.array__val = *(data->args.hipArrayDestroy.array);
       break;
+// hipArrayGetDescriptor[('HIP_ARRAY_DESCRIPTOR*', 'pArrayDescriptor'), ('hipArray*', 'array')]
+    case HIP_API_ID_hipArrayGetDescriptor:
+      if (data->args.hipArrayGetDescriptor.pArrayDescriptor) data->args.hipArrayGetDescriptor.pArrayDescriptor__val = *(data->args.hipArrayGetDescriptor.pArrayDescriptor);
+      if (data->args.hipArrayGetDescriptor.array) data->args.hipArrayGetDescriptor.array__val = *(data->args.hipArrayGetDescriptor.array);
+      break;
+// hipArrayGetInfo[('hipChannelFormatDesc*', 'desc'), ('hipExtent*', 'extent'), ('unsigned int*', 'flags'), ('hipArray*', 'array')]
+    case HIP_API_ID_hipArrayGetInfo:
+      if (data->args.hipArrayGetInfo.desc) data->args.hipArrayGetInfo.desc__val = *(data->args.hipArrayGetInfo.desc);
+      if (data->args.hipArrayGetInfo.extent) data->args.hipArrayGetInfo.extent__val = *(data->args.hipArrayGetInfo.extent);
+      if (data->args.hipArrayGetInfo.flags) data->args.hipArrayGetInfo.flags__val = *(data->args.hipArrayGetInfo.flags);
+      if (data->args.hipArrayGetInfo.array) data->args.hipArrayGetInfo.array__val = *(data->args.hipArrayGetInfo.array);
+      break;
 // hipChooseDevice[('int*', 'device'), ('const hipDeviceProp_t*', 'prop')]
     case HIP_API_ID_hipChooseDevice:
       if (data->args.hipChooseDevice.device) data->args.hipChooseDevice.device__val = *(data->args.hipChooseDevice.device);
@@ -6881,6 +6940,14 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
       else { oss << ", pAllocateArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DCreate.pAllocateArray__val); }
       oss << ")";
     break;
+    case HIP_API_ID_hipArray3DGetDescriptor:
+      oss << "hipArray3DGetDescriptor(";
+      if (data->args.hipArray3DGetDescriptor.pArrayDescriptor == NULL) oss << "pArrayDescriptor=NULL";
+      else { oss << "pArrayDescriptor="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DGetDescriptor.pArrayDescriptor__val); }
+      if (data->args.hipArray3DGetDescriptor.array == NULL) oss << ", array=NULL";
+      else { oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArray3DGetDescriptor.array__val); }
+      oss << ")";
+    break;
     case HIP_API_ID_hipArrayCreate:
       oss << "hipArrayCreate(";
       if (data->args.hipArrayCreate.pHandle == NULL) oss << "pHandle=NULL";
@@ -6895,6 +6962,26 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
       else { oss << "array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayDestroy.array__val); }
       oss << ")";
     break;
+    case HIP_API_ID_hipArrayGetDescriptor:
+      oss << "hipArrayGetDescriptor(";
+      if (data->args.hipArrayGetDescriptor.pArrayDescriptor == NULL) oss << "pArrayDescriptor=NULL";
+      else { oss << "pArrayDescriptor="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetDescriptor.pArrayDescriptor__val); }
+      if (data->args.hipArrayGetDescriptor.array == NULL) oss << ", array=NULL";
+      else { oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetDescriptor.array__val); }
+      oss << ")";
+    break;
+    case HIP_API_ID_hipArrayGetInfo:
+      oss << "hipArrayGetInfo(";
+      if (data->args.hipArrayGetInfo.desc == NULL) oss << "desc=NULL";
+      else { oss << "desc="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.desc__val); }
+      if (data->args.hipArrayGetInfo.extent == NULL) oss << ", extent=NULL";
+      else { oss << ", extent="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.extent__val); }
+      if (data->args.hipArrayGetInfo.flags == NULL) oss << ", flags=NULL";
+      else { oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.flags__val); }
+      if (data->args.hipArrayGetInfo.array == NULL) oss << ", array=NULL";
+      else { oss << ", array="; roctracer::hip_support::detail::operator<<(oss, data->args.hipArrayGetInfo.array__val); }
+      oss << ")";
+    break;
     case HIP_API_ID_hipChooseDevice:
       oss << "hipChooseDevice(";
       if (data->args.hipChooseDevice.device == NULL) oss << "device=NULL";
diff --git a/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
index 0c492b7c66..a06f9b45b9 100644
--- a/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
+++ b/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
@@ -3186,6 +3186,21 @@ inline static hipError_t hipArray3DCreate(hiparray* pHandle,
     return hipCUResultTohipError(cuArray3DCreate(pHandle, pAllocateArray));
 }
 
+inline static hipError_t hipArrayGetInfo(hipChannelFormatDesc* desc, hipExtent* extent,
+                                          unsigned int* flags, hipArray* array) {
+    return hipCUDAErrorTohipError(cudaArrayGetInfo(desc, extent, flags, array));
+}
+
+inline static hipError_t hipArrayGetDescriptor(HIP_ARRAY_DESCRIPTOR* pArrayDescriptor,
+                                               hipArray* array) {
+    return hipCUResultTohipError(cuArrayGetDescriptor(pArrayDescriptor, (CUarray)array));
+}
+
+inline static hipError_t hipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor,
+                                                 hipArray* array) {
+    return hipCUResultTohipError(cuArray3DGetDescriptor(pArrayDescriptor, (CUarray)array));
+}
+
 inline static hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) {
     return hipCUDAErrorTohipError(cudaStreamBeginCapture(stream, mode));
 }
diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def
index 9a1d3e6e49..70279e21fe 100644
--- a/hipamd/src/amdhip.def
+++ b/hipamd/src/amdhip.def
@@ -103,6 +103,9 @@ hipMemPoolImportPointer
 hipArrayCreate
 hipArray3DCreate
 hipArrayDestroy
+hipArrayGetInfo
+hipArrayGetDescriptor
+hipArray3DGetDescriptor
 hipMallocArray
 hipMemAdvise
 hipMemAllocPitch
diff --git a/hipamd/src/hip_hcc.def.in b/hipamd/src/hip_hcc.def.in
index 187c0fdd89..129fa7c4d9 100644
--- a/hipamd/src/hip_hcc.def.in
+++ b/hipamd/src/hip_hcc.def.in
@@ -103,6 +103,9 @@ hipMemPoolImportPointer
 hipArrayCreate
 hipArray3DCreate
 hipArrayDestroy
+hipArrayGetInfo
+hipArrayGetDescriptor
+hipArray3DGetDescriptor
 hipMallocArray
 hipMemAdvise
 hipMemAllocPitch
diff --git a/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in
index a9adf3722a..81251ccaaf 100644
--- a/hipamd/src/hip_hcc.map.in
+++ b/hipamd/src/hip_hcc.map.in
@@ -517,3 +517,12 @@ global:
 local:
     *;
 } hip_5.3;
+
+hip_5.6 {
+global:
+    hipArrayGetInfo;
+    hipArrayGetDescriptor;
+    hipArray3DGetDescriptor;
+local:
+    *;
+} hip_5.5;
\ No newline at end of file
diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index 2fdb28b537..40f5a52845 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -1047,6 +1047,7 @@ hipError_t ihipArrayCreate(hipArray** array,
   (*array)->depth = pAllocateArray->Depth;
   (*array)->Format = pAllocateArray->Format;
   (*array)->NumChannels = pAllocateArray->NumChannels;
+  (*array)->flags = pAllocateArray->Flags;
   {
     amd::ScopedLock lock(hip::hipArraySetLock);
     hip::hipArraySet.insert(*array);
@@ -3541,22 +3542,105 @@ hipError_t hipArrayDestroy(hipArray* array) {
   HIP_RETURN(ihipArrayDestroy(array));
 }
 
-hipError_t hipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor,
-                                   hipArray* array) {
-  HIP_INIT_API(hipArray3DGetDescriptor, pArrayDescriptor, array);
+hipError_t ihipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* desc,
+                                    hipArray* array) {
+  {
+    amd::ScopedLock lock(hip::hipArraySetLock);
+    if (hip::hipArraySet.find(array) == hip::hipArraySet.end()) {
+      return hipErrorInvalidHandle;
+    }
+  }
 
-  assert(false && "Unimplemented");
+  desc->Width = array->width;
+  desc->Height = array->height;
+  desc->Depth = array->depth;
+  desc->Format = array->Format;
+  desc->NumChannels = array->NumChannels;
+  desc->Flags = array->flags;
 
-  HIP_RETURN(hipSuccess);
+  return hipSuccess;
+}
+
+hipError_t hipArrayGetInfo(hipChannelFormatDesc* desc,
+                           hipExtent* extent,
+                           unsigned int* flags,
+                           hipArray* array) {
+  HIP_INIT_API(hipArrayGetInfo, desc, extent, flags, array);
+  CHECK_STREAM_CAPTURE_SUPPORTED();
+
+  if (array == nullptr) {
+    HIP_RETURN(hipErrorInvalidHandle);
+  }
+
+  // If all output parameters are nullptr, then no need to proceed further
+  if ((desc == nullptr) && (extent == nullptr) && (flags == nullptr)) {
+    HIP_RETURN(hipSuccess);
+  }
+
+  HIP_ARRAY3D_DESCRIPTOR array3DDescriptor;
+  hipError_t status = ihipArray3DGetDescriptor(&array3DDescriptor, array);
+
+  // Fill each output parameter
+  if (status == hipSuccess) {
+    if (desc != nullptr) {
+      *desc = hip::getChannelFormatDesc(array3DDescriptor.NumChannels, array3DDescriptor.Format);
+    }
+
+    if (extent != nullptr) {
+      extent->width = array3DDescriptor.Width;
+      extent->height = array3DDescriptor.Height;
+      extent->depth = array3DDescriptor.Depth;
+    }
+
+    if (flags != nullptr) {
+      *flags = array3DDescriptor.Flags;
+    }
+  }
+
+  HIP_RETURN(status);
 }
 
 hipError_t hipArrayGetDescriptor(HIP_ARRAY_DESCRIPTOR* pArrayDescriptor,
                                  hipArray* array) {
   HIP_INIT_API(hipArrayGetDescriptor, pArrayDescriptor, array);
+  CHECK_STREAM_CAPTURE_SUPPORTED();
 
-  assert(false && "Unimplemented");
+  if (array == nullptr) {
+    HIP_RETURN(hipErrorInvalidHandle);
+  }
 
-  HIP_RETURN(hipSuccess);
+  if (pArrayDescriptor == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  HIP_ARRAY3D_DESCRIPTOR array3DDescriptor;
+  hipError_t status = ihipArray3DGetDescriptor(&array3DDescriptor, array);
+
+  // Fill each output parameter
+  if (status == hipSuccess) {
+    pArrayDescriptor->Width = array3DDescriptor.Width;
+    pArrayDescriptor->Height = array3DDescriptor.Height;
+    pArrayDescriptor->Format = array3DDescriptor.Format;
+    pArrayDescriptor->NumChannels = array3DDescriptor.NumChannels;
+  }
+
+  HIP_RETURN(status);
+}
+
+hipError_t hipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor,
+                                   hipArray* array) {
+  HIP_INIT_API(hipArray3DGetDescriptor, pArrayDescriptor, array);
+  CHECK_STREAM_CAPTURE_SUPPORTED();
+
+  if (array == nullptr) {
+    HIP_RETURN(hipErrorInvalidHandle);
+  }
+
+  if (pArrayDescriptor == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  HIP_RETURN(ihipArray3DGetDescriptor(pArrayDescriptor, array));
 }
 
 hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy,