Merge branch 'SYCL-2020' of https://github.com/KhronosGroup/SYCL-CTS …

…into sycl2020/device_global/functional_tests/one_kernel

.github/workflows/cts_ci.yml

@@ -55,7 +55,7 @@ jobs:
           # - sycl-impl: computecpp
           # version: 2.8.0
           - sycl-impl: dpcpp
-            version: 1f3f9b9
+            version: ec97c57
           - sycl-impl: hipsycl
             version: b836149
     steps:
@@ -94,7 +94,7 @@ jobs:
           - sycl-impl: computecpp
             version: 2.8.0
           - sycl-impl: dpcpp
-            version: 1f3f9b9
+            version: ec97c57
           - sycl-impl: hipsycl
             version: b836149
     env:

CMakeLists.txt

@@ -43,6 +43,15 @@ else()
 endif()
 # ------------------
 
+# ------------------
+# Extensive mode for running extension oneAPI compile-time property list tests
+option(SYCL_ENABLE_EXT_ONEAPI_PROPERTY_LIST_TESTS
+       "Enable extension oneAPI compile-time property list tests" OFF)
+if(SYCL_ENABLE_EXT_ONEAPI_PROPERTY_LIST_TESTS)
+    message(STATUS "oneAPI extension compile-time property list tests mode: ON")
+endif()
+# ------------------
+
 # ------------------
 # Extensive mode for running legacy tests
 option(SYCL_CTS_TEST_DEPRECATED_FEATURES

ci/dpcpp.filter

@@ -1,17 +1,6 @@
-accessor
-atomic
 buffer
 exceptions
-handler
-hierarchical
-id
-invoke
-item
-math_builtin_api
 multi_ptr
-nd_item
 opencl_interop
-range
 reduction
-specialization_constants
 usm

cmake/FindDPCPP.cmake

@@ -38,6 +38,9 @@ else()
 endif()
 set(CMAKE_CXX_FLAGS "${DPCPP_FP_FLAG} ${CMAKE_CXX_FLAGS}")
 
+# Disable range rounding feature to reduce # of SYCL kernels.
+set(CMAKE_CXX_FLAGS "-D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ ${CMAKE_CXX_FLAGS}")
+
 option(DPCPP_DISABLE_SYCL2020_DEPRECATION_WARNINGS
   "Disable SYCL 2020 deprecation warnings" ON)
 if(DPCPP_DISABLE_SYCL2020_DEPRECATION_WARNINGS)

tests/accessor/accessor_api_image_common.h

@@ -312,7 +312,7 @@ struct use_normalization_coefficient<sycl::cl_half4> : std::true_type {};
 template <typename T>
 typename image_format_channel<T>::storage_t get_expected_image_value() {
   using storage_t = typename image_format_channel<T>::storage_t;
-  if (use_normalization_coefficient<T>::value) {
+  if constexpr (use_normalization_coefficient<T>::value) {
     return static_cast<storage_t>(0.2f);
   } else {
     return static_cast<storage_t>(17);
@@ -984,10 +984,26 @@ class image_accessor_api_sampled_r {
        *    ftp://ftp.icsi.berkeley.edu/pub/theory/priest-thesis.ps.Z
        *  for details
        */
-      return get_expected_value_linear(idx + 1);
+      return get_expected_value_linear(next_id(idx));
     }
   }
 
+  /**
+   *  @brief Computes the next id for getting an expected upper value for the
+   *  linear filtration. If the target is target::image_array, the last
+   *  component is not a coordinate, it is the index of an image in the array
+   *  to read from, so the last component should not be incremented.
+   */
+  template <int dims>
+  sycl::id<dims> next_id(sycl::id<dims> idx) const {
+    static_assert(dims > 0);
+    ++idx;
+    if constexpr (target == sycl::target::image_array) {
+      --idx[dims-1];
+    }
+    return idx;
+  }
+
   /**
    *  @brief Error for floating point image data
    */
@@ -1104,16 +1120,30 @@ class image_accessor_api_sampled_r {
     const bool useNormalized =
         m_sampler.coordinate_normalization_mode ==
             sycl::coordinate_normalization_mode::normalized;
+    const bool useNearest =
+        m_sampler.filtering_mode == sycl::filtering_mode::nearest;
+
     if (useNormalized) {
       const bool worksForLower =
         check_read<acc_coord_tag::use_normalized_lower>(idx);
       if (worksForLower)
         check_read<acc_coord_tag::use_normalized_upper>(idx);
     } else {
-      const bool worksForInteger =
-        check_read<acc_coord_tag::use_int>(idx);
-      if (worksForInteger)
+      const bool worksForFloat =
         check_read<acc_coord_tag::use_float>(idx);
+      if (worksForFloat && useNearest) {
+        /*
+          SYCL 1.2.1 mentions the following: "SYCL overlays the OpenCL
+          specification and inherits its capabilities and restrictions as well
+          as the additional features it provides on top of OpenCL 1.2."
+          
+          OpenCL spec rules that when int is used as the coordinates, the
+          filter mode must be NEAREST: "The read_imagef calls that take integer
+          coordinates must use a sampler with filter mode set to
+          CLK_FILTER_NEAREST ...; otherwise the values returned are undefined."
+         */
+        check_read<acc_coord_tag::use_int>(idx);
+      }
     }
   }
 };

tests/common/get_cts_string.h

@@ -14,8 +14,76 @@
 namespace sycl_cts {
 namespace get_cts_string {
 
+/**
+ * @brief Enum class for verbosity level of output in the for_bundle_state
+ * function
+ *
+ */
+enum class verbosity { brief = 0, detail };
+
+/** @brief Stringify bool value
+ *  @tparam bool flag Value to convert
+ *  @retval String with interpretation of bool
+ */
+constexpr std::string_view for_bool(const bool flag) {
+  return flag ? "true" : "false";
+}
+
+/** @brief Return string's description depending on the type of bundle_state and
+ * verbosity level
+ *  @tparam State of kernel bundle
+ *  @tparam Level of verbosity
+ *  @retval string's description of bundle state
+ */
+template <sycl::bundle_state State, verbosity Level = verbosity::brief>
+inline std::string for_bundle_state() {
+  std::string result;
+  if constexpr (State == sycl::bundle_state::input) {
+    result += "input";
+  } else if constexpr (State == sycl::bundle_state::object) {
+    result += "object";
+  } else if constexpr (State == sycl::bundle_state::executable) {
+    result += "executable";
+  } else {
+    static_assert(State != State, "incorrect kernel bundle state");
+  }
+
+  if constexpr (Level > verbosity::brief) {
+    result += " kernel bundle state";
+  }
+
+  return result;
+}
+
 /** @brief Return string's description depending on the type of address space
- *  @tparam AddressSpace address space type
+ *  @tparam AddressSpace type of address space that needs to be converted to
+ * string
+ *  @retval String description of given address space
+ */
+template <sycl::access::address_space AddressSpace>
+constexpr std::string_view for_address_space() {
+  if constexpr (AddressSpace == sycl::access::address_space::global_space) {
+    return "global_space";
+  } else if constexpr (AddressSpace ==
+                       sycl::access::address_space::local_space) {
+    return "local_space";
+  } else if constexpr (AddressSpace ==
+                       sycl::access::address_space::private_space) {
+    return "private_space";
+  } else if constexpr (AddressSpace ==
+                       sycl::access::address_space::generic_space) {
+    return "generic_space";
+  } else if constexpr (AddressSpace ==
+                       sycl::access::address_space::constant_space) {
+    return "constant_space(deprecated)";
+  } else {
+    static_assert(AddressSpace != AddressSpace,
+                  "Unknown sycl::access::address_space type");
+  }
+}
+
+/** @brief Return string's description depending on the type of decorated
+ *  @tparam Decorated value that needs to be converted to string
  *  @retval String description of address space
  */
 template <sycl::access::decorated Decorated>
@@ -27,7 +95,74 @@ constexpr std::string_view for_decorated() {
   } else if constexpr (Decorated == sycl::access::decorated::legacy) {
     return "legacy";
   } else {
-    static_assert(Decorated != Decorated, "Unknown decorated type");
+    static_assert(Decorated != Decorated,
+                  "Unknown sycl::access::decorated type");
+  }
+}
+
+/** @brief Return string's description depending on the type of Mode
+ *  @tparam Mode value that needs to be converted to string
+ *  @retval String description of mode
+ */
+template <sycl::access::mode Mode>
+constexpr std::string_view for_mode() {
+  if constexpr (Mode == sycl::access::mode::read) {
+    return "read";
+  } else if constexpr (Mode == sycl::access::mode::read_write) {
+    return "read_write";
+  } else if constexpr (Mode == sycl::access::mode::write) {
+    return "write";
+  } else if constexpr (Mode == sycl::access::mode::atomic) {
+    return "atomic(deprecated)";
+  } else if constexpr (Mode == sycl::access::mode::discard_read_write) {
+    return "discard_read_write(deprecated)";
+  } else if constexpr (Mode == sycl::access::mode::discard_write) {
+    return "discard_write(deprecated)";
+  } else {
+    static_assert(Mode != Mode, "Unknown sycl::access_mode type");
+  }
+}
+
+/** @brief Return string's description depending on the type of Target
+ *  @tparam Target value that needs to be converted to string
+ *  @retval String description of target
+ */
+template <sycl::access::target Target>
+constexpr std::string_view for_target() {
+  if constexpr (Target == sycl::access::target::constant_buffer) {
+    return "constant_buffer";
+  } else if constexpr (Target == sycl::access::target::device) {
+    return "device";
+  } else if constexpr (Target == sycl::access::target::host_buffer) {
+    return "host_buffer";
+  } else if constexpr (Target == sycl::access::target::host_image) {
+    return "host_image";
+  } else if constexpr (Target == sycl::access::target::image) {
+    return "image";
+  } else if constexpr (Target == sycl::access::target::image_array) {
+    return "image_array";
+  } else if constexpr (Target == sycl::access::target::local) {
+    return "local";
+  } else if (Target == sycl::access::target::global_buffer) {
+    return "global_buffer(deprecated)";
+  } else {
+    static_assert(Target != Target, "Unknown sycl::access::target type");
+  }
+}
+
+/** @brief Return string's description depending on the type of isPlaceHolder
+ *  @tparam PlaceHolder value that needs to be converted to string
+ *  @retval String description of placeholder
+ */
+template <sycl::access::placeholder Placeholder>
+constexpr std::string_view for_placeholder() {
+  if constexpr (Placeholder == sycl::access::placeholder::false_t) {
+    return "false_t";
+  } else if constexpr (Placeholder == sycl::access::placeholder::true_t) {
+    return "true_t";
+  } else {
+    static_assert(Placeholder != Placeholder,
+                  "Unknown sycl::access::placeholder type");
   }
 }
 

tests/extension/oneapi_property_list/CMakeLists.txt

@@ -0,0 +1,5 @@
+if(SYCL_ENABLE_EXT_ONEAPI_PROPERTY_LIST_TESTS)
+  file(GLOB test_cases_list *.cpp)
+
+  add_cts_test(${test_cases_list})
+endif()

tests/extension/oneapi_property_list/property_list_prop_eq_op.cpp

@@ -0,0 +1,75 @@
+/*******************************************************************************
+//
+//  SYCL 2020 Extension Conformance Test
+//
+//  Provides test for equality and inequality operators for properties
+//
+*******************************************************************************/
+
+#include "../../common/common.h"
+
+#define TEST_NAME property_list_prop_eq_op
+
+namespace TEST_NAMESPACE {
+
+using namespace sycl_cts;
+
+template <typename T, typename U>
+constexpr void check_equal(util::logger &log, T prop1, U prop2,
+                           bool expected_equal) {
+  if ((prop1 == prop2) != expected_equal) {
+    FAIL(log, "wrong result for equality operator");
+  }
+
+  if ((prop1 != prop2) == expected_equal) {
+    FAIL(log, "wrong result for inequality operator");
+  }
+}
+
+/** test sycl::ext::oneapi::property_value equality operators
+ */
+class TEST_NAME : public util::test_base {
+ public:
+  /** return information about this test
+   */
+  void get_info(test_base::info &out) const override {
+    set_test_info(out, TOSTRING(TEST_NAME), TEST_FILE);
+  }
+
+  /** execute the test
+   */
+  void run(util::logger &log) override {
+#if !defined(SYCL_EXT_ONEAPI_PROPERTY_LIST)
+    log.skip("SYCL_EXT_ONEAPI_PROPERTY_LIST is not defined, test is skipped");
+#elif !defined(SYCL_EXT_ONEAPI_DEVICE_GLOBAL)
+    log.skip("SYCL_EXT_ONEAPI_DEVICE_GLOBAL is not defined, test is skipped");
+#else
+    {
+      using namespace sycl::ext::oneapi;
+      property_list props{device_image_scope_v};
+      constexpr auto prop_value_device_image_scope =
+          props.get_property<device_image_scope>();
+
+      property_list props2{implement_in_csr_v<true>};
+      constexpr auto prop_value_implement_in_csr_true =
+          props2.get_property<implement_in_csr>();
+
+      property_list props3{implement_in_csr_v<false>};
+      constexpr auto prop_value_implement_in_csr_false =
+          props3.get_property<implement_in_csr>();
+
+      check_equal(log, prop_value_device_image_scope,
+                  prop_value_device_image_scope, true);
+      check_equal(log, prop_value_implement_in_csr_true,
+                  prop_value_implement_in_csr_true, true);
+      check_equal(log, prop_value_implement_in_csr_true,
+                  prop_value_implement_in_csr_false, false);
+    }
+#endif
+  }
+};
+
+// register this test with the test_collection.
+util::test_proxy<TEST_NAME> proxy;
+
+} /* namespace TEST_NAMESPACE */