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

@@ -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*,

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);

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

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);

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

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<amd::Memory>(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<amd::Memory>(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);

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,