diff --git a/rocclr/hip_hcc.def.in b/rocclr/hip_hcc.def.in
index 53bdfb3065..a3476c7f33 100755
--- a/rocclr/hip_hcc.def.in
+++ b/rocclr/hip_hcc.def.in
@@ -253,3 +253,4 @@ hipTexObjectGetResourceDesc
 hipTexObjectGetResourceViewDesc
 hipTexObjectGetTextureDesc
 hipExtStreamCreateWithCUMask
+hipStreamGetPriority
diff --git a/rocclr/hip_hcc.map.in b/rocclr/hip_hcc.map.in
index 394a5f5dd3..11637f2696 100755
--- a/rocclr/hip_hcc.map.in
+++ b/rocclr/hip_hcc.map.in
@@ -260,6 +260,7 @@ global:
     hipEnableActivityCallback*;
     hipGetCmdName*;
     hipExtStreamCreateWithCUMask;
+    hipStreamGetPriority;
     };
 local:
     *;
diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp
index 400aa8b99a..702b93c8b3 100755
--- a/rocclr/hip_internal.hpp
+++ b/rocclr/hip_internal.hpp
@@ -109,6 +109,8 @@ namespace hip {
     amd::Monitor& Lock() const { return lock_; }
     /// Returns the creation flags for the current stream
     unsigned int Flags() const { return flags_; }
+    /// Returns the priority for the current stream
+    amd::CommandQueue::Priority Priority() const { return priority_; }
 
     /// Sync all non-blocking streams
     static void syncNonBlockingStreams();
diff --git a/rocclr/hip_stream.cpp b/rocclr/hip_stream.cpp
index bfe80cd6bb..43fce39299 100755
--- a/rocclr/hip_stream.cpp
+++ b/rocclr/hip_stream.cpp
@@ -339,3 +339,15 @@ hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize
   HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal, cuMaskv));
 }
 
+// ================================================================================================
+hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
+  HIP_INIT_API(hipStreamGetPriority, stream, priority);
+  if ((priority != nullptr) && (stream != nullptr)) {
+    *priority = static_cast<int>(reinterpret_cast<hip::Stream*>(stream)->Priority());
+  } else {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  HIP_RETURN(hipSuccess);
+
+}