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

SWDEV-307184 - Move local size calculation

Browse Source 
With COv5 local size calculation must occur before
runtime programs kernel arguments

Change-Id: I0726c6529bde69b8fcf5360aa83986cf84e04168
This commit is contained in:
German Andryeyev
2022-04-05 11:16:35 -04:00
parent e5a791217e
commit caa6110c29
1 changed files with 2 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocclr/device/rocm/rocvirtual.cpp
+2 -1
View File
@@ -2733,6 +2733,8 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
amd::NDRange local(sizes.local());
address hidden_arguments = const_cast<address>(parameters);
// Calculate local size if it wasn't provided
devKernel->FindLocalWorkSize(sizes.dimensions(), sizes.global(), local);
// Check if runtime has to setup hidden arguments
for (uint32_t i = signature.numParameters(); i < signature.numParametersAll(); ++i) {
@@ -2931,7 +2933,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;
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;