From 75e3a7c7919eebedb75c1e01d0f01742db2576fb Mon Sep 17 00:00:00 2001 From: foreman Date: Tue, 2 Jan 2018 11:19:58 -0500 Subject: [PATCH] P4 to Git Change 1498909 by wchau@wchau_OCL_boltzmann on 2018/01/02 11:06:42 SWDEV-137270 - Add findLocalWorkSize to OpenCL runtime ROC device Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#48 edit [ROCm/clr commit: 52c7e598b212679f0ae52272d3c41a3284a1aeb0] --- .../rocclr/runtime/device/rocm/rocvirtual.cpp | 89 +++++++++++++++---- 1 file changed, 72 insertions(+), 17 deletions(-) diff --git a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp index 377c58a0da..58466e091a 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp @@ -1398,8 +1398,11 @@ static inline address addArg(address dst, const void* src, size_t size) { // Over rides the workgroup size fields in the packet with runtime/compiler set sizes void setRuntimeCompilerLocalSize(hsa_kernel_dispatch_packet_t& dispatchPacket, - amd::NDRangeContainer sizes, const size_t* compile_size, + amd::NDRangeContainer sizes, device::Kernel* devKernel, const roc::Device& dev) { + + const size_t* compile_size = devKernel->workGroupInfo()->compileSize_; + // Todo (sramalin) need to check if compile_size is set to 0 if dimension is not valid // else this error check is incorrect if (compile_size[0] || compile_size[1] || compile_size[2]) { @@ -1407,20 +1410,74 @@ void setRuntimeCompilerLocalSize(hsa_kernel_dispatch_packet_t& dispatchPacket, dispatchPacket.workgroup_size_y = sizes.dimensions() > 1 ? compile_size[1] : 1; dispatchPacket.workgroup_size_z = sizes.dimensions() > 2 ? compile_size[2] : 1; } else { - // Runtime must set the group size - dispatchPacket.workgroup_size_x = 1; - dispatchPacket.workgroup_size_y = 1; - dispatchPacket.workgroup_size_z = 1; + size_t thrPerGrp; + bool b1DOverrideSet = !flagIsDefault(GPU_MAX_WORKGROUP_SIZE); + bool b2DOverrideSet = !flagIsDefault(GPU_MAX_WORKGROUP_SIZE_2D_X) || + !flagIsDefault(GPU_MAX_WORKGROUP_SIZE_2D_Y); + bool b3DOverrideSet = !flagIsDefault(GPU_MAX_WORKGROUP_SIZE_3D_X) || + !flagIsDefault(GPU_MAX_WORKGROUP_SIZE_3D_Y) || + !flagIsDefault(GPU_MAX_WORKGROUP_SIZE_3D_Z); - if (sizes.dimensions() == 1) { - dispatchPacket.workgroup_size_x = dev.settings().preferredWorkGroupSize_; - } else if (sizes.dimensions() == 2) { - dispatchPacket.workgroup_size_x = dev.settings().maxWorkGroupSize2DX_; - dispatchPacket.workgroup_size_y = dev.settings().maxWorkGroupSize2DY_; - } else if (sizes.dimensions() == 3) { - dispatchPacket.workgroup_size_x = dev.settings().maxWorkGroupSize3DX_; - dispatchPacket.workgroup_size_y = dev.settings().maxWorkGroupSize3DY_; - dispatchPacket.workgroup_size_z = dev.settings().maxWorkGroupSize3DZ_; + bool overrideSet = ((sizes.dimensions() == 1) && b1DOverrideSet) || + ((sizes.dimensions() == 2) && b2DOverrideSet) || + ((sizes.dimensions() == 3) && b3DOverrideSet); + if (!overrideSet) { + // Find threads per group + thrPerGrp = devKernel->workGroupInfo()->size_; + + size_t tmp = thrPerGrp; + // Split the local workgroup into the most efficient way + for (uint d = 0; d < sizes.dimensions(); ++d) { + size_t div = tmp; + for (; (sizes.global()[d] % div) != 0; div--) + ; + sizes.local()[d] = div; + tmp /= div; + } + + // Check if partial dispatch is enabled and + if (dev.settings().partialDispatch_ && + // we couldn't find optimal workload + (sizes.local().product() % devKernel->workGroupInfo()->wavefrontSize_) != 0) { + size_t maxSize = 0; + size_t maxDim = 0; + for (uint d = 0; d < sizes.dimensions(); ++d) { + if (maxSize < sizes.global()[d]) { + maxSize = sizes.global()[d]; + maxDim = d; + } + } + // Check if a local workgroup has the most optimal size + if (thrPerGrp > maxSize) { + thrPerGrp = maxSize; + } + sizes.local()[maxDim] = thrPerGrp; + for (uint d = 0; d < sizes.dimensions(); ++d) { + if (d != maxDim) { + sizes.local()[d] = 1; + } + } + } + + dispatchPacket.workgroup_size_x = sizes.dimensions() > 0 ? sizes.local()[0] : 1; + dispatchPacket.workgroup_size_y = sizes.dimensions() > 1 ? sizes.local()[1] : 1; + dispatchPacket.workgroup_size_z = sizes.dimensions() > 2 ? sizes.local()[2] : 1; + } else { + // Runtime must set the group size + dispatchPacket.workgroup_size_x = 1; + dispatchPacket.workgroup_size_y = 1; + dispatchPacket.workgroup_size_z = 1; + + if (sizes.dimensions() == 1) { + dispatchPacket.workgroup_size_x = dev.settings().preferredWorkGroupSize_; + } else if (sizes.dimensions() == 2) { + dispatchPacket.workgroup_size_x = dev.settings().maxWorkGroupSize2DX_; + dispatchPacket.workgroup_size_y = dev.settings().maxWorkGroupSize2DY_; + } else if (sizes.dimensions() == 3) { + dispatchPacket.workgroup_size_x = dev.settings().maxWorkGroupSize3DX_; + dispatchPacket.workgroup_size_y = dev.settings().maxWorkGroupSize3DY_; + dispatchPacket.workgroup_size_z = dev.settings().maxWorkGroupSize3DZ_; + } } } } @@ -1716,7 +1773,6 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const dispatchPacket.grid_size_y = sizes.dimensions() > 1 ? newGlobalSize[1] : 1; dispatchPacket.grid_size_z = sizes.dimensions() > 2 ? newGlobalSize[2] : 1; - const size_t* compile_size = devKernel->workGroupInfo()->compileSize_; if (sizes.local().product() != 0) { dispatchPacket.workgroup_size_x = sizes.dimensions() > 0 ? sizes.local()[0] : 1; dispatchPacket.workgroup_size_y = sizes.dimensions() > 1 ? sizes.local()[1] : 1; @@ -1724,8 +1780,7 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const } else { amd::NDRangeContainer tmpSizes(sizes.dimensions(), &newOffset[0], &newGlobalSize[0], &(const_cast(sizes).local()[0])); - - setRuntimeCompilerLocalSize(dispatchPacket, tmpSizes, compile_size, dev()); + setRuntimeCompilerLocalSize(dispatchPacket, tmpSizes, devKernel, dev()); } dispatchPacket.kernarg_address = argBuffer; dispatchPacket.group_segment_size = ldsUsage;