From 3165cb60e41fcf24345269b8571d0132fd584ca4 Mon Sep 17 00:00:00 2001 From: foreman Date: Thu, 5 Jan 2017 17:46:40 -0500 Subject: [PATCH] P4 to Git Change 1358998 by gandryey@gera-w8 on 2017/01/05 17:36:02 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 [ROCm/clr commit: 0673572c7909a48c642a46cfb30ed22b3902d300] --- .../clr/rocclr/runtime/device/rocm/rocdevice.cpp | 14 +++++++------- .../clr/rocclr/runtime/device/rocm/rocsettings.cpp | 5 +++++ 2 files changed, 12 insertions(+), 7 deletions(-) diff --git a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp b/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp index ca33e9cead..a11471fb4d 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp @@ -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(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(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; diff --git a/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp b/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp index 0d1bdebcd5..d3cb3a0b17 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocsettings.cpp @@ -101,6 +101,11 @@ Settings::create(bool doublePrecision) void Settings::override() { + // Limit reported workgroup size + if (GPU_MAX_WORKGROUP_SIZE != 0) { + maxWorkGroupSize_ = GPU_MAX_WORKGROUP_SIZE; + } + if (!flagIsDefault(GPU_MAX_COMMAND_QUEUES)) { commandQueues_ = GPU_MAX_COMMAND_QUEUES; }