Merge branch 'SYCL-2020' into oneapi_enqueue_functions_tests

Signed-off-by: Michael Aziz <[email protected]>

test_plans/sycl_ext_oneapi_enqueue_functions.asciidoc

@@ -0,0 +1,219 @@
+:sectnums:
+:xrefstyle: short
+
+= Test plan for sycl_ext_oneapi_enqueue_functions
+
+This is a test plan for the API described in
+https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions].
+
+
+== Testing scope
+
+=== Device coverage
+
+All of the tests described below are performed only on the default device that
+is selected on the CTS command line.
+
+=== Feature test macro
+
+All of the tests should use `#ifdef SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS` so they can be skipped
+if feature is not supported.
+
+== Tests
+
+* All following tests run with either a queue or handler.
+* Tests that require a handler should create one as follows:
+```C++
+using syclex = sycl::ext::oneapi::experimental;
+
+syclex::submit(q, [&](sycl::handler& h) {
+    // ...
+}
+```
+
+=== Single Task
+
+Define a simple task kernel to compute a value. For each `single_task` overload, launch this kernel using the free-function and the equivalent member function. Assert that the outputs computed from the two launches are the same. The `single_task` overloads are the following:
+
+```C++
+namespace sycl::ext::oneapi::experimental {
+
+template <typename KernelName, typename KernelType>
+void single_task(sycl::queue q, const KernelType& k);
+
+template <typename KernelName, typename KernelType>
+void single_task(sycl::handler h, const KernelType& k);
+
+template <typename Args...>
+void single_task(sycl::queue q, const sycl::kernel& k, Args&&... args);
+
+template <typename Args...>
+void single_task(sycl::handler h, const sycl::kernel& k, Args&&... args);
+
+}
+```
+
+=== Basic Kernel
+
+Define a basic kernel that computes a set of values. Launch this kernel using each `parallel_for` overload and the equivalent `parallel_for` member function. Assert that the output for both kernel launches are the same. The `parallel_for` overloads are the following:
+
+```C++
+namespace sycl::ext::oneapi::experimental {
+
+template <typename KernelName, int Dimensions,
+          typename KernelType, typename... Reductions>
+void parallel_for(sycl::queue q, sycl::range<Dimensions> r,
+                  const KernelType& k, Reductions&&... reductions);
+
+template <typename KernelName, int Dimensions,
+          typename KernelType, typename... Reductions>
+void parallel_for(sycl::handler h, sycl::range<Dimensions> r,
+                  const KernelType& k, Reductions&&... reductions);
+
+template <typename KernelName, int Dimensions,
+          typename Properties,
+          typename KernelType, typename... Reductions>
+void parallel_for(sycl::queue q,
+                  launch_config<sycl::range<Dimensions>, Properties> c,
+                  const KernelType& k, Reductions&&... reductions);
+
+template <typename KernelName, int Dimensions,
+          typename Properties, typename KernelType, typename... Reductions>
+void parallel_for(sycl::handler h,
+                  launch_config<sycl::range<Dimensions>, Properties> c,
+                  const KernelType& k, Reductions&&... reductions);
+
+template <typename KernelName, int Dimensions, typename... Args>
+void parallel_for(sycl::queue q, sycl::range<Dimensions> r,
+                  const sycl::kernel& k, Args&&... args);
+
+template <typename KernelName, int Dimensions, typename... Args>
+void parallel_for(sycl::handler h, sycl::range<Dimensions> r,
+                  const sycl::kernel& k, Args&&... args);
+
+template <typename KernelName, int Dimensions,
+          typename Properties, typename... Args>
+void parallel_for(sycl::queue q,
+                  launch_config<sycl::range<Dimensions>, Properties> c,
+                  const sycl::kernel& k, Args&& args...);
+
+template <typename KernelName, int Dimensions,
+          typename Properties, typename... Args>
+void parallel_for(sycl::handler h,
+                  launch_config<sycl::range<Dimensions>, Properties> c,
+                  const sycl::kernel& k, Args&& args...);
+
+}
+```
+
+=== ND-range Kernel
+
+Define an ND-range kernel that computes a set of values. Launch this kernel using each `nd_launch` overload and the equivalent `parallel_for` member function. Assert that the output for both kernel launches are the same. The `nd_launch` overloads are the following:
+
+```C++
+namespace sycl::ext::oneapi::experimental {
+
+template <typename KernelName, int Dimensions,
+          typename KernelType, typename... Reductions>
+void nd_launch(sycl::queue q, sycl::nd_range<Dimensions> r,
+               const KernelType& k, Reductions&&... reductions);
+
+template <typename KernelName, int Dimensions,
+          typename KernelType, typename... Reductions>
+void nd_launch(sycl::handler h, sycl::nd_range<Dimensions> r,
+               const KernelType& k, Reductions&&... reductions);
+
+template <typename KernelName, int Dimensions,
+          typename Properties,
+          typename KernelType, typename... Reductions>
+void nd_launch(sycl::queue q,
+               launch_config<sycl::nd_range<Dimensions>, Properties> c,
+               const KernelType& k, Reductions&&... reductions);
+
+template <typename KernelName, int Dimensions,
+          typename Properties,
+          typename KernelType, typename... Reductions>
+void nd_launch(sycl::handler h,
+               launch_config<sycl::nd_range<Dimensions>, Properties> c,
+               const KernelType& k, Reductions&&... reductions);
+
+template <typename KernelName, int Dimensions, typename... Args>
+void nd_launch(sycl::queue q, sycl::nd_range<Dimensions> r,
+               const sycl::kernel& k, Args&&... args);
+
+template <typename KernelName, int Dimensions, typename... Args>
+void nd_launch(sycl::handler h, sycl::nd_range<Dimensions> r,
+               const sycl::kernel& k, Args&&... args);
+
+template <typename KernelName, int Dimensions,
+          typename Properties, typename... Args>
+void nd_launch(sycl::queue q,
+               launch_config<sycl::nd_range<Dimensions>, Properties> c,
+               const sycl::kernel& k, Args&& args...);
+
+template <typename KernelName, int Dimensions,
+          typename Properties, typename... Args>
+void nd_launch(sycl::handler h,
+               launch_config<sycl::nd_range<Dimensions>, Properties> c,
+               const sycl::kernel& k, Args&& args...);
+
+}
+```
+
+=== Memory Operations
+
+For the `memcpy`, `copy`, `memset`, and `fill` memory operations, create one or more test buffers and assert that they have the correct values after the operation completes. For the `prefetch` and `mem_advise` operations, assert that they can be called without throwing an exception. The list of memory operations to test are the following:
+
+```C++
+namespace sycl::ext::oneapi::experimental {
+
+void memcpy(sycl::queue q, void* dest, const void* src, size_t numBytes);
+
+void memcpy(sycl::handler h, void* dest, const void* src, size_t numBytes);
+
+template <typename T>
+void copy(sycl::queue q, const T* src, T* dest, size_t count);
+
+template <typename T>
+void copy(sycl::handler h, const T* src, T* dest, size_t count);
+
+void memset(sycl::queue q, void* ptr, int value, size_t numBytes);
+
+void memset(sycl::handler h, void* ptr, int value, size_t numBytes);
+
+template <typename T>
+void fill(sycl::queue q, T* ptr, const T& pattern, size_t count);
+
+template <typename T>
+void fill(sycl::handler h, T* ptr, const T& pattern, size_t count);
+
+void prefetch(sycl::queue q, void* ptr, size_t numBytes);
+
+void prefetch(sycl::handler h, void* ptr, size_t numBytes);
+
+void mem_advise(sycl::queue q, void* ptr, size_t numBytes, int advice);
+
+void mem_advise(sycl::handler h, void* ptr, size_t numBytes, int advice);
+
+}
+```
+
+=== Command Barriers
+
+These tests should use `#ifdef SYCL_EXT_ONEAPI_ENQUEUE_BARRIER` so they can be skipped
+if feature is not supported. For each barrier function, enqueue a some commands before and after enqueuing the barrier. Assert that the commands enqueued after the barrier do not execute until those enqueued before the barrier have completed. The barrier functions are the following:
+
+```C++
+namespace sycl::ext::oneapi::experimental {
+
+void barrier(sycl::queue q);
+
+void barrier(sycl::handler h);
+
+void partial_barrier(sycl::queue q, const std::vector<sycl::event>& events);
+
+void partial_barrier(sycl::handler h, const std::vector<sycl::event>& events);
+
+}
+```
+

tests/accessor/generic_accessor_api_common.h

@@ -58,9 +58,29 @@ void test_accessor_methods(const AccT &accessor,
 #endif
 }
 
+template <typename Accessor>
+extern const sycl::target accessor_target_v;
+
+template <typename T, int Dims, sycl::access_mode Mode, sycl::target Target,
+          sycl::access::placeholder IsPlaceholder>
+inline constexpr sycl::target
+    accessor_target_v<sycl::accessor<T, Dims, Mode, Target, IsPlaceholder>> =
+        Target;
+
+template <typename T, int Dims>
+inline constexpr sycl::target accessor_target_v<sycl::local_accessor<T, Dims>> =
+    sycl::target::device;
+
+template <typename T, int Dims, sycl::access_mode Mode>
+inline constexpr sycl::target
+    accessor_target_v<sycl::host_accessor<T, Dims, Mode>> =
+        sycl::target::host_buffer;
+
 template <typename T, typename AccT>
 void test_accessor_ptr_host(AccT &accessor, T expected_data) {
-  {
+  // get_multi_ptr is only defined for device accessors and local accessors
+  if constexpr (accessor_target_v<std::remove_cv_t<AccT>> ==
+                sycl::target::device) {
     INFO("check get_multi_ptr() method");
     auto acc_multi_ptr_no =
         accessor.template get_multi_ptr<sycl::access::decorated::no>();

tests/accessor_legacy/accessor_api_buffer_common.h

@@ -108,7 +108,10 @@ class check_buffer_accessor_api_methods {
             "accessor does not properly report placeholder status");
       }
     }
-    {
+    // legacy accessors do not have the size() member function
+    if constexpr (target != sycl::access::target::constant_buffer &&
+                  target != sycl::access::target::local &&
+                  target != sycl::access::target::host_buffer) {
       /** check size() method
        */
       auto accessorCount = accessor.size();

tests/accessor_legacy/accessor_api_local_common.h

@@ -67,19 +67,8 @@ class check_local_accessor_api_methods {
           util::get_cts_object::range<data_dim<dims>::value>::get(1, 1, 1);
       error_buffer_t errorBuffer(errors.get(), sycl::range<1>(2));
 
-      queue.submit([&](sycl::handler &h) {
+      queue.submit([&](sycl::handler& h) {
         auto acc = make_local_accessor_generic<T, dims, mode>(range, h);
-        {
-          /** check size() method
-          */
-          auto accessorCount = acc.size();
-          check_return_type<size_t>(log, accessorCount, "size()");
-          const auto expectedCount = ((dims == 0) ? 1 : count);
-          if (accessorCount != expectedCount) {
-            fail_for_accessor<T, dims, mode, target>(log, typeName,
-                "accessor does not return the correct count");
-          }
-        }
         {
           /** check get_count() method
            */

tests/atomic_ref/atomic_ref_common.h

@@ -253,7 +253,8 @@ inline bool memory_order_is_supported(sycl::queue& q,
   return it != memory_orders_supported.end();
 }
 
-inline bool memory_scope_is_suppoted(sycl::queue& q, sycl::memory_scope scope) {
+inline bool memory_scope_is_supported(sycl::queue& q,
+                                      sycl::memory_scope scope) {
   std::vector<sycl::memory_scope> memory_scopes_supported =
       q.get_device()
           .get_info<sycl::info::device::atomic_memory_scope_capabilities>();
@@ -266,7 +267,7 @@ inline bool memory_order_and_scope_are_supported(sycl::queue& q,
                                                  sycl::memory_order order,
                                                  sycl::memory_scope scope) {
   return memory_order_is_supported(q, order) &&
-         memory_scope_is_suppoted(q, scope);
+         memory_scope_is_supported(q, scope);
 }
 
 inline bool memory_order_and_scope_are_not_supported(sycl::queue& q,

tests/atomic_ref_stress/atomic_ref_stress_common.h

@@ -172,15 +172,15 @@ class aquire_release {
           refA.store(0);
           refB.store(0);
           sycl::group_barrier(item.get_group());
-          if (item.get_local_id() == 0) {
+          if (item.get_local_id() == sycl::id(0)) {
             x = refA.load();
             refB.store(1);
           } else {
             y = refB.load();
             refA.store(1);
           }
           sycl::group_barrier(item.get_group());
-          if (item.get_local_id() == 0)
+          if (item.get_local_id() == sycl::id(0))
             res_acc[item.get_group_linear_id()] = !(x == 1 && y == 1);
         });
       });