From de10e7e1e6e0dbbc2f5dca0e6b57f24e4bc3be14 Mon Sep 17 00:00:00 2001 From: Anusha Godavarthy Surya Date: Thu, 16 Jul 2020 15:17:51 -0400 Subject: [PATCH] SWDEV-244600 - HIP BLIT code object needs to have reserved symbol name Change-Id: I8401fea5eab71c0f7414eec0666066d9553a6622 [ROCm/clr commit: 093f7fa3cad924cf53df827f42ec493a8ff4e905] --- projects/clr/rocclr/device/blitcl.cpp | 20 +++++++++---------- projects/clr/rocclr/device/gpu/gpublit.hpp | 10 +++++----- projects/clr/rocclr/device/pal/palblit.hpp | 10 +++++----- projects/clr/rocclr/device/pal/palschedcl.cpp | 4 ++-- projects/clr/rocclr/device/rocm/rocblit.hpp | 10 +++++----- .../clr/rocclr/device/rocm/rocschedcl.cpp | 4 ++-- 6 files changed, 29 insertions(+), 29 deletions(-) diff --git a/projects/clr/rocclr/device/blitcl.cpp b/projects/clr/rocclr/device/blitcl.cpp index 12097b9853..c328c30c3d 100644 --- a/projects/clr/rocclr/device/blitcl.cpp +++ b/projects/clr/rocclr/device/blitcl.cpp @@ -34,27 +34,27 @@ const char* BlitSourceCode = BLIT_KERNELS( extern void __amd_fillBuffer(__global uchar*, __global uint*, __constant uchar*, uint, ulong, ulong); - __kernel void copyBufferRect(__global uchar* src, __global uchar* dst, ulong4 srcRect, + __kernel void __amd_rocclr_copyBufferRect(__global uchar* src, __global uchar* dst, ulong4 srcRect, ulong4 dstRect, ulong4 size) { __amd_copyBufferRect(src, dst, srcRect, dstRect, size); } - __kernel void copyBufferRectAligned(__global uint* src, __global uint* dst, ulong4 srcRect, + __kernel void __amd_rocclr_copyBufferRectAligned(__global uint* src, __global uint* dst, ulong4 srcRect, ulong4 dstRect, ulong4 size) { __amd_copyBufferRectAligned(src, dst, srcRect, dstRect, size); } - __kernel void copyBuffer(__global uchar* srcI, __global uchar* dstI, ulong srcOrigin, + __kernel void __amd_rocclr_copyBuffer(__global uchar* srcI, __global uchar* dstI, ulong srcOrigin, ulong dstOrigin, ulong size, uint remain) { __amd_copyBuffer(srcI, dstI, srcOrigin, dstOrigin, size, remain); } - __kernel void copyBufferAligned(__global uint* src, __global uint* dst, ulong srcOrigin, + __kernel void __amd_rocclr_copyBufferAligned(__global uint* src, __global uint* dst, ulong srcOrigin, ulong dstOrigin, ulong size, uint alignment) { __amd_copyBufferAligned(src, dst, srcOrigin, dstOrigin, size, alignment); } - __kernel void fillBuffer(__global uchar* bufUChar, __global uint* bufUInt, + __kernel void __amd_rocclr_fillBuffer(__global uchar* bufUChar, __global uint* bufUInt, __constant uchar* pattern, uint patternSize, ulong offset, ulong size) { __amd_fillBuffer(bufUChar, bufUInt, pattern, patternSize, offset, size); @@ -75,13 +75,13 @@ const char* BlitSourceCode = BLIT_KERNELS( uint); - __kernel void copyBufferToImage(__global uint* src, __write_only image2d_array_t dst, + __kernel void __amd_rocclr_copyBufferToImage(__global uint* src, __write_only image2d_array_t dst, ulong4 srcOrigin, int4 dstOrigin, int4 size, uint4 format, ulong4 pitch) { __amd_copyBufferToImage(src, dst, srcOrigin, dstOrigin, size, format, pitch); } - __kernel void copyImageToBuffer(__read_only image2d_array_t src, __global uint* dstUInt, + __kernel void __amd_rocclr_copyImageToBuffer(__read_only image2d_array_t src, __global uint* dstUInt, __global ushort* dstUShort, __global uchar* dstUChar, int4 srcOrigin, ulong4 dstOrigin, int4 size, uint4 format, ulong4 pitch) { @@ -89,16 +89,16 @@ const char* BlitSourceCode = BLIT_KERNELS( pitch); } - __kernel void copyImage(__read_only image2d_array_t src, __write_only image2d_array_t dst, + __kernel void __amd_rocclr_copyImage(__read_only image2d_array_t src, __write_only image2d_array_t dst, int4 srcOrigin, int4 dstOrigin, int4 size) { __amd_copyImage(src, dst, srcOrigin, dstOrigin, size); } - __kernel void copyImage1DA(__read_only image2d_array_t src, __write_only image2d_array_t dst, + __kernel void __amd_rocclr_copyImage1DA(__read_only image2d_array_t src, __write_only image2d_array_t dst, int4 srcOrigin, int4 dstOrigin, int4 size) { __amd_copyImage1DA(src, dst, srcOrigin, dstOrigin, size); } - __kernel void fillImage(__write_only image2d_array_t image, float4 patternFLOAT4, + __kernel void __amd_rocclr_fillImage(__write_only image2d_array_t image, float4 patternFLOAT4, int4 patternINT4, uint4 patternUINT4, int4 origin, int4 size, uint type) { __amd_fillImage(image, patternFLOAT4, patternINT4, patternUINT4, origin, size, type); diff --git a/projects/clr/rocclr/device/gpu/gpublit.hpp b/projects/clr/rocclr/device/gpu/gpublit.hpp index b9b8f70023..577bd3f59b 100644 --- a/projects/clr/rocclr/device/gpu/gpublit.hpp +++ b/projects/clr/rocclr/device/gpu/gpublit.hpp @@ -413,11 +413,11 @@ class KernelBlitManager : public DmaBlitManager { amd::Monitor* lockXferOps_; //!< Lock transfer operation }; -static constexpr const char* BlitName[KernelBlitManager::BlitTotal] = { - "copyImage", "copyImage1DA", "copyImageToBuffer", - "copyBufferToImage", "copyBufferRect", "copyBufferRectAligned", - "copyBuffer", "copyBufferAligned", "fillBuffer", - "fillImage", "scheduler", +static const char* BlitName[KernelBlitManager::BlitTotal] = { + "__amd_rocclr_copyImage", "__amd_rocclr_copyImage1DA", "__amd_rocclr_copyImageToBuffer", + "__amd_rocclr_copyBufferToImage", "__amd_rocclr_copyBufferRect", "__amd_rocclr_copyBufferRectAligned", + "__amd_rocclr_copyBuffer", "__amd_rocclr_copyBufferAligned", "__amd_rocclr_fillBuffer", + "__amd_rocclr_fillImage", "__amd_rocclr_scheduler", }; /*@}*/} // namespace gpu diff --git a/projects/clr/rocclr/device/pal/palblit.hpp b/projects/clr/rocclr/device/pal/palblit.hpp index cdca86fa16..afb282f661 100644 --- a/projects/clr/rocclr/device/pal/palblit.hpp +++ b/projects/clr/rocclr/device/pal/palblit.hpp @@ -424,11 +424,11 @@ class KernelBlitManager : public DmaBlitManager { mutable amd::Monitor lockXferOps_; //!< Lock transfer operation }; -static constexpr const char* BlitName[KernelBlitManager::BlitTotal] = { - "copyImage", "copyImage1DA", "copyImageToBuffer", - "copyBufferToImage", "copyBufferRect", "copyBufferRectAligned", - "copyBuffer", "copyBufferAligned", "fillBuffer", - "fillImage", "scheduler", "gwsInit" +static const char* BlitName[KernelBlitManager::BlitTotal] = { + "__amd_rocclr_copyImage", "__amd_rocclr_copyImage1DA", "__amd_rocclr_copyImageToBuffer", + "__amd_rocclr_copyBufferToImage", "__amd_rocclr_copyBufferRect", "__amd_rocclr_copyBufferRectAligned", + "__amd_rocclr_copyBuffer", "__amd_rocclr_copyBufferAligned", "__amd_rocclr_fillBuffer", + "__amd_rocclr_fillImage", "__amd_rocclr_scheduler", "__amd_rocclr_gwsInit" }; /*@}*/ // namespace pal diff --git a/projects/clr/rocclr/device/pal/palschedcl.cpp b/projects/clr/rocclr/device/pal/palschedcl.cpp index 26e878621a..672144d8e1 100644 --- a/projects/clr/rocclr/device/pal/palschedcl.cpp +++ b/projects/clr/rocclr/device/pal/palschedcl.cpp @@ -27,7 +27,7 @@ const char* SchedulerSourceCode = BLIT_KERNEL( \n extern void __amd_scheduler(__global void*, __global void*, uint); \n -__kernel void scheduler(__global void* queue, __global void* params, uint paramIdx) { +__kernel void __amd_rocclr_scheduler(__global void* queue, __global void* params, uint paramIdx) { __amd_scheduler(queue, params, paramIdx); } \n); @@ -36,7 +36,7 @@ const char* GwsInitSourceCode = BLIT_KERNEL( \n extern void __ockl_gws_init(uint nwm1, uint rid); \n -__kernel void gwsInit(uint value) { +__kernel void __amd_rocclr_gwsInit(uint value) { __ockl_gws_init(value, 0); } \n); diff --git a/projects/clr/rocclr/device/rocm/rocblit.hpp b/projects/clr/rocclr/device/rocm/rocblit.hpp index b12614bd76..ee6c3420a1 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.hpp +++ b/projects/clr/rocclr/device/rocm/rocblit.hpp @@ -471,11 +471,11 @@ class KernelBlitManager : public DmaBlitManager { mutable amd::Monitor lockXferOps_; //!< Lock transfer operation }; -static constexpr const char* BlitName[KernelBlitManager::BlitTotal] = { - "copyImage", "copyImage1DA", "copyImageToBuffer", - "copyBufferToImage", "copyBufferRect", "copyBufferRectAligned", - "copyBuffer", "copyBufferAligned", "fillBuffer", - "fillImage", "scheduler", "gwsInit" +static const char* BlitName[KernelBlitManager::BlitTotal] = { + "__amd_rocclr_copyImage", "__amd_rocclr_copyImage1DA", "__amd_rocclr_copyImageToBuffer", + "__amd_rocclr_copyBufferToImage", "__amd_rocclr_copyBufferRect", "__amd_rocclr_copyBufferRectAligned", + "__amd_rocclr_copyBuffer", "__amd_rocclr_copyBufferAligned", "__amd_rocclr_fillBuffer", + "__amd_rocclr_fillImage", "__amd_rocclr_scheduler", "__amd_rocclr_gwsInit" }; inline void KernelBlitManager::setArgument(amd::Kernel* kernel, size_t index, diff --git a/projects/clr/rocclr/device/rocm/rocschedcl.cpp b/projects/clr/rocclr/device/rocm/rocschedcl.cpp index 94b3a76940..f06d300556 100644 --- a/projects/clr/rocclr/device/rocm/rocschedcl.cpp +++ b/projects/clr/rocclr/device/rocm/rocschedcl.cpp @@ -26,7 +26,7 @@ const char* SchedulerSourceCode = BLIT_KERNEL( \n extern void __amd_scheduler_rocm(__global void*); \n -__kernel void scheduler(__global void* params) { +__kernel void __amd_rocclr_scheduler(__global void* params) { __amd_scheduler_rocm(params); } \n); @@ -35,7 +35,7 @@ const char* GwsInitSourceCode = BLIT_KERNEL( \n extern void __ockl_gws_init(uint nwm1, uint rid); \n -__kernel void gwsInit(uint value) { +__kernel void __amd_rocclr_gwsInit(uint value) { __ockl_gws_init(value, 0); } \n);