SWDEV-459758 - Pass workgroup size explicitly (#185)

It's easier for compiler to move explicit kernel arguments into user SGPRs

[ROCm/clr commit: 3fd7650fe3]

projects/clr/rocclr/device/blitcl.cpp

@@ -51,8 +51,10 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
 
   __kernel void __amd_rocclr_fillBufferAligned(
     __global void* buf, __constant uchar* pattern,
-    uint pattern_size, uint alignment, ulong end_ptr, uint next_chunk) {
-    int id = get_global_id(0);
+    uint pattern_size, uint alignment, ulong end_ptr, uint next_chunk, uint workgroup_size) {
+    uint l = __builtin_amdgcn_workitem_id_x();
+    uint g = __builtin_amdgcn_workgroup_id_x();
+    ulong id = (g * workgroup_size + l);
     long cur_id = id * pattern_size;
     if (alignment == sizeof(ulong2)) {
       __global ulong2* bufULong2 = (__global ulong2*)buf;
@@ -122,9 +124,12 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
   }
 
   __kernel void __amd_rocclr_copyBuffer(__global uchar* src, __global uchar* dst,
-                                          ulong size, uint remainder,
-                                          uint aligned_size, ulong end_ptr, uint next_chunk) {
-    ulong id = get_global_id(0);
+                                        ulong size, uint remainder,
+                                        uint aligned_size, ulong end_ptr,
+                                        uint next_chunk, uint workgroup_size) {
+    uint l = __builtin_amdgcn_workitem_id_x();
+    uint g = __builtin_amdgcn_workgroup_id_x();
+    ulong id = (g * workgroup_size + l);
     ulong id_remainder = id;
 
     if (aligned_size == sizeof(ulong2)) {

projects/clr/rocclr/device/pal/palblit.cpp

@@ -2222,6 +2222,8 @@ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern,
       setArgument(kernels_[kFillType], 4, sizeof(end_ptr), &end_ptr);
       uint32_t next_chunk = globalWorkSize * kpattern_size;
       setArgument(kernels_[kFillType], 5, sizeof(uint32_t), &next_chunk);
+      uint32_t lws = localWorkSize;
+      setArgument(kernels_[kFillType], 6, sizeof(lws), &lws);
 
       // Create ND range object for the kernel's execution
       amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize);
@@ -2285,6 +2287,8 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds
 
     uint32_t next_chunk = globalWorkSize;
     setArgument(kernels_[kBlitType], 6, sizeof(next_chunk), &next_chunk);
+    uint32_t lws = localWorkSize;
+    setArgument(kernels_[kBlitType], 7, sizeof(lws), &lws);
 
     // Create ND range object for the kernel's execution
     amd::NDRangeContainer ndrange(1, nullptr, &globalWorkSize, &localWorkSize);

projects/clr/rocclr/device/rocm/rocblit.cpp

@@ -2023,6 +2023,8 @@ bool KernelBlitManager::fillBuffer1D(device::Memory& memory, const void* pattern
       setArgument(kernels_[kFillType], 4, sizeof(kfill_size), &kfill_size);
       uint32_t next_chunk = globalWorkSize * kpattern_size;
       setArgument(kernels_[kFillType], 5, sizeof(uint32_t), &next_chunk);
+      uint32_t lws = localWorkSize;
+      setArgument(kernels_[kFillType], 6, sizeof(lws), &lws);
 
       // Create ND range object for the kernel's execution
       amd::NDRangeContainer ndrange(1, globalWorkOffset, &globalWorkSize, &localWorkSize);
@@ -2188,6 +2190,8 @@ bool KernelBlitManager::shaderCopyBuffer(address dst, address src,
 
   uint32_t next_chunk = globalWorkSize;
   setArgument(kernels_[kBlitType], 6, sizeof(next_chunk), &next_chunk);
+  uint32_t lws = localWorkSize;
+  setArgument(kernels_[kBlitType], 7, sizeof(lws), &lws);
 
   // Create ND range object for the kernel's execution
   amd::NDRangeContainer ndrange(1, nullptr, &globalWorkSize, &localWorkSize);