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
This commit is contained in:
@@ -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 {
|
||||
|
||||
@@ -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<uint>(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;
|
||||
}
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -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<uint>(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;
|
||||
}
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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_;
|
||||
|
||||
Reference in New Issue
Block a user