P4 to Git Change 1456735 by wchau@wchau_OCL_boltzmann on 2017/09/08 11:04:08

SWDEV-130808 - Resubmit the CL#1456363 after finding the AMD SDK 2.9.1 Apps test failure is an app issue. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_device.cpp#66 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/khronos/headers/opencl1.2/CL/cl.hpp#6 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/khronos/headers/opencl1.2/CL/cl_ext.h#17 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/khronos/headers/opencl2.0/CL/cl.hpp#8 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/khronos/headers/opencl2.0/CL/cl_ext.h#33 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/khronos/headers/opencl2.1/CL/cl.hpp#5 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/khronos/headers/opencl2.1/CL/cl_ext.h#10 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/khronos/headers/opencl2.2/CL/cl.hpp#4 edit ... //depot/stg/opencl/drivers/opencl/api/opencl/khronos/headers/opencl2.2/CL/cl_ext.h#4 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#290 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.cpp#572 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpusettings.cpp#355 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpusettings.hpp#101 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#58 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palsettings.cpp#32 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palsettings.hpp#13 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#60 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsettings.cpp#23 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsettings.hpp#11 edit ... //depot/stg/opencl/drivers/opencl/runtime/utils/flags.hpp#277 edit [ROCm/clr commit: 703f2f6900]
2017-09-08 11:17:38 -04:00
@@ -283,6 +283,10 @@ struct Info : public amd::EmbeddedObject {
  //  using the data-parallel execution model.
  size_t maxWorkGroupSize_;

+  //! Preferred number of work-items in a work-group executing a kernel
+  //  using the data-parallel execution model.
+  size_t preferredWorkGroupSize_;
+
  //! Number of shader engines in physical GPU
  size_t numberOfShaderEngines;

@@ -487,6 +487,7 @@ void NullDevice::fillDeviceInfo(const CALdeviceattribs& calAttr, const gslMemInf
  info_.maxWorkItemSizes_[0] = info_.maxWorkGroupSize_;
  info_.maxWorkItemSizes_[1] = info_.maxWorkGroupSize_;
  info_.maxWorkItemSizes_[2] = info_.maxWorkGroupSize_;
+  info_.preferredWorkGroupSize_ = settings().preferredWorkGroupSize_;

  if (settings().hwLDSSize_ != 0) {
    info_.localMemType_ = CL_LOCAL;
@@ -57,8 +57,8 @@ Settings::Settings() {
  doublePrecision_ = ::CL_KHR_FP64;

  // Fill workgroup info size
-  // @todo: revisit the 256 limitation on workgroup size
-  maxWorkGroupSize_ = 256;
+  maxWorkGroupSize_ = 1024;
+  preferredWorkGroupSize_ = 256;

  hostMemDirectAccess_ = HostMemDisable;

@@ -431,6 +431,10 @@ void Settings::override() {
    maxWorkGroupSize_ = GPU_MAX_WORKGROUP_SIZE;
  }

+  if (GPU_PREFERRED_WORKGROUP_SIZE != 0) {
+    preferredWorkGroupSize_ = GPU_PREFERRED_WORKGROUP_SIZE;
+  }
+
  // Override blit engine type
  if (GPU_BLIT_ENGINE_TYPE != BlitEngineDefault) {
    blitEngine_ = GPU_BLIT_ENGINE_TYPE;
@@ -78,6 +78,7 @@ class Settings : public device::Settings {
  uint maxRenameSize_;                //!< Maximum size for all renames
  uint hwLDSSize_;                    //!< HW local data store size
  uint maxWorkGroupSize_;             //!< Requested workgroup size for this device
+  uint preferredWorkGroupSize_;       //!< Requested preferred workgroup size for this device
  uint hostMemDirectAccess_;          //!< Enables direct access to the host memory
  amd::LibrarySelector libSelector_;  //!< Select linking libraries for compiler
  uint workloadSplitSize_;            //!< Workload split size
@@ -445,6 +445,7 @@ void NullDevice::fillDeviceInfo(const Pal::DeviceProperties& palProp,
  info_.maxWorkItemSizes_[0] = info_.maxWorkGroupSize_;
  info_.maxWorkItemSizes_[1] = info_.maxWorkGroupSize_;
  info_.maxWorkItemSizes_[2] = info_.maxWorkGroupSize_;
+  info_.preferredWorkGroupSize_ = settings().preferredWorkGroupSize_;

  info_.localMemType_ = CL_LOCAL;
  info_.localMemSize_ = settings().hwLDSSize_;
@@ -57,8 +57,8 @@ Settings::Settings() {
  doublePrecision_ = ::CL_KHR_FP64;

  // Fill workgroup info size
-  // @todo: revisit the 256 limitation on workgroup size
-  maxWorkGroupSize_ = 256;
+  maxWorkGroupSize_ = 1024;
+  preferredWorkGroupSize_ = 256;

  hostMemDirectAccess_ = HostMemDisable;

@@ -431,6 +431,10 @@ void Settings::override() {
    maxWorkGroupSize_ = GPU_MAX_WORKGROUP_SIZE;
  }

+  if (GPU_PREFERRED_WORKGROUP_SIZE != 0) {
+    preferredWorkGroupSize_ = GPU_PREFERRED_WORKGROUP_SIZE;
+  }
+
  // Override blit engine type
  if (GPU_BLIT_ENGINE_TYPE != BlitEngineDefault) {
    blitEngine_ = GPU_BLIT_ENGINE_TYPE;
@@ -74,6 +74,7 @@ class Settings : public device::Settings {
  uint maxRenameSize_;         //!< Maximum size for all renames
  uint hwLDSSize_;             //!< HW local data store size
  uint maxWorkGroupSize_;      //!< Requested workgroup size for this device
+  uint preferredWorkGroupSize_;//!< Requested preferred workgroup size for this device
  uint workloadSplitSize_;     //!< Workload split size
  uint minWorkloadTime_;       //!< Minimal workload time in 0.1 ms
  uint maxWorkloadTime_;       //!< Maximum workload time in 0.1 ms
@@ -916,6 +916,7 @@ bool Device::populateOCLDeviceConstants() {
  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_.preferredWorkGroupSize_ = settings().preferredWorkGroupSize_;

  info_.nativeVectorWidthChar_ = info_.preferredVectorWidthChar_ = 4;
  info_.nativeVectorWidthShort_ = info_.preferredVectorWidthShort_ = 2;
@@ -22,7 +22,9 @@ Settings::Settings() {
  enableLocalMemory_ = HSA_LOCAL_MEMORY_ENABLE;
  enableImageHandle_ = true;

-  maxWorkGroupSize_ = 256;
+  maxWorkGroupSize_ = 1024;
+  preferredWorkGroupSize_ = 256;
+
  maxWorkGroupSize2DX_ = 16;
  maxWorkGroupSize2DY_ = 16;
  maxWorkGroupSize3DX_ = 4;
@@ -132,6 +134,10 @@ void Settings::override() {
    maxWorkGroupSize_ = GPU_MAX_WORKGROUP_SIZE;
  }

+  if (GPU_PREFERRED_WORKGROUP_SIZE != 0) {
+    preferredWorkGroupSize_ = GPU_PREFERRED_WORKGROUP_SIZE;
+  }
+
  if (!flagIsDefault(GPU_MAX_COMMAND_QUEUES)) {
    commandQueues_ = GPU_MAX_COMMAND_QUEUES;
  }
@@ -37,6 +37,9 @@ class Settings : public device::Settings {
  //! Default max workgroup size for 1D
  int maxWorkGroupSize_;

+  //! Preferred workgroup size
+  uint preferredWorkGroupSize_;
+
  //! Default max workgroup sizes for 2D
  int maxWorkGroupSize2DX_;
  int maxWorkGroupSize2DY_;
@@ -24,6 +24,8 @@ debug(bool, CPU_USE_ALIGNMENT_MAP, false,                                     \
        "Use flag to enable alignment mapping for parameters for CPU")        \
 release(int, GPU_MAX_WORKGROUP_SIZE, 0,                                       \
        "Maximum number of workitems in a workgroup for GPU, 0 -use default") \
+release(int, GPU_PREFERRED_WORKGROUP_SIZE, 0,                                 \
+        "Preferred number of workitems in a workgroup for GPU, 0 -use default") \
 release(int, GPU_MAX_WORKGROUP_SIZE_2D_X, 0,                                  \
        "Maximum number of workitems in a 2D workgroup for GPU, x component, 0 -use default") \
 release(int, GPU_MAX_WORKGROUP_SIZE_2D_Y, 0,                                  \