Merge pull request intel#1653 from bb-sycl/xmain-cand

Auto pulldown and update tc files for xmain-cand branch on 20230317

Author: myler

SYCL/Basic/accessor/accessor.cpp

@@ -39,7 +39,9 @@ template <typename T> struct InheritedAccessor : public AccAlias<T> {
   using AccAlias<T>::AccAlias;
 };
 
-template <typename Acc> struct AccWrapper { Acc accessor; };
+template <typename Acc> struct AccWrapper {
+  Acc accessor;
+};
 
 template <typename Acc1, typename Acc2> struct AccsWrapper {
   int a;
@@ -58,7 +60,9 @@ template <typename Acc> struct Wrapper2 {
   AccWrapper<Acc> wrapped;
 };
 
-template <typename Acc> struct Wrapper3 { Wrapper2<Acc> w2; };
+template <typename Acc> struct Wrapper3 {
+  Wrapper2<Acc> w2;
+};
 
 template <typename T> void TestAccSizeFuncs(const std::vector<T> &vec) {
   auto test = [=](auto &Res, const auto &Acc) {
@@ -1051,5 +1055,46 @@ int main() {
       assert(v[i] == ((i * 2 + 1) + i));
   }
 
+  // Assignment operator test for 0-dim buffer accessor
+  {
+    sycl::queue Queue;
+    int Data = 32;
+
+    // Explicit block to prompt copy-back to Data
+    {
+      sycl::buffer<int, 1> DataBuffer(&Data, sycl::range<1>(1));
+
+      Queue.submit([&](sycl::handler &CGH) {
+        sycl::accessor<int, 0> Acc(DataBuffer, CGH);
+        CGH.single_task<class acc_0_dim_assignment>([=]() { Acc = 64; });
+      });
+      Queue.wait();
+    }
+
+    assert(Data == 64);
+  }
+
+  // Assignment operator test for 0-dim local accessor
+  {
+    sycl::queue Queue;
+    int Data = 0;
+
+    // Explicit block to prompt copy-back to Data
+    {
+      sycl::buffer<int, 1> DataBuffer(&Data, sycl::range<1>(1));
+
+      Queue.submit([&](sycl::handler &CGH) {
+        sycl::accessor<int, 0> Acc(DataBuffer, CGH);
+        sycl::local_accessor<int, 0> LocalAcc(CGH);
+        CGH.single_task<class local_acc_0_dim_assignment>([=]() {
+          LocalAcc = 64;
+          Acc = LocalAcc;
+        });
+      });
+    }
+
+    assert(Data == 64);
+  }
+
   std::cout << "Test passed" << std::endl;
 }

SYCL/Basic/gpu_max_wgs_error.cpp

@@ -0,0 +1,40 @@
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -fno-sycl-id-queries-fit-in-int
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+#include <numeric>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+const std::string expected_msg =
+    "Total number of work-items in a work-group cannot exceed";
+
+template <int N>
+void check(range<N> global, range<N> local, bool expect_fail = false) {
+  queue q;
+  try {
+    q.submit([&](handler &cgh) {
+      cgh.parallel_for(nd_range<N>(global, local), [=](nd_item<N> item) {});
+    });
+    assert(!expect_fail && "Exception expected");
+  } catch (nd_range_error e) {
+    if (expect_fail) {
+      std::string msg = e.what();
+      assert(msg.rfind(expected_msg, 0) == 0 && "Unexpected error message");
+    } else {
+      throw e;
+    }
+  }
+}
+
+int main() {
+  queue q;
+  device d = q.get_device();
+  id<2> max_2 = d.get_info<sycl::info::device::max_work_item_sizes<2>>();
+  check(range<2>(max_2[0], max_2[1]), range<2>(max_2[0], max_2[1]), true);
+
+  id<3> max_3 = d.get_info<sycl::info::device::max_work_item_sizes<3>>();
+  check(range<3>(max_3[0], max_3[1], max_3[2]),
+        range<3>(max_3[0], max_3[1], max_3[2]), true);
+}

SYCL/Basic/stream/stream.cpp

@@ -24,19 +24,37 @@ int main() {
     context Context = Queue.get_context();
 
     // Check constructor and getters
-    Queue.submit([&](handler &CGH) {
-      stream Out(1024, 80, CGH,
-                 property_list{property::buffer::context_bound{Context}});
-      assert(Out.size() == 1024);
-      assert(Out.get_work_item_buffer_size() == 80);
-      assert(Out.has_property<property::buffer::context_bound>());
-      assert(!Out.has_property<property::queue::in_order>());
-      assert(
-          Out.get_property<property::buffer::context_bound>().get_context() ==
-          Context);
-
-      CGH.single_task<class DummyTask1>([=]() {});
-    });
+    size_t sizeInKernel = 0;
+    size_t workItemBufferSizeInKernel = 0;
+    {
+      sycl::buffer<size_t> bufSize(&sizeInKernel, 1);
+      sycl::buffer<size_t> bufWorkItemBufferSize(&workItemBufferSizeInKernel,
+                                                 1);
+
+      // Check constructor and getters
+      Queue.submit([&](handler &CGH) {
+        stream Out(1024, 80, CGH,
+                   property_list{property::buffer::context_bound{Context}});
+        assert(Out.size() == 1024);
+        assert(Out.get_work_item_buffer_size() == 80);
+        assert(Out.has_property<property::buffer::context_bound>());
+        assert(!Out.has_property<property::queue::in_order>());
+        assert(
+            Out.get_property<property::buffer::context_bound>().get_context() ==
+            Context);
+
+        sycl::accessor accSize(bufSize, CGH, sycl::write_only);
+        sycl::accessor accWorkItemBufferSize(bufWorkItemBufferSize, CGH,
+                                             sycl::write_only);
+
+        CGH.single_task<class DummyTask1>([=]() {
+          accSize[0] = Out.size();
+          accWorkItemBufferSize[0] = Out.get_work_item_buffer_size();
+        });
+      });
+    }
+    assert(sizeInKernel == 1024);
+    assert(workItemBufferSizeInKernel == 80);
 
     // Check common reference semantics
     std::hash<stream> Hasher;

SYCL/Complex/sycl_complex_helper.hpp

@@ -34,6 +34,7 @@ template <> const char *get_typename<float>() { return "float"; }
 template <> const char *get_typename<sycl::half>() { return "sycl::half"; }
 
 // Helper to test each complex specilization
+// Overload for cplx_test_cases
 template <template <typename> typename action, typename... argsT>
 bool test_valid_types(sycl::queue &Q, argsT... args) {
   bool test_passes = true;
@@ -56,6 +57,27 @@ bool test_valid_types(sycl::queue &Q, argsT... args) {
   return test_passes;
 }
 
+// Overload for deci_test_cases
+template <template <typename, typename> typename action, typename... argsT>
+bool test_valid_types(sycl::queue &Q, argsT... args) {
+  bool test_passes = true;
+
+  if (Q.get_device().has(sycl::aspect::fp64)) {
+    test_passes &= action<double, bool>{}(Q, args...);
+    test_passes &= action<double, char>{}(Q, args...);
+    test_passes &= action<double, int>{}(Q, args...);
+    test_passes &= action<double, double>{}(Q, args...);
+  }
+
+  { test_passes &= action<float, float>{}(Q, args...); }
+
+  if (Q.get_device().has(sycl::aspect::fp16)) {
+    test_passes &= action<sycl::half, sycl::half>{}(Q, args...);
+  }
+
+  return test_passes;
+}
+
 // Overload for host only tests
 template <template <typename> typename action, typename... argsT>
 bool test_valid_types(argsT... args) {