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

P4 to Git Change 1358998 by gandryey@gera-w8 on 2017/01/05 17:36:02

Browse Source 
SWDEV-110504 - OpenCL on ROCm: MAX_WORK_GROUP_SIZE is 256 but max work items per dimension is 1024
	- Report Orca stack capabilities. OCL runtime currently can't report 1024 due to possible regressions in the apps.

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#32 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsettings.cpp#12 edit
This commit is contained in:
foreman
2017-01-05 17:46:40 -05:00
parent 46b8820454
commit 0673572c79
2 changed files with 12 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocclr/runtime/device/rocm/rocdevice.cpp
+7 -7
View File
@@ -823,15 +823,15 @@ Device::populateOCLDeviceConstants()
info_.maxParameterSize_ = 1024; // [TODO]: CAL stack values: 1024*
// constant
uint32_t max_work_group_size = settings().maxWorkGroupSize_;
/*
uint32_t max_work_group_size = 0;
if (HSA_STATUS_SUCCESS !=
hsa_agent_get_info(
_bkendDevice, HSA_AGENT_INFO_WORKGROUP_MAX_SIZE, &max_work_group_size)) {
return false;
}
*/
assert(max_work_group_size > 0);
max_work_group_size = std::min(max_work_group_size,
static_cast<uint32_t>(settings().maxWorkGroupSize_));
info_.maxWorkGroupSize_ = max_work_group_size;
uint16_t max_workgroup_size[3] = { 0, 0, 0 };
@@ -840,13 +840,13 @@ Device::populateOCLDeviceConstants()
_bkendDevice, HSA_AGENT_INFO_WORKGROUP_MAX_DIM, &max_workgroup_size)) {
return false;
}
assert(max_workgroup_size[0] != 0 && max_workgroup_size[1] != 0 &&
max_workgroup_size[2] != 0);
info_.maxWorkItemSizes_[0] = max_workgroup_size[0];
info_.maxWorkItemSizes_[1] = max_workgroup_size[1];
info_.maxWorkItemSizes_[2] = max_workgroup_size[2];
uint16_t max_work_item_size = static_cast<uint16_t>(max_work_group_size);
info_.maxWorkItemSizes_[0] = std::min(max_workgroup_size[0], max_work_item_size);
info_.maxWorkItemSizes_[1] = std::min(max_workgroup_size[1], max_work_item_size);
info_.maxWorkItemSizes_[2] = std::min(max_workgroup_size[2], max_work_item_size);
info_.nativeVectorWidthChar_ = info_.preferredVectorWidthChar_ = 4;
info_.nativeVectorWidthShort_ = info_.preferredVectorWidthShort_ = 2;