From fe739047e5299143af0b24c4df8c4353fc7d4d3e Mon Sep 17 00:00:00 2001 From: Ioannis Assiouras Date: Mon, 27 Nov 2023 13:25:15 +0000 Subject: [PATCH] SWDEV-433745 - Added demangled symbol for hipExtModuleLaunchKernel Change-Id: Ib8051fcbdf0c8c99e97c845b777d6562220701a0 --- hipamd/src/CMakeLists.txt | 3 +- hipamd/src/hip_table_interface.cpp | 13 +++++-- hipamd/src/hip_table_interface_c.cpp | 53 ++++++++++++++++++++++++++++ 3 files changed, 66 insertions(+), 3 deletions(-) create mode 100644 hipamd/src/hip_table_interface_c.cpp diff --git a/hipamd/src/CMakeLists.txt b/hipamd/src/CMakeLists.txt index 116c4d459a..fad9d9dbf8 100644 --- a/hipamd/src/CMakeLists.txt +++ b/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 diff --git a/hipamd/src/hip_table_interface.cpp b/hipamd/src/hip_table_interface.cpp index 5d1f6806fa..7989885bbe 100644 --- a/hipamd/src/hip_table_interface.cpp +++ b/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); } diff --git a/hipamd/src/hip_table_interface_c.cpp b/hipamd/src/hip_table_interface_c.cpp new file mode 100644 index 0000000000..5f765300af --- /dev/null +++ b/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 +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); +}