[SYCL] Add tests from intel/llvm (intel#50)

The end to end tests with OpenCL/CUDA/Level_Zero dependencies are moved out in-source LIT tests.
aelovikov-intel · Nov 9, 2020 · 6897dad · 6897dad
1 parent 14012dc
commit 6897dad
Show file tree

Hide file tree

Showing 3 changed files with 159 additions and 0 deletions.
diff --git a/SYCL/AOT/Inputs/aot.cpp b/SYCL/AOT/Inputs/aot.cpp
@@ -0,0 +1,75 @@
+//==--- 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 <CL/sycl.hpp>
+
+#include <array>
+#include <iostream>
+
+constexpr cl::sycl::access::mode sycl_read = cl::sycl::access::mode::read;
+constexpr cl::sycl::access::mode sycl_write = cl::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) {
+  cl::sycl::queue deviceQueue([](cl::sycl::exception_list ExceptionList) {
+    for (cl::sycl::exception_ptr_class ExceptionPtr : ExceptionList) {
+      try {
+        std::rethrow_exception(ExceptionPtr);
+      } catch (cl::sycl::exception &E) {
+        std::cerr << E.what();
+      } catch (...) {
+        std::cerr << "Unknown async exception was caught." << std::endl;
+      }
+    }
+  });
+
+  cl::sycl::range<1> numOfItems{N};
+  cl::sycl::buffer<T, 1> bufferA(VA.data(), numOfItems);
+  cl::sycl::buffer<T, 1> bufferB(VB.data(), numOfItems);
+  cl::sycl::buffer<T, 1> bufferC(VC.data(), numOfItems);
+
+  deviceQueue.submit([&](cl::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,
+    [=](cl::sycl::id<1> wiID) {
+      accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
+    });
+  });
+
+  deviceQueue.wait_and_throw();
+}
+
+int main() {
+  const size_t array_size = 4;
+  std::array<cl::sycl::cl_int, array_size> A = {{1, 2, 3, 4}},
+                                           B = {{1, 2, 3, 4}}, C;
+  std::array<cl::sycl::cl_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;
+}
diff --git a/SYCL/AOT/multiple-devices.cpp b/SYCL/AOT/multiple-devices.cpp
@@ -0,0 +1,67 @@
+//==-- 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, aoc, cpu, gpu, accelerator
+// UNSUPPORTED: cuda
+// CUDA is not compatible with SPIR.
+
+// 1-command compilation case
+// Targeting CPU, GPU, FPGA
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice,spir64_gen-unknown-unknown-sycldevice,spir64_fpga-unknown-unknown-sycldevice -Xsycl-target-backend=spir64_gen-unknown-unknown-sycldevice "-device *" %S/Inputs/aot.cpp -o %t_all.out
+// RUN: %HOST_RUN_PLACEHOLDER %t_all.out
+// RUN: %CPU_RUN_PLACEHOLDER %t_all.out
+// RUN: %GPU_RUN_PLACEHOLDER %t_all.out
+// RUN: %ACC_RUN_PLACEHOLDER %t_all.out
+
+// Produce object file, spirv, device images to combine these differently
+// at link-time, thus testing various AOT-compiled images configurations
+// RUN: %clangxx -fsycl %S/Inputs/aot.cpp -c -o %t.o
+// RUN: %clangxx -fsycl -fsycl-link-targets=spir64-unknown-unknown-sycldevice %t.o -o %t.spv
+// AOT-compile device binary images
+// RUN: opencl-aot %t.spv -o=%t_cpu.ir --device=cpu
+// RUN: ocloc -file %t.spv -spirv_input -output %t_gen.out -output_no_suffix -device cfl
+// RUN: aoc %t.spv -o %t_fpga.aocx -sycl -dep-files=%t.d
+
+// CPU, GPU
+// RUN: %clangxx -fsycl -fsycl-add-targets=spir64_x86_64:%t_cpu.ir,spir64_gen:%t_gen.out %t.o -o %t_cpu_gpu.out
+// RUN: %HOST_RUN_PLACEHOLDER %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-add-targets=spir64_x86_64:%t_cpu.ir,spir64_fpga:%t_fpga.aocx %t.o -o %t_cpu_fpga.out
+// RUN: %HOST_RUN_PLACEHOLDER %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-add-targets=spir64_gen:%t_gen.out,spir64_fpga:%t_fpga.aocx %t.o -o %t_gpu_fpga.out
+// RUN: %HOST_RUN_PLACEHOLDER %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-add-targets=spir64:%t.spv,spir64_gen:%t_gen.out,spir64_fpga:%t_fpga.aocx %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-add-targets=spir64:%t.spv,spir64_x86_64:%t_cpu.ir,spir64_fpga:%t_fpga.aocx %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-add-targets=spir64:%t.spv,spir64_x86_64:%t_cpu.ir,spir64_gen:%t_gen.out %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
diff --git a/SYCL/AOT/with-llvm-bc.cpp b/SYCL/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-unknown-unknown-sycldevice -c %S/Inputs/aot.cpp -o %t.o
+// RUN: %clangxx -fsycl -fsycl-link-targets=spir64-unknown-unknown-sycldevice %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