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; }