AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

P4 to Git Change 1599472 by gandryey@gera-w8 on 2018/08/29 12:25:34

Browse Source 
SWDEV-79445 - OCL generic changes and code clean-up
	- Move FindLocalWorkSize() logic to the abstraction layer
	- Replace the ROCr path with the common FindLocalWorkSize() functionality

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/device.cpp#227 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#314 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/devkernel.cpp#3 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/devkernel.hpp#3 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpukernel.cpp#330 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpukernel.hpp#132 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palkernel.cpp#63 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palkernel.hpp#22 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.cpp#42 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsettings.cpp#36 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsettings.hpp#16 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#67 edit


[ROCm/clr commit: 5ee211e801]
This commit is contained in:
foreman
2018-08-29 12:35:08 -04:00
parent a48fcb65b2
commit 5f5ea9af64
12 changed files with 216 additions and 408 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp
+6 -135
View File
@@ -1764,132 +1764,6 @@ void VirtualGPU::submitMigrateMemObjects(amd::MigrateMemObjectsCommand& vcmd) {
profilingEnd(vcmd);
}
// 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, device::Kernel* devKernel,
const roc::Device& dev) {
Kernel& gpuKernel = static_cast<Kernel&>(*devKernel);
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]) {
dispatchPacket.workgroup_size_x = sizes.dimensions() > 0 ? compile_size[0] : 1;
dispatchPacket.workgroup_size_y = sizes.dimensions() > 1 ? compile_size[1] : 1;
dispatchPacket.workgroup_size_z = sizes.dimensions() > 2 ? compile_size[2] : 1;
} else {
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);
bool overrideSet = ((sizes.dimensions() == 1) && b1DOverrideSet) ||
((sizes.dimensions() == 2) && b2DOverrideSet) ||
((sizes.dimensions() == 3) && b3DOverrideSet);
if (!overrideSet) {
// Find threads per group
thrPerGrp = devKernel->workGroupInfo()->size_;
if (gpuKernel.imageEnable() &&
// and thread group is a multiple value of wavefronts
((thrPerGrp % devKernel->workGroupInfo()->wavefrontSize_) == 0) &&
// and it's 2 or 3-dimensional workload
(sizes.dimensions() > 1) &&
((dev.settings().partialDispatch_) ||
(((sizes.global()[0] % 16) == 0) && ((sizes.global()[1] % 16) == 0)))) {
// Use 8x8 workgroup size if kernel has image writes)
if (gpuKernel.imageWrite() || (thrPerGrp != dev.settings().preferredWorkGroupSize_)) {
sizes.local()[0] = 8;
sizes.local()[1] = 8;
}
else {
sizes.local()[0] = 16;
sizes.local()[1] = 16;
}
if (sizes.dimensions() == 3) {
sizes.local()[2] = 1;
}
}
else {
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;
}
// Assuming DWORD access
const uint cacheLineMatch = dev.info().globalMemCacheLineSize_ >> 2;
// Check if partial dispatch is enabled and
if (dev.settings().partialDispatch_ &&
// we couldn't find optimal workload
((sizes.local().product() % devKernel->workGroupInfo()->wavefrontSize_) != 0 ||
// or size is too small for the cache line
(sizes.local()[0] < cacheLineMatch))) {
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;
}
}
if ((maxDim != 0) && (sizes.global()[0] >= (cacheLineMatch / 2))) {
sizes.local()[0] = cacheLineMatch;
thrPerGrp /= cacheLineMatch;
sizes.local()[maxDim] = thrPerGrp;
for (uint d = 1; d < sizes.dimensions(); ++d) {
if (d != maxDim) {
sizes.local()[d] = 1;
}
}
}
else {
// 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_;
}
}
}
}
bool VirtualGPU::createSchedulerParam()
{
if (nullptr != schedulerParam_) {
@@ -2235,15 +2109,12 @@ 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;
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;
dispatchPacket.workgroup_size_z = sizes.dimensions() > 2 ? sizes.local()[2] : 1;
} else {
amd::NDRangeContainer tmpSizes(sizes.dimensions(), &newOffset[0], &newGlobalSize[0],
&(const_cast<amd::NDRangeContainer&>(sizes).local()[0]));
setRuntimeCompilerLocalSize(dispatchPacket, tmpSizes, devKernel, dev());
}
amd::NDRange local(sizes.local());
devKernel->FindLocalWorkSize(sizes.dimensions(), sizes.global(), local);
dispatchPacket.workgroup_size_x = sizes.dimensions() > 0 ? local[0] : 1;
dispatchPacket.workgroup_size_y = sizes.dimensions() > 1 ? local[1] : 1;
dispatchPacket.workgroup_size_z = sizes.dimensions() > 2 ? local[2] : 1;
dispatchPacket.kernarg_address = argBuffer;
dispatchPacket.group_segment_size = ldsUsage;
dispatchPacket.private_segment_size = devKernel->workGroupInfo()->privateMemSize_;