From 8aef16e13c447643c3fe98b1c4b6202a900a4cb9 Mon Sep 17 00:00:00 2001 From: foreman Date: Tue, 12 Sep 2017 17:13:01 -0400 Subject: [PATCH] P4 to Git Change 1458035 by wchau@wchau_OCL_boltzmann on 2017/09/12 16:31:24 SWDEV-130808 - set the local sizes to preferredWorkGroupSize_ when clEnqueueNDRange is not given and the kernel does not have required workgroup sizes. Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpukernel.cpp#320 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuvirtual.cpp#411 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palkernel.cpp#36 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#56 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.cpp#26 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#43 edit --- rocclr/runtime/device/gpu/gpukernel.cpp | 8 ++++---- rocclr/runtime/device/gpu/gpuvirtual.cpp | 2 +- rocclr/runtime/device/pal/palkernel.cpp | 6 +++--- rocclr/runtime/device/pal/palvirtual.cpp | 2 +- rocclr/runtime/device/rocm/rockernel.cpp | 4 ++-- rocclr/runtime/device/rocm/rocvirtual.cpp | 2 +- 6 files changed, 12 insertions(+), 12 deletions(-) diff --git a/rocclr/runtime/device/gpu/gpukernel.cpp b/rocclr/runtime/device/gpu/gpukernel.cpp index 663098a07e..15368f6f98 100644 --- a/rocclr/runtime/device/gpu/gpukernel.cpp +++ b/rocclr/runtime/device/gpu/gpukernel.cpp @@ -1053,7 +1053,7 @@ void Kernel::findLocalWorkSize(size_t workDim, const amd::NDRange& gblWorkSize, (workDim > 1) && ((dev().settings().partialDispatch_) || (((gblWorkSize[0] % 16) == 0) && ((gblWorkSize[1] % 16) == 0)))) { // Use 8x8 workgroup size if kernel has image writes - if ((flags() & ImageWrite) || (thrPerGrp != nullDev().info().maxWorkGroupSize_)) { + if ((flags() & ImageWrite) || (thrPerGrp != nullDev().info().preferredWorkGroupSize_)) { lclWorkSize[0] = 8; lclWorkSize[1] = 8; } else { @@ -2174,7 +2174,7 @@ const MetadataVersion MetadataTypeQualifiers = MetadataVersion(3, 1, 103); bool NullKernel::parseArguments(const std::string& metaData, uint* uavRefCount) { // Initialize workgroup info - workGroupInfo_.size_ = nullDev().info().maxWorkGroupSize_; + workGroupInfo_.size_ = nullDev().info().preferredWorkGroupSize_; MetadataVersion mdVersion; // Find first tag @@ -3208,7 +3208,7 @@ bool HSAILKernel::init(amd::hsa::loader::Symbol* sym, bool finalize) { workGroupInfo_.size_ = workGroupInfo_.compileSize_[0] * workGroupInfo_.compileSize_[1] * workGroupInfo_.compileSize_[2]; } else { - workGroupInfo_.size_ = dev().info().maxWorkGroupSize_; + workGroupInfo_.size_ = dev().info().preferredWorkGroupSize_; } // Pull out printf metadata from the ELF @@ -3345,7 +3345,7 @@ void HSAILKernel::findLocalWorkSize(size_t workDim, const amd::NDRange& gblWorkS (workDim > 1) && ((dev().settings().partialDispatch_) || (((gblWorkSize[0] % 16) == 0) && ((gblWorkSize[1] % 16) == 0)))) { // Use 8x8 workgroup size if kernel has image writes - if (flags_.imageWriteEna_ || (thrPerGrp != dev().info().maxWorkGroupSize_)) { + if (flags_.imageWriteEna_ || (thrPerGrp != dev().info().preferredWorkGroupSize_)) { lclWorkSize[0] = 8; lclWorkSize[1] = 8; } else { diff --git a/rocclr/runtime/device/gpu/gpuvirtual.cpp b/rocclr/runtime/device/gpu/gpuvirtual.cpp index 90c7d8c145..1efbb37e5b 100644 --- a/rocclr/runtime/device/gpu/gpuvirtual.cpp +++ b/rocclr/runtime/device/gpu/gpuvirtual.cpp @@ -142,7 +142,7 @@ void VirtualGPU::DmaFlushMgmt::findSplitSize(const Device& dev, uint64_t threads uint64_t workload = threads * instructions; if (maxDispatchWorkload_ < workload) { dispatchSplitSize_ = static_cast(maxDispatchWorkload_ / instructions); - uint fullLoad = dev.info().maxComputeUnits_ * dev.info().maxWorkGroupSize_; + uint fullLoad = dev.info().maxComputeUnits_ * dev.info().preferredWorkGroupSize_; if ((dispatchSplitSize_ % fullLoad) != 0) { dispatchSplitSize_ = (dispatchSplitSize_ / fullLoad + 1) * fullLoad; } diff --git a/rocclr/runtime/device/pal/palkernel.cpp b/rocclr/runtime/device/pal/palkernel.cpp index d16aaf89b4..526e5199a0 100644 --- a/rocclr/runtime/device/pal/palkernel.cpp +++ b/rocclr/runtime/device/pal/palkernel.cpp @@ -659,7 +659,7 @@ bool HSAILKernel::init(amd::hsa::loader::Symbol* sym, bool finalize) { workGroupInfo_.size_ = workGroupInfo_.compileSize_[0] * workGroupInfo_.compileSize_[1] * workGroupInfo_.compileSize_[2]; } else { - workGroupInfo_.size_ = dev().info().maxWorkGroupSize_; + workGroupInfo_.size_ = dev().info().preferredWorkGroupSize_; } // Pull out printf metadata from the ELF @@ -793,7 +793,7 @@ void HSAILKernel::findLocalWorkSize(size_t workDim, const amd::NDRange& gblWorkS (workDim > 1) && ((dev().settings().partialDispatch_) || (((gblWorkSize[0] % 16) == 0) && ((gblWorkSize[1] % 16) == 0)))) { // Use 8x8 workgroup size if kernel has image writes - if (flags_.imageWriteEna_ || (thrPerGrp != dev().info().maxWorkGroupSize_)) { + if (flags_.imageWriteEna_ || (thrPerGrp != dev().info().preferredWorkGroupSize_)) { lclWorkSize[0] = 8; lclWorkSize[1] = 8; } else { @@ -1540,7 +1540,7 @@ bool LightningKernel::init(amd::hsa::loader::Symbol* symbol) { workGroupInfo_.size_ = workGroupInfo_.compileSize_[0] * workGroupInfo_.compileSize_[1] * workGroupInfo_.compileSize_[2]; } else { - workGroupInfo_.size_ = dev().info().maxWorkGroupSize_; + workGroupInfo_.size_ = dev().info().preferredWorkGroupSize_; } initPrintf(programMD->mPrintf); diff --git a/rocclr/runtime/device/pal/palvirtual.cpp b/rocclr/runtime/device/pal/palvirtual.cpp index 33a65eaf5a..1cf28b1bac 100644 --- a/rocclr/runtime/device/pal/palvirtual.cpp +++ b/rocclr/runtime/device/pal/palvirtual.cpp @@ -505,7 +505,7 @@ void VirtualGPU::DmaFlushMgmt::findSplitSize(const Device& dev, uint64_t threads uint64_t workload = threads * instructions; if (maxDispatchWorkload_ < workload) { dispatchSplitSize_ = static_cast(maxDispatchWorkload_ / instructions); - uint fullLoad = dev.info().maxComputeUnits_ * dev.info().maxWorkGroupSize_; + uint fullLoad = dev.info().maxComputeUnits_ * dev.info().preferredWorkGroupSize_; if ((dispatchSplitSize_ % fullLoad) != 0) { dispatchSplitSize_ = (dispatchSplitSize_ / fullLoad + 1) * fullLoad; } diff --git a/rocclr/runtime/device/rocm/rockernel.cpp b/rocclr/runtime/device/rocm/rockernel.cpp index 62edec2d1d..6097376ebe 100644 --- a/rocclr/runtime/device/rocm/rockernel.cpp +++ b/rocclr/runtime/device/rocm/rockernel.cpp @@ -720,7 +720,7 @@ bool Kernel::init_LC() { workGroupInfo_.size_ = workGroupInfo_.compileSize_[0] * workGroupInfo_.compileSize_[1] * workGroupInfo_.compileSize_[2]; } else { - workGroupInfo_.size_ = program_->dev().info().maxWorkGroupSize_; + workGroupInfo_.size_ = program_->dev().info().preferredWorkGroupSize_; } initPrintf_LC(programMD->mPrintf); @@ -807,7 +807,7 @@ bool Kernel::init() { workGroupInfo_.size_ = workGroupInfo_.compileSize_[0] * workGroupInfo_.compileSize_[1] * workGroupInfo_.compileSize_[2]; } else { - workGroupInfo_.size_ = program_->dev().info().maxWorkGroupSize_; + workGroupInfo_.size_ = program_->dev().info().preferredWorkGroupSize_; } // Pull out printf metadata from the ELF diff --git a/rocclr/runtime/device/rocm/rocvirtual.cpp b/rocclr/runtime/device/rocm/rocvirtual.cpp index 6887ebdeab..8592f7e08b 100644 --- a/rocclr/runtime/device/rocm/rocvirtual.cpp +++ b/rocclr/runtime/device/rocm/rocvirtual.cpp @@ -1413,7 +1413,7 @@ void setRuntimeCompilerLocalSize(hsa_kernel_dispatch_packet_t& dispatchPacket, dispatchPacket.workgroup_size_z = 1; if (sizes.dimensions() == 1) { - dispatchPacket.workgroup_size_x = dev.settings().maxWorkGroupSize_; + 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_;