Merge branch 'intel' of https://github.com/intel/llvm-test-suite into esimd_emu_test_updates

Author: dongkyunahn-intel

SYCL/Basic/accessor/accessor.cpp

@@ -657,5 +657,53 @@ int main() {
       return 1;
     }
   }
+
+  // exceptions with illegal ranges or no_init
+  {
+    const size_t bufSize = 10;
+    std::vector<int> res(bufSize);
+    sycl::range<1> r(bufSize);
+    sycl::buffer<int, 1> b(res.data(), r);
+    sycl::range<1> illegalR(bufSize + 1);
+    sycl::id<1> offset(bufSize);
+
+    // illegal ranges
+    try {
+      auto acc = b.get_access<sycl::access::mode::read_write>(illegalR, offset);
+      assert(false && "operation should not have succeeded");
+    } catch (sycl::exception &e) {
+      assert(e.code() == sycl::errc::invalid && "errc should be errc::invalid");
+    }
+    try {
+      sycl::queue q;
+      q.submit([&](sycl::handler &cgh) {
+        auto acc = b.get_access<sycl::access::mode::read_write>(cgh, illegalR);
+      });
+      q.wait_and_throw();
+      assert(false &&
+             "we should not be here. operation should not have succeeded");
+    } catch (sycl::exception &e) {
+      assert(e.code() == sycl::errc::invalid && "errc should be errc::invalid");
+    }
+
+    // no_init incompatible with read_only
+    try {
+      sycl::host_accessor out{b, sycl::read_only, sycl::no_init};
+      assert(false && "operation should have failed");
+    } catch (sycl::exception &e) {
+      assert(e.code() == sycl::errc::invalid && "errc should be errc::invalid");
+    }
+    try {
+      sycl::queue q;
+      q.submit([&](sycl::handler &cgh) {
+        sycl::accessor out{b, cgh, sycl::read_only, sycl::no_init};
+      });
+      q.wait_and_throw();
+      assert(false && "we should not be here. operation should have failed");
+    } catch (sycl::exception &e) {
+      assert(e.code() == sycl::errc::invalid && "errc should be errc::invalid");
+    }
+  }
+
   std::cout << "Test passed" << std::endl;
 }

SYCL/Basic/buffer/buffer.cpp

@@ -394,7 +394,7 @@ int main() {
       myQueue.submit([&](handler &cgh) {
         accessor<int, 2, access::mode::write, access::target::device,
                  access::placeholder::false_t>
-            B(Buffer, cgh, range<2>(20, 20), id<2>(10, 10));
+            B(Buffer, cgh, range<2>(10, 10), id<2>(10, 10));
         cgh.parallel_for<class bufferByRangeOffset>(
             range<2>{10, 5}, [=](id<2> index) { B[index] = 1; });
       });

SYCL/Basic/buffer/buffer_full_copy.cpp

@@ -220,12 +220,39 @@ void check_copy_host_to_device(cl::sycl::queue &Queue) {
   }
 }
 
+void check_exception_code() {
+  sycl::queue q;
+
+  const size_t bufSize = 10;
+  std::vector<int> res(bufSize);
+  // std::iota(res.begin(), res.end(), 1);
+  sycl::range<1> r(bufSize);
+  sycl::buffer<int, 1> b(res.data(), r);
+  sycl::range<1> smallRange(bufSize / 2);
+  sycl::id<1> offset(bufSize);
+
+  try {
+    q.submit([&](sycl::handler &cgh) {
+      sycl::accessor src(b, cgh);
+      sycl::accessor destToSmall(b, cgh, smallRange);
+      cgh.copy(src, destToSmall);
+    });
+    q.wait_and_throw();
+
+    assert(false &&
+           "copy with too small Dest arg should have thrown an exception");
+  } catch (sycl::exception e) {
+    assert(e.code() == sycl::errc::invalid);
+  }
+}
+
 int main() {
   try {
     cl::sycl::queue Queue;
     check_copy_host_to_device(Queue);
     check_copy_device_to_host(Queue);
     check_fill(Queue);
+    check_exception_code();
   } catch (cl::sycl::exception &ex) {
     std::cerr << ex.what() << std::endl;
     return 1;

SYCL/Basic/buffer/subbuffer.cpp

@@ -3,6 +3,9 @@
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// UNSUPPORTED: (opencl && gpu)
+
 //
 //==---------- subbuffer.cpp --- sub-buffer basic test ---------------------==//
 //
@@ -71,26 +74,34 @@ void check1DSubBuffer(cl::sycl::queue &q) {
   size *= 2;
 
   std::size_t offset = size / 2, subbuf_size = 10, offset_inside_subbuf = 3,
-              subbuffer_access_range = 10;
+              subbuffer_access_range = subbuf_size - offset_inside_subbuf; // 7.
   std::vector<int> vec(size);
   std::vector<int> vec2(subbuf_size, 0);
   std::iota(vec.begin(), vec.end(), 0);
 
+  std::cout << "buffer size: " << size << ", subbuffer start: " << offset
+            << std::endl;
+
   try {
     cl::sycl::buffer<int, 1> buf(vec.data(), size);
     cl::sycl::buffer<int, 1> buf2(vec2.data(), subbuf_size);
+    // subbuffer is 10 elements, starting at midpoint. (typically 32)
     cl::sycl::buffer<int, 1> subbuf(buf, cl::sycl::id<1>(offset),
                                     cl::sycl::range<1>(subbuf_size));
 
+    // test offset accessor against a subbuffer
     q.submit([&](cl::sycl::handler &cgh) {
+      // accessor starts at the third element of the subbuffer
+      // and can read for 7 more (ie to the end of the subbuffer)
       auto acc = subbuf.get_access<cl::sycl::access::mode::read_write>(
           cgh, cl::sycl::range<1>(subbuffer_access_range),
           cl::sycl::id<1>(offset_inside_subbuf));
-      cgh.parallel_for<class foobar>(
-          cl::sycl::range<1>(subbuffer_access_range - offset_inside_subbuf),
-          [=](cl::sycl::id<1> i) { acc[i] *= -1; });
+      // subrange is made negative. ( 32 33 34 -35 -36 -37 -38 -39 -40 -41)
+      cgh.parallel_for<class foobar>(cl::sycl::range<1>(subbuffer_access_range),
+                                     [=](cl::sycl::id<1> i) { acc[i] *= -1; });
     });
 
+    // copy results of last operation back to buf2/vec2
     q.submit([&](cl::sycl::handler &cgh) {
       auto acc_sub = subbuf.get_access<cl::sycl::access::mode::read>(cgh);
       auto acc_buf = buf2.get_access<cl::sycl::access::mode::write>(cgh);
@@ -99,27 +110,48 @@ void check1DSubBuffer(cl::sycl::queue &q) {
           [=](cl::sycl::id<1> i) { acc_buf[i] = acc_sub[i]; });
     });
 
+    // multiple entire subbuffer by 10.
+    // now original buffer will be
+    // (..29 30 31 | 320 330 340 -350 -360 -370 -380 -390 -400 -410 | 42 43 44
+    // ...)
     q.submit([&](cl::sycl::handler &cgh) {
       auto acc_sub = subbuf.get_access<cl::sycl::access::mode::read_write>(
-          cgh, cl::sycl::range<1>(subbuffer_access_range));
+          cgh, cl::sycl::range<1>(subbuf_size));
       cgh.parallel_for<class foobar_1>(
-          cl::sycl::range<1>(subbuffer_access_range),
+          cl::sycl::range<1>(subbuf_size),
           [=](cl::sycl::id<1> i) { acc_sub[i] *= 10; });
     });
     q.wait_and_throw();
 
+    // buffers go out of scope. data must be copied back to vector no later than
+    // this.
   } catch (const cl::sycl::exception &e) {
     std::cerr << e.what() << std::endl;
     assert(false && "Exception was caught");
   }
 
+  // check buffer data in the area of the subbuffer
+  // OCL:GPU confused   =>  320 330 340 -350 -360 -370 -380 39   40   41
+  // every other device =>  320 330 340 -350 -360 -370 -380 -390 -400 -410
   for (int i = offset; i < offset + subbuf_size; ++i)
     assert(vec[i] == (i < offset + offset_inside_subbuf ? i * 10 : i * -10) &&
-           "Invalid result in 1d sub buffer");
+           "Invalid result in buffer overlapped by 1d sub buffer");
+
+  // check buffer data in the area OUTSIDE the subbuffer
+  for (int i = 0; i < size; i++) {
+    if (i < offset)
+      assert(vec[i] == i && "data preceding subbuffer incorrectly altered");
+
+    if (i > offset + subbuf_size)
+      assert(vec[i] == i && "data following subbuffer incorrectly altered");
+  }
 
+  // check the copy of the subbuffer data after the first operation
+  // OCL:GPU        => 32 33 34 -35 -36 -37 -38 0   0   0
+  // everyone else  => 32 33 34 -35 -36 -37 -38 -39 -40 -41
   for (int i = 0; i < subbuf_size; ++i)
     assert(vec2[i] == (i < 3 ? (offset + i) : (offset + i) * -1) &&
-           "Invalid result in 1d sub buffer");
+           "Invalid result in captured 1d sub buffer, vec2");
 }
 
 void checkExceptions() {

SYCL/Basic/host_platform_avail.cpp

@@ -4,6 +4,9 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t1.out
 // RUN: env SYCL_DEVICE_FILTER=acc,host %t1.out
 
+// Temporarily disable on L0 due to fails in CI
+// UNSUPPORTED: level_zero
+
 //==------ host_platform_avail.cpp - Host Platform Availability test -------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

SYCL/Basic/intel-ext-device.cpp

@@ -5,6 +5,8 @@
 // REQUIRES: gpu
 // UNSUPPORTED: cuda
 // UNSUPPORTED: hip
+// Temporarily disable on L0 due to fails in CI
+// UNSUPPORTED: level_zero
 
 //==--------- intel-ext-device.cpp - SYCL device test ------------==//
 //

SYCL/Basic/interop/get_native_ze.cpp

@@ -2,6 +2,9 @@
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.ze.out
 // RUN: %t.ze.out
 
+// Temporarily disable on L0 due to fails in CI
+// UNSUPPORTED: level_zero
+
 #include <level_zero/ze_api.h>
 
 #include <sycl/ext/oneapi/backend/level_zero.hpp>
@@ -32,5 +35,14 @@ int main() {
   ze_result_t Err = zeKernelGetProperties(Handle, &KernelProperties);
   assert(Err == ZE_RESULT_SUCCESS);
 
+  // SYCL2020 4.5.1.2 - verify exception errc
+  try {
+    // this test is L0 only, so we ask for an unavailable backend.
+    auto BE2 = sycl::get_native<sycl::backend::opencl>(Q);
+    assert(false && "we should not be here.");
+  } catch (sycl::exception e) {
+    assert(e.code() == sycl::errc::backend_mismatch && "wrong error code");
+  }
+
   return 0;
 }

SYCL/Basic/partition_supported.cpp

@@ -75,6 +75,24 @@ int main() {
         return -1;
       }
     }
+    auto maxUnits = dev.get_info<sycl::info::device::max_compute_units>();
+    try {
+      std::vector<sycl::device> v = dev.create_sub_devices<
+          sycl::info::partition_property::partition_equally>(maxUnits + 1);
+      std::cerr << "create_sub_devices with more than max_compute_units should "
+                   "have thrown an error"
+                << std::endl;
+      return -1;
+    } catch (sycl::exception &ex) {
+      if (ex.code() != sycl::errc::invalid) {
+        std::cerr << "while an exception was correctly thrown, it has the "
+                     "wrong error code"
+                     "we should have received"
+                  << sycl::errc::invalid << "but instead got"
+                  << ex.code().value() << std::endl;
+        return -1;
+      }
+    }
   } else {
     try {
       auto subDevices = dev.create_sub_devices<
@@ -85,6 +103,12 @@ int main() {
                 << std::endl;
       return -1;
     } catch (const cl::sycl::feature_not_supported &e) {
+      if (e.code() != sycl::errc::feature_not_supported) {
+        std::cerr
+            << "error code should be errc::feature_not_supported instead of "
+            << e.code().value() << std::endl;
+        return -1;
+      }
     } catch (...) {
       std::cerr << "device::create_sub_device(info::partition_affinity_domain) "
                    "should have thrown cl::sycl::feature_not_supported"

SYCL/Basic/query_emulate_subdevice.cpp

@@ -3,4 +3,6 @@
 // RUN: NEOReadDebugKeys=1 SYCL_DEVICE_FILTER="gpu" %t.out
 
 // UNSUPPORTED: gpu-intel-dg1,cuda,hip
+// Temporarily disable on L0 due to fails in CI
+// UNSUPPORTED: level_zero
 #include "query.hpp"

SYCL/Basic/queue/queue.cpp

@@ -124,9 +124,11 @@ int main() {
       context Context(deviceA);
       queue Queue(Context, deviceB);
       assert(Context == Queue.get_context());
-    } catch (std::exception &e) {
-      std::cout << "Exception check passed: " << e.what() << std::endl;
-      GotException = true;
+    } catch (sycl::exception &e) {
+      if (e.code() == sycl::errc::invalid) {
+        std::cout << "Exception check passed: " << e.what() << std::endl;
+        GotException = true;
+      }
     }
     assert(GotException);
   }

SYCL/ESIMD/api/esimd_merge.cpp

@@ -7,8 +7,6 @@
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
-// TODO: esimd_emulator fails due to SEGFAULT error from __esimd_svm_scatter
-// XFAIL: esimd_emulator
 // RUN: %clangxx -fsycl %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
@@ -34,10 +32,11 @@ template <class T> void prn(T *arr, int size, const char *title) {
 }
 
 int main(void) {
+  constexpr unsigned NUM_THREADS = 2;
   constexpr unsigned VL = 16;
   constexpr unsigned FACTOR = 2;
   constexpr unsigned SUB_VL = VL / FACTOR / FACTOR;
-  constexpr unsigned Size = VL * 2;
+  constexpr unsigned Size = VL * NUM_THREADS;
 
   queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
 
@@ -66,7 +65,7 @@ int main(void) {
 
   try {
     auto e = q.submit([&](handler &cgh) {
-      cgh.parallel_for<class Test>(Size, [=](id<1> i) SYCL_ESIMD_KERNEL {
+      cgh.parallel_for<class Test>(NUM_THREADS, [=](id<1> i) SYCL_ESIMD_KERNEL {
         simd<int, VL> va(A + i * VL);
         simd<int, VL> vb(B + i * VL);
         simd_mask<SUB_VL> m(M + i * VL);

SYCL/ESIMD/api/simd_view_copy_move_assign.cpp

@@ -38,9 +38,9 @@ bool test(queue q, std::string str, F funcUnderTest) {
   // The expected result gets the first half of values from B,
   int gold[VL];
   for (int i = 0; i < VL; ++i) {
-    A[i] = -i;
-    B[i] = i;
-    gold[i] = (i < HalfVL) ? B[i] : A[i];
+    A[i] = -i - 1;
+    B[i] = i + 1;
+    gold[i] = ((VL > 1) && (i < HalfVL)) ? B[i] : A[i];
   }
 
   try {
@@ -57,8 +57,8 @@ bool test(queue q, std::string str, F funcUnderTest) {
          simd<T, VL> va;
          simd<T, VL> vb;
          if constexpr (VL == 1) {
-           va[0] = PA[0];
-           vb[0] = PB[0];
+           va[0] = scalar_load<T>(PA, 0);
+           vb[0] = scalar_load<T>(PB, 0);
          } else {
            va.copy_from(PA, offset);
            vb.copy_from(PB, offset);
@@ -69,7 +69,7 @@ bool test(queue q, std::string str, F funcUnderTest) {
          funcUnderTest(va_view, vb_view);
 
          if constexpr (VL == 1) {
-           PA[0] = va[0];
+           scalar_store(PB, 0, (T)va[0]);
          } else {
            va.copy_to(PA, offset);
          }

SYCL/Plugin/sycl-ls-gpu-default-any.cpp

@@ -21,3 +21,5 @@
 // The test crashed on CUDA CI machines with the latest OpenCL GPU RT
 // (21.19.19792).
 // UNSUPPORTED: cuda || hip
+// Temporarily disable on L0 due to fails in CI
+// UNSUPPORTED: level_zero

SYCL/Plugin/sycl-ls-gpu-level-zero.cpp

@@ -13,3 +13,5 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
+// Temporarily disable on L0 due to fails in CI
+// UNSUPPORTED: level_zero

SYCL/Plugin/sycl-ls.cpp

@@ -12,3 +12,5 @@
 // The test crashed on CUDA CI machines with the latest OpenCL GPU RT
 // (21.19.19792).
 // UNSUPPORTED: cuda
+// Temporarily disable on L0 due to fails in CI
+// UNSUPPORTED: level_zero

SYCL/Regression/device_num.cpp

@@ -8,6 +8,8 @@
 // The test is using all available BEs but CUDA machine in CI does not have
 // functional OpenCL RT
 // UNSUPPORTED: cuda || hip
+// Temporarily disable on L0 due to fails in CI
+// UNSUPPORTED: level_zero
 
 #include <CL/sycl.hpp>
 #include <iostream>