SWDEV-433745 - Added demangled symbol for hipExtModuleLaunchKernel

Change-Id: Ib8051fcbdf0c8c99e97c845b777d6562220701a0


[ROCm/clr commit: fe739047e5]

projects/clr/hipamd/src/CMakeLists.txt

@@ -124,7 +124,8 @@ target_sources(amdhip64 PRIVATE
   hip_gl.cpp
   hip_vm.cpp
   hip_api_trace.cpp
-  hip_table_interface.cpp)
+  hip_table_interface.cpp
+  hip_table_interface_c.cpp)
 
 if(WIN32)
   target_sources(amdhip64 PRIVATE

projects/clr/hipamd/src/hip_table_interface.cpp

@@ -1518,7 +1518,14 @@ extern "C" hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w,
                                                      hipChannelFormatKind f) {
   return hip::GetHipDispatchTable()->hipCreateChannelDesc_fn(x, y, z, w, f);
 }
-hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
+
+#ifdef _WIN32
+# define DllExport __declspec(dllexport)
+#else // !_WIN32
+# define DllExport
+#endif // !_WIN32
+
+DllExport hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                     uint32_t localWorkSizeX, uint32_t localWorkSizeY,
                                     uint32_t localWorkSizeZ, size_t sharedMemBytes,
@@ -1528,7 +1535,8 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
       f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
       localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags);
 }
-hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
+
+DllExport hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                     uint32_t localWorkSizeX, uint32_t localWorkSizeY,
                                     uint32_t localWorkSizeZ, size_t sharedMemBytes,
@@ -1538,6 +1546,7 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
       f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
       localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent);
 }
+
 hipError_t hipMemcpy_spt(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
   return hip::GetHipDispatchTable()->hipMemcpy_spt_fn(dst, src, sizeBytes, kind);
 }

projects/clr/hipamd/src/hip_table_interface_c.cpp

@@ -0,0 +1,53 @@
+/*
+    Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+    Permission is hereby granted, free of charge, to any person obtaining a copy
+    of this software and associated documentation files (the "Software"), to deal
+    in the Software without restriction, including without limitation the rights
+    to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+    copies of the Software, and to permit persons to whom the Software is
+    furnished to do so, subject to the following conditions:
+
+    The above copyright notice and this permission notice shall be included in
+    all copies or substantial portions of the Software.
+
+    THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+    IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+    FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+    AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+    LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+    OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+    THE SOFTWARE.
+   */
+#include <hip/amd_detail/hip_api_trace.hpp>
+namespace hip {
+const HipDispatchTable* GetHipDispatchTable();
+const HipCompilerDispatchTable* GetHipCompilerDispatchTable();
+}  // namespace hip
+
+#ifdef _WIN32
+# define DllExport extern "C" __declspec(dllexport)
+#else // !_WIN32
+# define DllExport extern "C"
+#endif // !_WIN32
+
+DllExport hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
+                                    uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
+                                    uint32_t localWorkSizeX, uint32_t localWorkSizeY,
+                                    uint32_t localWorkSizeZ, size_t sharedMemBytes,
+                                    hipStream_t hStream, void** kernelParams, void** extra,
+                                    hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) {
+  return hip::GetHipDispatchTable()->hipExtModuleLaunchKernel_fn(
+      f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
+      localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags);
+}
+DllExport hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
+                                    uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
+                                    uint32_t localWorkSizeX, uint32_t localWorkSizeY,
+                                    uint32_t localWorkSizeZ, size_t sharedMemBytes,
+                                    hipStream_t hStream, void** kernelParams, void** extra,
+                                    hipEvent_t startEvent, hipEvent_t stopEvent) {
+  return hip::GetHipDispatchTable()->hipHccModuleLaunchKernel_fn(
+      f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
+      localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent);
+}