diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in
index 8522f3b7fb..ab86e6e35b 100644
--- a/projects/clr/hipamd/src/hip_hcc.map.in
+++ b/projects/clr/hipamd/src/hip_hcc.map.in
@@ -261,13 +261,6 @@ global:
     hipTexObjectGetResourceDesc;
     hipTexObjectGetResourceViewDesc;
     hipTexObjectGetTextureDesc;
-    extern "C++" {
-    hip_impl::hipLaunchKernelGGLImpl*;
-    hip_impl::demangle*;
-    hipCreateSurfaceObject*;
-    hipDestroySurfaceObject*;
-    hipHccModuleLaunchKernel*;
-    hipExtModuleLaunchKernel*;
     hipGetCmdName*;
     hipExtStreamCreateWithCUMask;
     hipStreamGetPriority;
@@ -278,6 +271,12 @@ global:
     hipMemcpyParam2DAsync;
     __hipGetPCH;
     hipExtStreamGetCUMask;
+    extern "C++" {
+    hipCreateSurfaceObject*;
+    hipDestroySurfaceObject*;
+    hipHccModuleLaunchKernel*;
+    hipExtModuleLaunchKernel*;
+
     };
 local:
     *;
diff --git a/projects/clr/hipamd/src/hip_platform.cpp b/projects/clr/hipamd/src/hip_platform.cpp
index 6a16d45ed2..8f532361bc 100644
--- a/projects/clr/hipamd/src/hip_platform.cpp
+++ b/projects/clr/hipamd/src/hip_platform.cpp
@@ -564,49 +564,6 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks,
 }
 }
 
-
-#if defined(ATI_OS_LINUX)
-
-namespace hip_impl {
-
-void hipLaunchKernelGGLImpl(uintptr_t function_address, const dim3& numBlocks,
-                            const dim3& dimBlocks, uint32_t sharedMemBytes, hipStream_t stream,
-                            void** kernarg) {
-  HIP_INIT_VOID();
-
-  hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
-  int deviceId = (s != nullptr) ? s->DeviceId() : ihipGetDevice();
-  if (deviceId == -1) {
-    LogPrintfError("Wrong Device Id: %d \n", deviceId);
-  }
-
-  hipFunction_t func = nullptr;
-  hipError_t hip_error = PlatformState::instance().getStatFunc(
-      &func, reinterpret_cast<void*>(function_address), deviceId);
-  if ((hip_error != hipSuccess) || (func == nullptr)) {
-    LogPrintfError("Cannot find the static function: 0x%x", function_address);
-  }
-
-  hip_error =
-      hipModuleLaunchKernel(func, numBlocks.x, numBlocks.y, numBlocks.z, dimBlocks.x, dimBlocks.y,
-                            dimBlocks.z, sharedMemBytes, stream, nullptr, kernarg);
-  assert(hip_error == hipSuccess);
-}
-
-void hipLaunchCooperativeKernelGGLImpl(uintptr_t function_address, const dim3& numBlocks,
-                                       const dim3& dimBlocks, uint32_t sharedMemBytes,
-                                       hipStream_t stream, void** kernarg) {
-  HIP_INIT_VOID();
-
-  hipError_t err = hipLaunchCooperativeKernel(reinterpret_cast<void*>(function_address), numBlocks,
-                                              dimBlocks, kernarg, sharedMemBytes, stream);
-  assert(err == hipSuccess);
-}
-
-}  // namespace hip_impl
-
-#endif  // defined(ATI_OS_LINUX)
-
 hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim, void** args,
                             size_t sharedMemBytes, hipStream_t stream, hipEvent_t startEvent,
                             hipEvent_t stopEvent, int flags) {