Merge intel/llvm-test-suite in-tree

Author: aelovikov-intel

sycl/test-e2e/.clang-format

@@ -0,0 +1,2 @@
+BasedOnStyle: LLVM
+CommentPragmas: "(RUN|FAIL|REQUIRES|UNSUPPORTED|CHECK[A-Za-z0-9_-]*) *:|expected-"

sycl/test-e2e/AOT/Inputs/aot.cpp

@@ -0,0 +1,73 @@
+//==--- aot.cpp - Simple vector addition (AOT compilation example) --------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===---------------------------------------------------------------------===//
+
+#include <sycl/sycl.hpp>
+
+#include <array>
+#include <iostream>
+
+constexpr sycl::access::mode sycl_read = sycl::access::mode::read;
+constexpr sycl::access::mode sycl_write = sycl::access::mode::write;
+
+template <typename T> class SimpleVadd;
+
+template <typename T, size_t N>
+void simple_vadd(const std::array<T, N> &VA, const std::array<T, N> &VB,
+                 std::array<T, N> &VC) {
+  sycl::queue deviceQueue([](sycl::exception_list ExceptionList) {
+    for (std::exception_ptr ExceptionPtr : ExceptionList) {
+      try {
+        std::rethrow_exception(ExceptionPtr);
+      } catch (sycl::exception &E) {
+        std::cerr << E.what();
+      } catch (...) {
+        std::cerr << "Unknown async exception was caught." << std::endl;
+      }
+    }
+  });
+
+  sycl::range<1> numOfItems{N};
+  sycl::buffer<T, 1> bufferA(VA.data(), numOfItems);
+  sycl::buffer<T, 1> bufferB(VB.data(), numOfItems);
+  sycl::buffer<T, 1> bufferC(VC.data(), numOfItems);
+
+  deviceQueue.submit([&](sycl::handler &cgh) {
+    auto accessorA = bufferA.template get_access<sycl_read>(cgh);
+    auto accessorB = bufferB.template get_access<sycl_read>(cgh);
+    auto accessorC = bufferC.template get_access<sycl_write>(cgh);
+
+    cgh.parallel_for<class SimpleVadd<T>>(numOfItems, [=](sycl::id<1> wiID) {
+      accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
+    });
+  });
+
+  deviceQueue.wait_and_throw();
+}
+
+int main() {
+  const size_t array_size = 4;
+  std::array<int, array_size> A = {{1, 2, 3, 4}}, B = {{1, 2, 3, 4}}, C;
+  std::array<float, array_size> D = {{1.f, 2.f, 3.f, 4.f}},
+                                E = {{1.f, 2.f, 3.f, 4.f}}, F;
+  simple_vadd(A, B, C);
+  simple_vadd(D, E, F);
+  for (unsigned int i = 0; i < array_size; i++) {
+    if (C[i] != A[i] + B[i]) {
+      std::cout << "The results are incorrect (element " << i << " is " << C[i]
+                << "!\n";
+      return 1;
+    }
+    if (F[i] != D[i] + E[i]) {
+      std::cout << "The results are incorrect (element " << i << " is " << F[i]
+                << "!\n";
+      return 1;
+    }
+  }
+  std::cout << "The results are correct!\n";
+  return 0;
+}

sycl/test-e2e/AOT/accelerator.cpp

@@ -0,0 +1,12 @@
+//=-- accelerator.cpp - compilation for fpga emulator dev using opencl-aot --=//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===---------------------------------------------------------------------===//
+
+// REQUIRES: opencl-aot, accelerator
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga %S/Inputs/aot.cpp -o %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out

sycl/test-e2e/AOT/cpu.cpp

@@ -0,0 +1,15 @@
+//==--- cpu.cpp - AOT compilation for cpu devices using opencl-aot --------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===---------------------------------------------------------------------===//
+
+// REQUIRES: opencl-aot, cpu
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 %S/Inputs/aot.cpp -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+
+// Test that opencl-aot can handle multiple build options.
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64  %S/Inputs/aot.cpp -Xsycl-target-backend "--bo=-g" -Xsycl-target-backend "--bo=-cl-opt-disable" -o %t2.out

sycl/test-e2e/AOT/gpu.cpp

@@ -0,0 +1,14 @@
+//==--- gpu.cpp - AOT compilation for gen devices using GEN compiler ------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===---------------------------------------------------------------------===//
+
+// REQUIRES: ocloc, gpu
+// UNSUPPORTED: cuda
+// CUDA is not compatible with SPIR.
+//
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %S/Inputs/aot.cpp -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out

sycl/test-e2e/AOT/multiple-devices.cpp

@@ -0,0 +1,56 @@
+//==-- multiple-devices.cpp - Appropriate AOT-compiled image selection -----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: opencl-aot, ocloc, cpu, gpu, accelerator
+// UNSUPPORTED: cuda
+// CUDA is not compatible with SPIR.
+
+// Produce a fat object for all targets (generic SPIR-V, CPU, GPU, FPGA)
+// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_x86_64,spir64_gen,spir64_fpga %S/Inputs/aot.cpp -c -o %t.o
+
+// CPU, GPU, FPGA
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen,spir64_fpga -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_all_aot.out
+// RUN: %CPU_RUN_PLACEHOLDER %t_all_aot.out
+// RUN: %GPU_RUN_PLACEHOLDER %t_all_aot.out
+// RUN: %ACC_RUN_PLACEHOLDER %t_all_aot.out
+
+// CPU, GPU
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_cpu_gpu.out
+// RUN: %CPU_RUN_PLACEHOLDER %t_cpu_gpu.out
+// RUN: %GPU_RUN_PLACEHOLDER %t_cpu_gpu.out
+
+// CPU, FPGA
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_fpga %t.o -o %t_cpu_fpga.out
+// RUN: %CPU_RUN_PLACEHOLDER %t_cpu_fpga.out
+// RUN: %ACC_RUN_PLACEHOLDER %t_cpu_fpga.out
+
+// GPU, FPGA
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen,spir64_fpga -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_gpu_fpga.out
+// RUN: %GPU_RUN_PLACEHOLDER %t_gpu_fpga.out
+// RUN: %ACC_RUN_PLACEHOLDER %t_gpu_fpga.out
+
+// No AOT-compiled image for CPU
+// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen,spir64_fpga -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_spv_gpu_fpga.out
+// RUN: %CPU_RUN_PLACEHOLDER %t_spv_gpu_fpga.out
+// Check that execution on AOT-compatible devices is unaffected
+// RUN: %GPU_RUN_PLACEHOLDER %t_spv_gpu_fpga.out
+// RUN: %ACC_RUN_PLACEHOLDER %t_spv_gpu_fpga.out
+
+// No AOT-compiled image for GPU
+// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_x86_64,spir64_fpga %t.o -o %t_spv_cpu_fpga.out
+// RUN: %GPU_RUN_PLACEHOLDER %t_spv_cpu_fpga.out
+// Check that execution on AOT-compatible devices is unaffected
+// RUN: %CPU_RUN_PLACEHOLDER %t_spv_cpu_fpga.out
+// RUN: %ACC_RUN_PLACEHOLDER %t_spv_cpu_fpga.out
+
+// No AOT-compiled image for FPGA
+// RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_spv_cpu_gpu.out
+// RUN: %ACC_RUN_PLACEHOLDER %t_spv_cpu_gpu.out
+// Check that execution on AOT-compatible devices is unaffected
+// RUN: %CPU_RUN_PLACEHOLDER %t_spv_cpu_gpu.out
+// RUN: %GPU_RUN_PLACEHOLDER %t_spv_cpu_gpu.out

sycl/test-e2e/AOT/with-llvm-bc.cpp

@@ -0,0 +1,17 @@
+//==----- with-llvm-bc.cpp - SYCL kernel with LLVM IR bitcode as binary ----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: cpu, dump_ir
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64 -c %S/Inputs/aot.cpp -o %t.o
+// RUN: %clangxx -fsycl -fsycl-link-targets=spir64 %t.o -o %t.spv
+// RUN: llvm-spirv -r %t.spv -o %t.bc
+// RUN: %clangxx -fsycl -fsycl-add-targets=spir64:%t.bc %t.o -o %t.out
+//
+// Only CPU supports LLVM IR bitcode as a binary
+// RUN: %CPU_RUN_PLACEHOLDER %t.out

sycl/test-e2e/Assert/Inputs/kernels_in_file2.cpp

@@ -0,0 +1,45 @@
+#include "kernels_in_file2.hpp"
+
+#ifdef DEFINE_NDEBUG_INFILE2
+#define NDEBUG
+#else
+#undef NDEBUG
+#endif
+
+#include <cassert>
+
+using namespace sycl;
+using namespace sycl::access;
+
+int calculus(int X) {
+  assert(X && "this message from calculus");
+  return X * 2;
+}
+
+void check_nil(int value) { assert(value && "this message from file2"); }
+
+static constexpr size_t BUFFER_SIZE = 4;
+
+void enqueueKernel_1_fromFile2(queue *Q) {
+  sycl::range<1> numOfItems{BUFFER_SIZE};
+  sycl::buffer<int, 1> Buf(numOfItems);
+
+  Q->submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class kernel1_from_separate_file>(
+        numOfItems, [=](sycl::id<1> wiID) { check_nil(Acc[wiID]); });
+  });
+}
+
+void enqueueKernel_2_fromFile2(queue *Q) {
+  sycl::range<1> numOfItems{BUFFER_SIZE};
+  sycl::buffer<int, 1> Buf(numOfItems);
+
+  Q->submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class kernel2_from_separate_file>(
+        numOfItems, [=](sycl::id<1> wiID) { check_nil(Acc[wiID]); });
+  });
+}

sycl/test-e2e/Assert/Inputs/kernels_in_file2.hpp

@@ -0,0 +1,7 @@
+#include <sycl/sycl.hpp>
+
+SYCL_EXTERNAL int calculus(int X);
+
+void enqueueKernel_1_fromFile2(sycl::queue *Q);
+
+void enqueueKernel_2_fromFile2(sycl::queue *Q);

sycl/test-e2e/Assert/assert_in_kernels.cpp

@@ -0,0 +1,24 @@
+// REQUIRES: linux
+
+// https://github.com/intel/llvm/issues/7634
+// UNSUPPORTED: hip
+
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.cpu.txt || true
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.cpu.txt
+// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.gpu.txt || true
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.gpu.txt
+// Shouldn't fail on ACC as fallback assert isn't enqueued there
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.acc.txt
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.acc.txt
+//
+// CHECK-NOT:  One shouldn't see this message
+// CHECK:      {{.*}}assert_in_kernels.hpp:25: void kernelFunc2(int *, int): {{.*}} [{{[0,2]}},0,0], {{.*}} [0,0,0]
+// CHECK-SAME: Assertion `Buf[wiID] == 0 && "from assert statement"` failed.
+// CHECK-NOT:  test aborts earlier, one shouldn't see this message
+// CHECK-NOT:  The test ended.
+//
+// CHECK-ACC-NOT: {{.*}}assert_in_kernels.hpp:25: void kernelFunc2(int *, int): {{.*}} [{{[0,2]}},0,0], {{.*}} [0,0,0]
+// CHECK-ACC: The test ended.
+
+#include "assert_in_kernels.hpp"

sycl/test-e2e/Assert/assert_in_kernels.hpp

@@ -0,0 +1,69 @@
+#include <cassert>
+#include <iostream>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::access;
+
+void kernelFunc1(int *Buf, int wiID) {
+  Buf[wiID] = 9;
+  assert(Buf[wiID] != 0 && "One shouldn't see this message");
+}
+
+void assertTest1(queue &Q, buffer<int, 1> &Buf) {
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class Kernel_1>(
+        Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc1(&Acc[0], wiID); });
+  });
+}
+
+void kernelFunc2(int *Buf, int wiID) {
+  if (wiID % 2 != 0)
+    Buf[wiID] = 0;
+  assert(Buf[wiID] == 0 && "from assert statement");
+}
+
+void assertTest2(queue &Q, buffer<int, 1> &Buf) {
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class Kernel_2>(
+        Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc2(&Acc[0], wiID); });
+  });
+}
+
+void kernelFunc3(int *Buf, int wiID) {
+  if (wiID == 0)
+    assert(false && "test aborts earlier, one shouldn't see this message");
+  Buf[wiID] = 9;
+}
+
+void assertTest3(queue &Q, buffer<int, 1> &Buf) {
+  Q.submit([&](handler &CGH) {
+    auto Acc = Buf.template get_access<mode::read_write>(CGH);
+
+    CGH.parallel_for<class Kernel_3>(
+        Buf.get_range(), [=](sycl::id<1> wiID) { kernelFunc3(&Acc[0], wiID); });
+  });
+}
+
+int main(int Argc, const char *Argv[]) {
+  std::array<int, 4> Vec = {1, 2, 3, 4};
+  sycl::range<1> numOfItems{Vec.size()};
+  sycl::buffer<int, 1> Buf(Vec.data(), numOfItems);
+
+  queue Q;
+  assertTest1(Q, Buf);
+  Q.wait();
+
+  assertTest2(Q, Buf);
+  Q.wait();
+
+  assertTest3(Q, Buf);
+  Q.wait();
+
+  std::cout << "The test ended." << std::endl;
+  return 0;
+}

sycl/test-e2e/Assert/assert_in_kernels_ndebug.cpp

@@ -0,0 +1,9 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DNDEBUG %S/assert_in_kernels.cpp -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
+// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
+// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
+//
+// CHECK-NOT: One shouldn't see this message
+// CHECK-NOT: from assert statement
+// CHECK-NOT: test aborts earlier, one shouldn't see this message
+// CHECK: The test ended.

sycl/test-e2e/Assert/assert_in_kernels_win.cpp

@@ -0,0 +1,22 @@
+// REQUIRES: windows
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// Shouldn't fail on ACC as fallback assert isn't enqueued there
+// RUN: %ACC_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --check-prefix=CHECK-ACC --input-file %t.txt
+//
+// CHECK-NOT:  One shouldn't see this message
+// FIXME Windows version prints '(null)' instead of '<unknown func>' once in a
+// while for some insane reason.
+// CHECK:      {{.*}}assert_in_kernels.hpp:25: {{<unknown func>|(null)}}: {{.*}} [{{[0,2]}},0,0], {{.*}} [0,0,0]
+// CHECK-SAME: Assertion `Buf[wiID] == 0 && "from assert statement"` failed.
+// CHECK-NOT:  test aborts earlier, one shouldn't see this message
+// CHECK-NOT:  The test ended.
+//
+// CHECK-ACC-NOT: {{.*}}assert_in_kernels.hpp:25: {{<unknown func>|(null)}}: {{.*}} [{{[0,2]}},0,0], {{.*}} [0,0,0]
+// CHECK-ACC: The test ended.
+
+#include "assert_in_kernels.hpp"