resolve latest merge conflicts

Signed-off-by: Chris Perkins <[email protected]>

Author: cperkinsintel

SYCL/Basic/access_to_subset.cpp

@@ -1,5 +1,4 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
-// RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out

SYCL/Basic/accessor/accessor.cpp

@@ -1,5 +1,4 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
-// RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
@@ -141,28 +140,26 @@ int main() {
   // Device accessor with 2-dimensional subscript operators.
   {
     sycl::queue Queue;
-    if (!Queue.is_host()) {
-      int array[2][3] = {0};
-      {
-        sycl::range<2> Range(2, 3);
-        sycl::buffer<int, 2> buf((int *)array, Range,
-                                 {sycl::property::buffer::use_host_ptr()});
+    int array[2][3] = {0};
+    {
+      sycl::range<2> Range(2, 3);
+      sycl::buffer<int, 2> buf((int *)array, Range,
+                               {sycl::property::buffer::use_host_ptr()});
 
-        Queue.submit([&](sycl::handler &cgh) {
-          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
-          cgh.parallel_for<class dim2_subscr>(Range, [=](sycl::item<2> itemID) {
-            acc[itemID.get_id(0)][itemID.get_id(1)] += itemID.get_linear_id();
-          });
+      Queue.submit([&](sycl::handler &cgh) {
+        auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+        cgh.parallel_for<class dim2_subscr>(Range, [=](sycl::item<2> itemID) {
+          acc[itemID.get_id(0)][itemID.get_id(1)] += itemID.get_linear_id();
         });
-        Queue.wait();
-      }
-      for (int i = 0; i < 2; i++) {
-        for (int j = 0; j < 3; j++) {
-          if (array[i][j] != i * 3 + j) {
-            std::cerr << array[i][j] << " != " << (i * 3 + j) << std::endl;
-            assert(0);
-            return 1;
-          }
+      });
+      Queue.wait();
+    }
+    for (int i = 0; i < 2; i++) {
+      for (int j = 0; j < 3; j++) {
+        if (array[i][j] != i * 3 + j) {
+          std::cerr << array[i][j] << " != " << (i * 3 + j) << std::endl;
+          assert(0);
+          return 1;
         }
       }
     }
@@ -172,52 +169,48 @@ int main() {
   // check compile error
   {
     sycl::queue queue;
-    if (!queue.is_host()) {
-      sycl::range<2> range(1, 1);
-      int Arr[] = {2};
-      {
-        sycl::buffer<int, 1> Buf(Arr, 1);
-        queue.submit([&](sycl::handler &cgh) {
-          auto acc = sycl::accessor<int, 2, sycl::access::mode::atomic,
-                                    sycl::target::local>(range, cgh);
-          cgh.parallel_for<class dim2_subscr_atomic>(
-              sycl::nd_range<2>{range, range}, [=](sycl::nd_item<2>) {
-                sycl::atomic<int, sycl::access::address_space::local_space>
-                    value = acc[0][0];
-              });
-        });
-      }
+    sycl::range<2> range(1, 1);
+    int Arr[] = {2};
+    {
+      sycl::buffer<int, 1> Buf(Arr, 1);
+      queue.submit([&](sycl::handler &cgh) {
+        auto acc = sycl::accessor<int, 2, sycl::access::mode::atomic,
+                                  sycl::target::local>(range, cgh);
+        cgh.parallel_for<class dim2_subscr_atomic>(
+            sycl::nd_range<2>{range, range}, [=](sycl::nd_item<2>) {
+              sycl::atomic<int, sycl::access::address_space::local_space>
+                  value = acc[0][0];
+            });
+      });
     }
   }
 
   // Device accessor with 3-dimensional subscript operators.
   {
     sycl::queue Queue;
-    if (!Queue.is_host()) {
-      int array[2][3][4] = {0};
-      {
-        sycl::range<3> Range(2, 3, 4);
-        sycl::buffer<int, 3> buf((int *)array, Range,
-                                 {sycl::property::buffer::use_host_ptr()});
+    int array[2][3][4] = {0};
+    {
+      sycl::range<3> Range(2, 3, 4);
+      sycl::buffer<int, 3> buf((int *)array, Range,
+                               {sycl::property::buffer::use_host_ptr()});
 
-        Queue.submit([&](sycl::handler &cgh) {
-          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
-          cgh.parallel_for<class dim3_subscr>(Range, [=](sycl::item<3> itemID) {
-            acc[itemID.get_id(0)][itemID.get_id(1)][itemID.get_id(2)] +=
-                itemID.get_linear_id();
-          });
+      Queue.submit([&](sycl::handler &cgh) {
+        auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+        cgh.parallel_for<class dim3_subscr>(Range, [=](sycl::item<3> itemID) {
+          acc[itemID.get_id(0)][itemID.get_id(1)][itemID.get_id(2)] +=
+              itemID.get_linear_id();
         });
-        Queue.wait();
-      }
-      for (int i = 0; i < 2; i++) {
-        for (int j = 0; j < 3; j++) {
-          for (int k = 0; k < 4; k++) {
-            int expected = k + 4 * (j + 3 * i);
-            if (array[i][j][k] != expected) {
-              std::cerr << array[i][j][k] << " != " << expected << std::endl;
-              assert(0);
-              return 1;
-            }
+      });
+      Queue.wait();
+    }
+    for (int i = 0; i < 2; i++) {
+      for (int j = 0; j < 3; j++) {
+        for (int k = 0; k < 4; k++) {
+          int expected = k + 4 * (j + 3 * i);
+          if (array[i][j][k] != expected) {
+            std::cerr << array[i][j][k] << " != " << expected << std::endl;
+            assert(0);
+            return 1;
           }
         }
       }
@@ -295,28 +288,26 @@ int main() {
   // Check that accessor is initialized when accessor is wrapped to some class.
   {
     sycl::queue queue;
-    if (!queue.is_host()) {
-      int array[10] = {0};
-      {
-        sycl::buffer<int, 1> buf((int *)array, sycl::range<1>(10),
-                                 {sycl::property::buffer::use_host_ptr()});
-        queue.submit([&](sycl::handler &cgh) {
-          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
-          auto acc_wrapped = AccWrapper<decltype(acc)>{acc};
-          cgh.parallel_for<class wrapped_access1>(
-              sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) {
-                auto idx = it.get_linear_id();
-                acc_wrapped.accessor[idx] = 333;
-              });
-        });
-        queue.wait();
-      }
-      for (int i = 0; i < 10; i++) {
-        if (array[i] != 333) {
-          std::cerr << array[i] << " != 333" << std::endl;
-          assert(0);
-          return 1;
-        }
+    int array[10] = {0};
+    {
+      sycl::buffer<int, 1> buf((int *)array, sycl::range<1>(10),
+                               {sycl::property::buffer::use_host_ptr()});
+      queue.submit([&](sycl::handler &cgh) {
+        auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+        auto acc_wrapped = AccWrapper<decltype(acc)>{acc};
+        cgh.parallel_for<class wrapped_access1>(
+            sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) {
+              auto idx = it.get_linear_id();
+              acc_wrapped.accessor[idx] = 333;
+            });
+      });
+      queue.wait();
+    }
+    for (int i = 0; i < 10; i++) {
+      if (array[i] != 333) {
+        std::cerr << array[i] << " != 333" << std::endl;
+        assert(0);
+        return 1;
       }
     }
   }
@@ -325,40 +316,38 @@ int main() {
   // initialized in proper way and value is assigned.
   {
     sycl::queue queue;
-    if (!queue.is_host()) {
-      int array1[10] = {0};
-      int array2[10] = {0};
-      {
-        sycl::buffer<int, 1> buf1((int *)array1, sycl::range<1>(10),
-                                  {sycl::property::buffer::use_host_ptr()});
-        sycl::buffer<int, 1> buf2((int *)array2, sycl::range<1>(10),
-                                  {sycl::property::buffer::use_host_ptr()});
-        queue.submit([&](sycl::handler &cgh) {
-          auto acc1 = buf1.get_access<sycl::access::mode::read_write>(cgh);
-          auto acc2 = buf2.get_access<sycl::access::mode::read_write>(cgh);
-          auto acc_wrapped =
-              AccsWrapper<decltype(acc1), decltype(acc2)>{10, acc1, 5, acc2};
-          cgh.parallel_for<class wrapped_access2>(
-              sycl::range<1>(10), [=](sycl::item<1> it) {
-                auto idx = it.get_linear_id();
-                acc_wrapped.accessor1[idx] = 333;
-                acc_wrapped.accessor2[idx] = 777;
-              });
-        });
-        queue.wait();
-      }
+    int array1[10] = {0};
+    int array2[10] = {0};
+    {
+      sycl::buffer<int, 1> buf1((int *)array1, sycl::range<1>(10),
+                                {sycl::property::buffer::use_host_ptr()});
+      sycl::buffer<int, 1> buf2((int *)array2, sycl::range<1>(10),
+                                {sycl::property::buffer::use_host_ptr()});
+      queue.submit([&](sycl::handler &cgh) {
+        auto acc1 = buf1.get_access<sycl::access::mode::read_write>(cgh);
+        auto acc2 = buf2.get_access<sycl::access::mode::read_write>(cgh);
+        auto acc_wrapped =
+            AccsWrapper<decltype(acc1), decltype(acc2)>{10, acc1, 5, acc2};
+        cgh.parallel_for<class wrapped_access2>(
+            sycl::range<1>(10), [=](sycl::item<1> it) {
+              auto idx = it.get_linear_id();
+              acc_wrapped.accessor1[idx] = 333;
+              acc_wrapped.accessor2[idx] = 777;
+            });
+      });
+      queue.wait();
+    }
+    for (int i = 0; i < 10; i++) {
       for (int i = 0; i < 10; i++) {
-        for (int i = 0; i < 10; i++) {
-          if (array1[i] != 333) {
-            std::cerr << array1[i] << " != 333" << std::endl;
-            assert(0);
-            return 1;
-          }
-          if (array2[i] != 777) {
-            std::cerr << array2[i] << " != 777" << std::endl;
-            assert(0);
-            return 1;
-          }
+        if (array1[i] != 333) {
+          std::cerr << array1[i] << " != 333" << std::endl;
+          assert(0);
+          return 1;
+        }
+        if (array2[i] != 777) {
+          std::cerr << array2[i] << " != 777" << std::endl;
+          assert(0);
+          return 1;
         }
       }
     }
@@ -367,31 +356,29 @@ int main() {
   // Several levels of wrappers for accessor.
   {
     sycl::queue queue;
-    if (!queue.is_host()) {
-      int array[10] = {0};
-      {
-        sycl::buffer<int, 1> buf((int *)array, sycl::range<1>(10),
-                                 {sycl::property::buffer::use_host_ptr()});
-        queue.submit([&](sycl::handler &cgh) {
-          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
-          auto acc_wrapped = AccWrapper<decltype(acc)>{acc};
-          Wrapper1 wr1;
-          auto wr2 = Wrapper2<decltype(acc)>{wr1, acc_wrapped};
-          auto wr3 = Wrapper3<decltype(acc)>{wr2};
-          cgh.parallel_for<class wrapped_access3>(
-              sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) {
-                auto idx = it.get_linear_id();
-                wr3.w2.wrapped.accessor[idx] = 333;
-              });
-        });
-        queue.wait();
-      }
-      for (int i = 0; i < 10; i++) {
-        if (array[i] != 333) {
-          std::cerr << array[i] << " != 333" << std::endl;
-          assert(0);
-          return 1;
-        }
+    int array[10] = {0};
+    {
+      sycl::buffer<int, 1> buf((int *)array, sycl::range<1>(10),
+                               {sycl::property::buffer::use_host_ptr()});
+      queue.submit([&](sycl::handler &cgh) {
+        auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+        auto acc_wrapped = AccWrapper<decltype(acc)>{acc};
+        Wrapper1 wr1;
+        auto wr2 = Wrapper2<decltype(acc)>{wr1, acc_wrapped};
+        auto wr3 = Wrapper3<decltype(acc)>{wr2};
+        cgh.parallel_for<class wrapped_access3>(
+            sycl::range<1>(buf.get_count()), [=](sycl::item<1> it) {
+              auto idx = it.get_linear_id();
+              wr3.w2.wrapped.accessor[idx] = 333;
+            });
+      });
+      queue.wait();
+    }
+    for (int i = 0; i < 10; i++) {
+      if (array[i] != 333) {
+        std::cerr << array[i] << " != 333" << std::endl;
+        assert(0);
+        return 1;
       }
     }
   }
@@ -563,31 +550,28 @@ int main() {
     sycl::queue q;
     // host device executes kernels via a different method and there
     // is no good way to throw an exception at this time.
-    if (!q.is_host()) {
-      sycl::range<1> r(4);
-      sycl::buffer<int, 1> b(r);
-      try {
-        sycl::accessor<int, 1, sycl::access::mode::read_write,
-                       sycl::access::target::device,
-                       sycl::access::placeholder::true_t>
-            acc(b);
+    sycl::range<1> r(4);
+    sycl::buffer<int, 1> b(r);
+    try {
+      sycl::accessor<int, 1, sycl::access::mode::read_write,
+                     sycl::access::target::device,
+                     sycl::access::placeholder::true_t>
+          acc(b);
 
-        q.submit([&](sycl::handler &cgh) {
-          // we do NOT call .require(acc) without which we should throw a
-          // synchronous exception with errc::kernel_argument
-          cgh.parallel_for<class ph>(
-              r, [=](sycl::id<1> index) { acc[index] = 0; });
-        });
-        q.wait_and_throw();
-        assert(false && "we should not be here, missing exception");
-      } catch (sycl::exception &e) {
-        std::cout << "exception received: " << e.what() << std::endl;
-        assert(e.code() == sycl::errc::kernel_argument &&
-               "incorrect error code");
-      } catch (...) {
-        std::cout << "some other exception" << std::endl;
-        return 1;
-      }
+      q.submit([&](sycl::handler &cgh) {
+        // we do NOT call .require(acc) without which we should throw a
+        // synchronous exception with errc::kernel_argument
+        cgh.parallel_for<class ph>(r,
+                                   [=](sycl::id<1> index) { acc[index] = 0; });
+      });
+      q.wait_and_throw();
+      assert(false && "we should not be here, missing exception");
+    } catch (sycl::exception &e) {
+      std::cout << "exception received: " << e.what() << std::endl;
+      assert(e.code() == sycl::errc::kernel_argument && "incorrect error code");
+    } catch (...) {
+      std::cout << "some other exception" << std::endl;
+      return 1;
     }
   }
 

SYCL/Basic/accessor/device_accessor_deduction.cpp

@@ -1,5 +1,4 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Daccessor_new_api_test %S/Inputs/device_accessor.cpp -o %t.out
-// RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out

SYCL/Basic/accessor/get_device_access_deduction.cpp

@@ -1,5 +1,4 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Dbuffer_new_api_test %S/Inputs/device_accessor.cpp -o %t.out
-// RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out

SYCL/Basic/accessor/get_host_access_deduction.cpp

@@ -1,5 +1,4 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Dbuffer_new_api_test %S/Inputs/host_accessor.cpp -o %t.out
-// RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out

SYCL/Basic/accessor/get_host_task_access_deduction.cpp

@@ -1,5 +1,4 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Dbuffer_new_api_test %S/Inputs/host_task_accessor.cpp -o %t.out
-// RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out