diff --git a/api/hip/hip_hcc.def.in b/api/hip/hip_hcc.def.in
index d2916f3885..bb9c179d3b 100644
--- a/api/hip/hip_hcc.def.in
+++ b/api/hip/hip_hcc.def.in
@@ -100,6 +100,7 @@ hipMemset2D
 hipMemsetAsync
 hipMemset2DAsync
 hipMemsetD8
+hipMemset3D
 hipModuleGetFunction
 hipModuleGetGlobal
 hipModuleLaunchKernel
diff --git a/api/hip/hip_hcc.map.in b/api/hip/hip_hcc.map.in
index 761641425e..452f260534 100644
--- a/api/hip/hip_hcc.map.in
+++ b/api/hip/hip_hcc.map.in
@@ -101,6 +101,7 @@ global:
     hipMemsetAsync;
     hipMemset2DAsync;
     hipMemsetD8;
+    hipMemset3D;
     hipModuleGetFunction;
     hipModuleGetGlobal;
     hipModuleLaunchKernel;
diff --git a/api/hip/hip_memory.cpp b/api/hip/hip_memory.cpp
index cd4c60c319..376b2b155f 100644
--- a/api/hip/hip_memory.cpp
+++ b/api/hip/hip_memory.cpp
@@ -363,6 +363,18 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) {
   return status;
 }
 
+hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) {
+  HIP_INIT_API(pitchedDevPtr, value, &extent);
+
+  void *dst = &pitchedDevPtr.ptr;
+  size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
+
+  hip::syncStreams();
+  amd::HostQueue* queue = hip::getNullStream();
+
+  return ihipMemset(&dst, value, sizeBytes, *queue);
+}
+
 hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {
   HIP_INIT_API(array, pAllocateArray);