From 21ba34d0fec9a60b7ae2d998efc2fce32597cfb9 Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Mon, 30 Aug 2021 12:49:30 -0700 Subject: [PATCH] SWDEV-297448 - Add 64bit and 16bit write support For the fillBuffer shader, if there are two 32bit writes to a MMIO register, it can get dropped. It has to be a single 64bit write. Add optimization to fillBuffer to write 64bit and 16bit writes. Change-Id: I3aa78e027898f8ae01e9c8f09004615673720c2b --- rocclr/device/blitcl.cpp | 21 ++++++++++------ rocclr/device/gpu/gpublit.cpp | 45 ++++++++++++++++++++++++---------- rocclr/device/gpu/gpublit.hpp | 9 ++++--- rocclr/device/pal/palblit.cpp | 45 ++++++++++++++++++++++++---------- rocclr/device/pal/palblit.hpp | 10 +++++--- rocclr/device/rocm/rocblit.cpp | 41 ++++++++++++++++++++++--------- rocclr/device/rocm/rocblit.hpp | 10 +++++--- 7 files changed, 125 insertions(+), 56 deletions(-) diff --git a/rocclr/device/blitcl.cpp b/rocclr/device/blitcl.cpp index 1520a4fb56..54dbb9c17c 100644 --- a/rocclr/device/blitcl.cpp +++ b/rocclr/device/blitcl.cpp @@ -31,8 +31,8 @@ const char* BlitSourceCode = BLIT_KERNELS( extern void __amd_copyBufferAligned(__global uint*, __global uint*, ulong, ulong, ulong, uint); - extern void __amd_fillBuffer(__global uchar*, __global uint*, __constant uchar*, uint, ulong, - ulong); + extern void __amd_fillBufferAligned(__global uchar*, __global ushort*, __global uint*, __global ulong*, + __constant uchar*, uint, ulong, ulong); __kernel void __amd_rocclr_copyBufferRect(__global uchar* src, __global uchar* dst, ulong4 srcRect, ulong4 dstRect, ulong4 size) { @@ -54,11 +54,18 @@ const char* BlitSourceCode = BLIT_KERNELS( __amd_copyBufferAligned(src, dst, srcOrigin, dstOrigin, size, alignment); } - __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); - } extern void __amd_copyBufferToImage(__global uint*, __write_only image2d_array_t, ulong4, + __kernel void __amd_rocclr_fillBufferAligned(__global uchar* bufUChar, + __global ushort* bufUShort, + __global uint* bufUInt, + __global ulong* bufULong, + __constant uchar* pattern, + uint patternSize, ulong offset, + ulong size) { + __amd_fillBufferAligned(bufUChar, bufUShort, bufUInt, bufULong, + pattern, patternSize, offset, size); + } + + extern void __amd_copyBufferToImage(__global uint*, __write_only image2d_array_t, ulong4, int4, int4, uint4, ulong4); extern void __amd_copyImageToBuffer(__read_only image2d_array_t, __global uint*, diff --git a/rocclr/device/gpu/gpublit.cpp b/rocclr/device/gpu/gpublit.cpp index 26bd37e9b1..56d179a2ce 100644 --- a/rocclr/device/gpu/gpublit.cpp +++ b/rocclr/device/gpu/gpublit.cpp @@ -2034,21 +2034,40 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, synchronize(); return result; } else { - uint fillType = FillBuffer; + uint fillType = FillBufferAligned; size_t globalWorkOffset[3] = {0, 0, 0}; uint64_t fillSize = size[0] / patternSize; size_t globalWorkSize = amd::alignUp(fillSize, 256); size_t localWorkSize = 256; - bool dwordAligned = ((patternSize % sizeof(uint32_t)) == 0) ? true : false; + uint32_t alignment = (patternSize & 0x7) == 0 ? + sizeof(uint64_t) : + (patternSize & 0x3) == 0 ? + sizeof(uint32_t) : + (patternSize & 0x1) == 0 ? + sizeof(uint16_t) : sizeof(uint8_t); // Program kernels arguments for the fill operation Memory* mem = &gpuMem(memory); - if (dwordAligned) { - setArgument(kernels_[fillType], 0, sizeof(cl_mem), NULL); + if (alignment == sizeof(uint64_t)) { + setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), &mem); + } else if (alignment == sizeof(uint32_t)) { + setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), &mem); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); + } else if (alignment == sizeof(uint16_t)) { + setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); setArgument(kernels_[fillType], 1, sizeof(cl_mem), &mem); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); } else { setArgument(kernels_[fillType], 0, sizeof(cl_mem), &mem); - setArgument(kernels_[fillType], 1, sizeof(cl_mem), NULL); + setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); } Memory* gpuCB = dev().getGpuMemory(constantBuffer_); if (gpuCB == NULL) { @@ -2057,15 +2076,15 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, void* constBuf = gpuCB->map(&gpu(), Resource::WriteOnly); memcpy(constBuf, pattern, patternSize); gpuCB->unmap(&gpu()); - setArgument(kernels_[fillType], 2, sizeof(cl_mem), &gpuCB); + setArgument(kernels_[fillType], 4, sizeof(cl_mem), &gpuCB); uint64_t offset = origin[0]; - if (dwordAligned) { - patternSize /= sizeof(uint32_t); - offset /= sizeof(uint32_t); - } - setArgument(kernels_[fillType], 3, sizeof(uint32_t), &patternSize); - setArgument(kernels_[fillType], 4, sizeof(offset), &offset); - setArgument(kernels_[fillType], 5, sizeof(fillSize), &fillSize); + + patternSize/= alignment; + offset /= alignment; + + setArgument(kernels_[fillType], 5, sizeof(uint32_t), &patternSize); + setArgument(kernels_[fillType], 6, sizeof(offset), &offset); + setArgument(kernels_[fillType], 7, sizeof(fillSize), &fillSize); // Create ND range object for the kernel's execution amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize); diff --git a/rocclr/device/gpu/gpublit.hpp b/rocclr/device/gpu/gpublit.hpp index 0458831625..3dc9e977f9 100644 --- a/rocclr/device/gpu/gpublit.hpp +++ b/rocclr/device/gpu/gpublit.hpp @@ -220,7 +220,7 @@ class KernelBlitManager : public DmaBlitManager { BlitCopyBufferRectAligned, BlitCopyBuffer, BlitCopyBufferAligned, - FillBuffer, + FillBufferAligned, FillImage, Scheduler, BlitTotal @@ -416,9 +416,10 @@ class KernelBlitManager : public DmaBlitManager { 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_copyBufferToImage", "__amd_rocclr_copyBufferRect", + "__amd_rocclr_copyBufferRectAligned", "__amd_rocclr_copyBuffer", + "__amd_rocclr_copyBufferAligned", "__amd_rocclr_fillBufferAligned", + "__amd_rocclr_fillImage", "__amd_rocclr_scheduler" }; /*@}*/} // namespace gpu diff --git a/rocclr/device/pal/palblit.cpp b/rocclr/device/pal/palblit.cpp index ad64ced842..10733f47a7 100644 --- a/rocclr/device/pal/palblit.cpp +++ b/rocclr/device/pal/palblit.cpp @@ -2114,36 +2114,55 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, synchronize(); return result; } else { - uint fillType = FillBuffer; + uint fillType = FillBufferAligned; size_t globalWorkOffset[3] = {0, 0, 0}; uint64_t fillSize = size[0] / patternSize; size_t globalWorkSize = amd::alignUp(fillSize, 256); size_t localWorkSize = 256; - bool dwordAligned = ((patternSize % sizeof(uint32_t)) == 0) ? true : false; + uint32_t alignment = (patternSize & 0x7) == 0 ? + sizeof(uint64_t) : + (patternSize & 0x3) == 0 ? + sizeof(uint32_t) : + (patternSize & 0x1) == 0 ? + sizeof(uint16_t) : sizeof(uint8_t); // Program kernels arguments for the fill operation Memory* mem = &gpuMem(memory); - if (dwordAligned) { - setArgument(kernels_[fillType], 0, sizeof(cl_mem), NULL); + if (alignment == sizeof(uint64_t)) { + setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), &mem); + } else if (alignment == sizeof(uint32_t)) { + setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), &mem); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); + } else if (alignment == sizeof(uint16_t)) { + setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); setArgument(kernels_[fillType], 1, sizeof(cl_mem), &mem); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); } else { setArgument(kernels_[fillType], 0, sizeof(cl_mem), &mem); - setArgument(kernels_[fillType], 1, sizeof(cl_mem), NULL); + setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); } Memory& gpuCB = gpu().xferWrite().Acquire(patternSize); void* constBuf = gpuCB.map(&gpu(), Resource::NoWait); memcpy(constBuf, pattern, patternSize); gpuCB.unmap(&gpu()); Memory* pGpuCB = &gpuCB; - setArgument(kernels_[fillType], 2, sizeof(cl_mem), &pGpuCB); + setArgument(kernels_[fillType], 4, sizeof(cl_mem), &pGpuCB); uint64_t offset = origin[0]; - if (dwordAligned) { - patternSize /= sizeof(uint32_t); - offset /= sizeof(uint32_t); - } - setArgument(kernels_[fillType], 3, sizeof(uint32_t), &patternSize); - setArgument(kernels_[fillType], 4, sizeof(offset), &offset); - setArgument(kernels_[fillType], 5, sizeof(fillSize), &fillSize); + + patternSize/= alignment; + offset /= alignment; + + setArgument(kernels_[fillType], 5, sizeof(uint32_t), &patternSize); + setArgument(kernels_[fillType], 6, sizeof(offset), &offset); + setArgument(kernels_[fillType], 7, sizeof(fillSize), &fillSize); // Create ND range object for the kernel's execution amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize); diff --git a/rocclr/device/pal/palblit.hpp b/rocclr/device/pal/palblit.hpp index 2a0bdd2725..b8c1d59b3f 100644 --- a/rocclr/device/pal/palblit.hpp +++ b/rocclr/device/pal/palblit.hpp @@ -219,7 +219,7 @@ class KernelBlitManager : public DmaBlitManager { BlitCopyBufferRectAligned, BlitCopyBuffer, BlitCopyBufferAligned, - FillBuffer, + FillBufferAligned, FillImage, Scheduler, GwsInit, @@ -427,9 +427,11 @@ class KernelBlitManager : public DmaBlitManager { 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" + "__amd_rocclr_copyBufferToImage", "__amd_rocclr_copyBufferRect", + "__amd_rocclr_copyBufferRectAligned", "__amd_rocclr_copyBuffer", + "__amd_rocclr_copyBufferAligned", "__amd_rocclr_fillBufferAligned", + "__amd_rocclr_fillImage", "__amd_rocclr_scheduler", + "__amd_rocclr_gwsInit" }; /*@}*/ // namespace pal diff --git a/rocclr/device/rocm/rocblit.cpp b/rocclr/device/rocm/rocblit.cpp index d3b295e49e..63e4c379f0 100644 --- a/rocclr/device/rocm/rocblit.cpp +++ b/rocclr/device/rocm/rocblit.cpp @@ -1938,21 +1938,40 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, synchronize(); return result; } else { - uint fillType = FillBuffer; + uint fillType = FillBufferAligned; size_t globalWorkOffset[3] = {0, 0, 0}; uint64_t fillSize = size[0] / patternSize; size_t globalWorkSize = amd::alignUp(fillSize, 256); size_t localWorkSize = 256; - bool dwordAligned = ((patternSize % sizeof(uint32_t)) == 0) ? true : false; + uint32_t alignment = (patternSize & 0x7) == 0 ? + sizeof(uint64_t) : + (patternSize & 0x3) == 0 ? + sizeof(uint32_t) : + (patternSize & 0x1) == 0 ? + sizeof(uint16_t) : sizeof(uint8_t); // Program kernels arguments for the fill operation cl_mem mem = as_cl(memory.owner()); - if (dwordAligned) { + if (alignment == sizeof(uint64_t)) { + setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), &mem); + } else if (alignment == sizeof(uint32_t)) { + setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), &mem); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); + } else if (alignment == sizeof(uint16_t)) { setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); setArgument(kernels_[fillType], 1, sizeof(cl_mem), &mem); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); } else { setArgument(kernels_[fillType], 0, sizeof(cl_mem), &mem); setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 2, sizeof(cl_mem), nullptr); + setArgument(kernels_[fillType], 3, sizeof(cl_mem), nullptr); } Memory* gpuCB = dev().getRocMemory(constantBuffer_); if (gpuCB == nullptr) { @@ -1964,15 +1983,15 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, memcpy(constBuf, pattern, patternSize); mem = as_cl(gpuCB->owner()); - setArgument(kernels_[fillType], 2, sizeof(cl_mem), &mem, constBufOffset); + setArgument(kernels_[fillType], 4, sizeof(cl_mem), &mem, constBufOffset); uint64_t offset = origin[0]; - if (dwordAligned) { - patternSize /= sizeof(uint32_t); - offset /= sizeof(uint32_t); - } - setArgument(kernels_[fillType], 3, sizeof(uint32_t), &patternSize); - setArgument(kernels_[fillType], 4, sizeof(offset), &offset); - setArgument(kernels_[fillType], 5, sizeof(fillSize), &fillSize); + + patternSize/= alignment; + offset /= alignment; + + setArgument(kernels_[fillType], 5, sizeof(uint32_t), &patternSize); + setArgument(kernels_[fillType], 6, sizeof(offset), &offset); + setArgument(kernels_[fillType], 7, sizeof(fillSize), &fillSize); // Create ND range object for the kernel's execution amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize); diff --git a/rocclr/device/rocm/rocblit.hpp b/rocclr/device/rocm/rocblit.hpp index 5c14bee2c1..c1a758f6bf 100644 --- a/rocclr/device/rocm/rocblit.hpp +++ b/rocclr/device/rocm/rocblit.hpp @@ -238,7 +238,7 @@ class KernelBlitManager : public DmaBlitManager { BlitCopyBufferRectAligned, BlitCopyBuffer, BlitCopyBufferAligned, - FillBuffer, + FillBufferAligned, FillImage, Scheduler, GwsInit, @@ -462,9 +462,11 @@ class KernelBlitManager : public DmaBlitManager { 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" + "__amd_rocclr_copyBufferToImage", "__amd_rocclr_copyBufferRect", + "__amd_rocclr_copyBufferRectAligned", "__amd_rocclr_copyBuffer", + "__amd_rocclr_copyBufferAligned", "__amd_rocclr_fillBufferAligned", + "__amd_rocclr_fillImage", "__amd_rocclr_scheduler", + "__amd_rocclr_gwsInit" }; inline void KernelBlitManager::setArgument(amd::Kernel* kernel, size_t index,