Restrict rounding-up optimization to GPU devices.

Author: rdeodhar

sycl/include/CL/sycl/handler.hpp

@@ -760,17 +760,20 @@ class __SYCL_EXPORT handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 
+    // The work group size preferred by this device.
     // A reasonable choice for rounding up the range is 32.
     constexpr size_t GoodLocalSizeX = 32;
 
     // Disable the rounding-up optimizations under these conditions:
-    // 1. The env var SYCL_OPT_PFWGS_DISABLE is set
-    // 2. When the string SYCL_OPT_PFWGS_DISABLE is in the kernel name.
-    // 3. The kernel is provided via an interoperability method.
-    // 4. The API "this_item" is used inside the kernel.
-    // 5. The range is already a multiple of the rounding factor.
+    // 1. The device is not a GPU. Only GPUs benefit from rounding.
+    // 2. The env var SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING is set.
+    // 3. The string SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING is in
+    //    the kernel name.
+    // 4. The kernel is provided via an interoperability method.
+    // 5. The API "this_item" is used inside the kernel.
+    // 6. The range is already a multiple of the rounding factor.
     //
-    // Cases 3 and 4 could be supported with extra effort.
+    // Cases 4 and 5 could be supported with extra effort.
     // As an optimization for the common case it is an
     // implementation choice to not support those scenarios.
     // Note that "this_item" is a free function, i.e. not tied to any
@@ -784,6 +787,7 @@ class __SYCL_EXPORT handler {
     std::string KName = typeid(NameT *).name();
     using KI = detail::KernelInfo<KernelName>;
     bool DisableRounding =
+        !is_gpu(MQueue) ||
         (getenv("SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING") != nullptr) ||
         (KName.find("SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING") !=
          std::string::npos) ||
@@ -1949,6 +1953,12 @@ class __SYCL_EXPORT handler {
   /// \param Count is a number of bytes to be prefetched.
   void prefetch(const void *Ptr, size_t Count);
 
+  /// Check if the queue being used is for a GPU device
+  ///
+  /// \param Queue is the queue for this handler.
+  /// \return Whether the device is a GPU.
+  bool is_gpu(shared_ptr_class<sycl::detail::queue_impl> Queue);
+
 private:
   shared_ptr_class<detail::queue_impl> MQueue;
   /// The storage for the arguments passed.

sycl/source/handler.cpp

@@ -343,5 +343,10 @@ void handler::prefetch(const void *Ptr, size_t Count) {
   MLength = Count;
   MCGType = detail::CG::PREFETCH_USM;
 }
+
+bool handler::is_gpu(shared_ptr_class<sycl::detail::queue_impl> Queue) {
+  device Dev = Queue->get_device();
+  return Dev.is_gpu();
+}
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/test/basic_tests/parallel_for_range_roundup.cpp

@@ -1,6 +1,5 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple  %s -o %t.out
 // RUN: env SYCL_PARALLEL_FOR_RANGE_ROUNDING_TRACE=1 %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
-// RUN: env SYCL_PARALLEL_FOR_RANGE_ROUNDING_TRACE=1 %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 
 #include <CL/sycl.hpp>