diff --git a/projects/hip/api/hip/hip_hcc.def.in b/projects/hip/api/hip/hip_hcc.def.in
index 7d359dbf28..e3ddb529a3 100644
--- a/projects/hip/api/hip/hip_hcc.def.in
+++ b/projects/hip/api/hip/hip_hcc.def.in
@@ -141,6 +141,8 @@ hipStreamGetFlags
 hipStreamQuery
 hipStreamSynchronize
 hipStreamWaitEvent
+__hipPopCallConfiguration
+__hipPushCallConfiguration
 __hipRegisterFatBinary
 __hipRegisterFunction
 __hipRegisterVar
@@ -148,6 +150,7 @@ __hipUnregisterFatBinary
 hipConfigureCall
 hipSetupArgument
 hipLaunchByPtr
+hipLaunchKernel
 hipCreateTextureObject
 hipDestroyTextureObject
 hipGetTextureObjectResourceDesc
diff --git a/projects/hip/api/hip/hip_hcc.map.in b/projects/hip/api/hip/hip_hcc.map.in
index afe47e8df1..d580b8b99b 100644
--- a/projects/hip/api/hip/hip_hcc.map.in
+++ b/projects/hip/api/hip/hip_hcc.map.in
@@ -141,6 +141,8 @@ global:
     hipStreamQuery;
     hipStreamSynchronize;
     hipStreamWaitEvent;
+    __hipPopCallConfiguration;
+    __hipPushCallConfiguration;
     __hipRegisterFatBinary;
     __hipRegisterFunction;
     __hipRegisterVar;
@@ -150,6 +152,7 @@ global:
     hipConfigureCall;
     hipSetupArgument;
     hipLaunchByPtr;
+    hipLaunchKernel;
     hipProfilerStart;
     hipProfilerStop;
     hiprtcCompileProgram;
diff --git a/projects/hip/api/hip/hip_platform.cpp b/projects/hip/api/hip/hip_platform.cpp
index ee05e7c52a..8e3fd32b66 100644
--- a/projects/hip/api/hip/hip_platform.cpp
+++ b/projects/hip/api/hip/hip_platform.cpp
@@ -411,6 +411,35 @@ extern "C" hipError_t hipConfigureCall(
   HIP_RETURN(hipSuccess);
 }
 
+extern "C" hipError_t __hipPushCallConfiguration(
+  dim3 gridDim,
+  dim3 blockDim,
+  size_t sharedMem,
+  hipStream_t stream)
+{
+  HIP_INIT_API(gridDim, blockDim, sharedMem, stream);
+
+  PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream);
+
+  HIP_RETURN(hipSuccess);
+}
+
+extern "C" hipError_t __hipPopCallConfiguration(dim3 *gridDim,
+                                                dim3 *blockDim,
+                                                size_t *sharedMem,
+                                                hipStream_t *stream) {
+  HIP_INIT_API(gridDim, blockDim, sharedMem, stream);
+
+  ihipExec_t exec;
+  PlatformState::instance().popExec(exec);
+  *gridDim = exec.gridDim_;
+  *blockDim = exec.blockDim_;
+  *sharedMem = exec.sharedMem_;
+  *stream = exec.hStream_;
+
+  HIP_RETURN(hipSuccess);
+}
+
 extern "C" hipError_t hipSetupArgument(
   const void *arg,
   size_t size,
@@ -449,6 +478,28 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction)
     exec.sharedMem_, exec.hStream_, nullptr, extra));
 }
 
+extern "C" hipError_t hipLaunchKernel(const void *hostFunction,
+                                      dim3 gridDim,
+                                      dim3 blockDim,
+                                      void** args,
+                                      size_t sharedMemBytes,
+                                      hipStream_t stream)
+{
+  HIP_INIT_API(hostFunction, gridDim, blockDim, args, sharedMemBytes,
+               stream);
+
+  int deviceId = ihipGetDevice();
+  hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId);
+  if (func == nullptr) {
+    HIP_RETURN(hipErrorUnknown);
+  }
+
+  HIP_RETURN(hipModuleLaunchKernel(func, gridDim.x, gridDim.y, gridDim.z,
+                                    blockDim.x, blockDim.y, blockDim.z,
+                                    sharedMemBytes, stream, args, nullptr));
+}
+
+
 hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
   size_t size = 0;
   if(!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), devPtr, &size)) {