AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-244600 - HIP BLIT code object needs to have reserved symbol name

Browse Source 
Change-Id: I8401fea5eab71c0f7414eec0666066d9553a6622


[ROCm/clr commit: 093f7fa3ca]
This commit is contained in:
Anusha Godavarthy Surya
2020-07-16 15:17:51 -04:00
parent e1b0edf35c
commit de10e7e1e6
6 changed files with 29 additions and 29 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/device/blitcl.cpp
+10 -10
View File
@@ -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);
projects/clr/rocclr/device/gpu/gpublit.hpp
+5 -5
View File
@@ -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
projects/clr/rocclr/device/pal/palblit.hpp
+5 -5
View File
@@ -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
projects/clr/rocclr/device/pal/palschedcl.cpp
+2 -2
View File
@@ -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);
projects/clr/rocclr/device/rocm/rocblit.hpp
+5 -5
View File
@@ -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,
projects/clr/rocclr/device/rocm/rocschedcl.cpp
+2 -2
View File
@@ -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);