From 6897dad9f8de38ff9fd5f81dbf2632d487bd132a Mon Sep 17 00:00:00 2001 From: vladimirlaz Date: Mon, 9 Nov 2020 14:00:10 +0300 Subject: [PATCH] [SYCL] Add tests from intel/llvm (#50) The end to end tests with OpenCL/CUDA/Level_Zero dependencies are moved out in-source LIT tests. --- SYCL/AOT/Inputs/aot.cpp | 75 +++++++++++++++++++++++++++++++++++ SYCL/AOT/multiple-devices.cpp | 67 +++++++++++++++++++++++++++++++ SYCL/AOT/with-llvm-bc.cpp | 17 ++++++++ 3 files changed, 159 insertions(+) create mode 100644 SYCL/AOT/Inputs/aot.cpp create mode 100644 SYCL/AOT/multiple-devices.cpp create mode 100644 SYCL/AOT/with-llvm-bc.cpp diff --git a/SYCL/AOT/Inputs/aot.cpp b/SYCL/AOT/Inputs/aot.cpp new file mode 100644 index 0000000000000..4c7224f8725df --- /dev/null +++ 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 + +#include +#include + +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 class SimpleVadd; + +template +void simple_vadd(const std::array &VA, const std::array &VB, + std::array &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 bufferA(VA.data(), numOfItems); + cl::sycl::buffer bufferB(VB.data(), numOfItems); + cl::sycl::buffer bufferC(VC.data(), numOfItems); + + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto accessorA = bufferA.template get_access(cgh); + auto accessorB = bufferB.template get_access(cgh); + auto accessorC = bufferC.template get_access(cgh); + + cgh.parallel_for>(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 A = {{1, 2, 3, 4}}, + B = {{1, 2, 3, 4}}, C; + std::array 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 new file mode 100644 index 0000000000000..f7ac3d89699c4 --- /dev/null +++ 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 new file mode 100644 index 0000000000000..79af5d5836823 --- /dev/null +++ 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