From a9df586812ff1ba2ccdf2b708a082854036e016e Mon Sep 17 00:00:00 2001 From: "Andryeyev, German" Date: Tue, 15 Apr 2025 15:22:15 -0400 Subject: [PATCH] SWDEV-459758 - Pass workgroup size explicitly (#185) It's easier for compiler to move explicit kernel arguments into user SGPRs [ROCm/clr commit: 3fd7650fe36469680ed5583426cab2d4a22911f7] --- projects/clr/rocclr/device/blitcl.cpp | 15 ++++++++++----- projects/clr/rocclr/device/pal/palblit.cpp | 4 ++++ projects/clr/rocclr/device/rocm/rocblit.cpp | 4 ++++ 3 files changed, 18 insertions(+), 5 deletions(-) diff --git a/projects/clr/rocclr/device/blitcl.cpp b/projects/clr/rocclr/device/blitcl.cpp index 9abfc1e795..969286ebdf 100644 --- a/projects/clr/rocclr/device/blitcl.cpp +++ b/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)) { diff --git a/projects/clr/rocclr/device/pal/palblit.cpp b/projects/clr/rocclr/device/pal/palblit.cpp index bd3d39c80f..090f286ba1 100644 --- a/projects/clr/rocclr/device/pal/palblit.cpp +++ b/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); diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index 4f878e3ad3..e3f471d4fc 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/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);