[SYCL] dpc++ tests

Author: vladimirlaz

SYCL/Basic/CMakeLists.txt

@@ -0,0 +1,16 @@
+set(LLVM_TOOLS_DIR "${LLVM_BINARY_DIR}/bin/")
+
+#get_target_property(SYCL_BINARY_DIR sycl-toolchain BINARY_DIR)
+
+set(SYCL_INCLUDE "${SYCL_INCLUDE_BUILD_DIR}")
+set(SYCL_TOOLS_SRC_DIR "${PROJECT_SOURCE_DIR}/tools/")
+set(LLVM_BUILD_BINARY_DIRS "${LLVM_BINARY_DIR}/bin/")
+set(LLVM_BUILD_LIBRARY_DIRS "${LLVM_BINARY_DIR}/lib/")
+
+set(RT_TEST_ARGS ${RT_TEST_ARGS} "-v")
+set(DEPLOY_RT_TEST_ARGS ${DEPLOY_RT_TEST_ARGS} "-v -D SYCL_TOOLS_DIR=${CMAKE_INSTALL_PREFIX}/bin -D SYCL_LIBS_DIR=${CMAKE_INSTALL_PREFIX}/lib${LLVM_LIBDIR_SUFFIX} -D SYCL_INCLUDE=${SYCL_INCLUDE_DEPLOY_DIR}")
+
+find_package(Threads REQUIRED)
+set(SYCL_THREADS_LIB ${CMAKE_THREAD_LIBS_INIT})
+
+configure_file("${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in" "${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg")

SYCL/Basic/README.md

@@ -0,0 +1,90 @@
+# Overview
+SYCL related test based on SYCL-LIT. These tests support
+execution on all supported devices and SYCL backends.
+
+# Table of contents
+ * [Execution](#execution)
+ * [Main parameters](#main-parameters)
+ * [LIT features which can be used to configure test execution](#lit-features-which-can-be-used-to-configure-test-execution)
+
+# Execution
+```
+git clone <GIT_REPO> # e.g. https://github.com/vladimirlaz/llvm-test-suite
+cd llvm-test-suite
+mkdir build
+cd build
+# configuring test execution (selecting compiler version, target BE and target device)
+cmake -G Ninja -DTEST_SUITE_SUBDIRS=SYCL -DTEST_SUITE_LIT=<PATH_TO_llvm-lit> -DSYCL_BE=<SYCL_BE> -DSYCL_TARGET_DEVICES=<TARGET_DEVICES> -C<CMAKE_CHASHED_CONFIG> ..
+# Building full list of tests in subdir
+ninja check
+# or
+llvm-lit .
+# Get list of available tests
+llvm-lit . --show-tests
+# Run specific test
+llvm-lit <path_to_test>
+```
+
+Notes:
+ - it is assumed that LIT framework, FileCheck and other LIT dependencies are available in the same directory with llvm-lit.
+ - compiler variant as well as compile/link options are defined in cashed cmake configurations:
+   - [dpcpp.cmake](../../cmake/caches/dpcpp.cmake)
+   - [clang_fsycl.cmake](../../cmake/cashes/clang_fsycl.cmake)
+   - [clang_fsycl_cuda.cmake](../../cmake/cashes/clang_fsycl_cuda.cmake)
+ - compiler is taken from environment.
+
+# Main parameters
+It is possible to change tets scope my specifying test directory/file in first
+argument to for thelit-runner.py script.
+
+***SYCL_TARGET_DEVICES*** should point to the directory containing DPCPP compiler
+
+***SYCL_TARGET_DEVICES*** defines comma separated target device types (default value is
+ cpu,gpu,acc,host). Supported target_devices values are:
+ - **cpu**  - CPU device available in OpenCL backend only;
+ - **gpu**  - GPU device available in OpenCL, Level Zero and CUDA backends;
+ - **acc**  - FPGA emulator device available in OpenCL backend only;
+ - **host** - SYCL Host device availabel with all backends.
+
+***SYCL_BE*** defined SYCL backend to be used for testing (default is PI_OPENCL).
+Supported sycl_be values:
+ - PI_OPENCL - for OpenCL backend;
+ - PI_CUDA - for CUDA backend;
+ - PI_LEVEL0 - Level Zero backend.
+
+It is asssumed that all dependencies (OpenCL runtimes,
+CUDA SDK, AOT compilers, etc) are available in the system.
+
+See examples below for configuring tests targetting different devices:
+ - SYCL host:
+```
+cmake -G Ninja  -DTEST_SUITE_COLLECT_CODE_SIZE=OFF  -DTEST_SUITE_COLLECT_COMPILE_TIME=OFF -DTEST_SUITE_SUBDIRS=SYCL  -DTEST_SUITE_LIT=<PATH_TO_llvm-lit> -DSYCL_BE=PI_OPENCL -DSYCL_TARGET_DEVICES="host" -C../cmake/caches/clang_fsycl.cmake  ..
+```
+ - OpenCL GPU
+```
+cmake -G Ninja  -DTEST_SUITE_COLLECT_CODE_SIZE=OFF  -DTEST_SUITE_COLLECT_COMPILE_TIME=OFF -DTEST_SUITE_SUBDIRS=SYCL  -DTEST_SUITE_LIT=<PATH_TO_llvm-lit> -DSYCL_BE=PI_OPENCL -DSYCL_TARGET_DEVICES="gpu" -C../cmake/caches/clang_fsycl.cmake  ..
+```
+ - OpenCL CPU
+```
+cmake -G Ninja  -DTEST_SUITE_COLLECT_CODE_SIZE=OFF  -DTEST_SUITE_COLLECT_COMPILE_TIME=OFF -DTEST_SUITE_SUBDIRS=SYCL  -DTEST_SUITE_LIT=<PATH_TO_llvm-lit> -DSYCL_BE=PI_OPENCL -DSYCL_TARGET_DEVICES="gpu" -C../cmake/caches/clang_fsycl.cmake  ..
+```
+ - OpenCL FPGA emulator
+```
+cmake -G Ninja  -DTEST_SUITE_COLLECT_CODE_SIZE=OFF  -DTEST_SUITE_COLLECT_COMPILE_TIME=OFF -DTEST_SUITE_SUBDIRS=SYCL  -DTEST_SUITE_LIT=<PATH_TO_llvm-lit> -DSYCL_BE=PI_OPENCL -DSYCL_TARGET_DEVICES="gpu" -C../cmake/caches/clang_fsycl.cmake  ..
+```
+ - CUDA GPU
+```
+cmake -G Ninja  -DTEST_SUITE_COLLECT_CODE_SIZE=OFF  -DTEST_SUITE_COLLECT_COMPILE_TIME=OFF -DTEST_SUITE_SUBDIRS=SYCL  -DTEST_SUITE_LIT=<PATH_TO_llvm-lit> -DSYCL_BE=PI_CUDA -DSYCL_TARGET_DEVICES="gpu" -C../cmake/caches/clang_fsycl_cuda.cmake  ..
+```
+ - Level Zero GPU
+```
+cmake -G Ninja  -DTEST_SUITE_COLLECT_CODE_SIZE=OFF  -DTEST_SUITE_COLLECT_COMPILE_TIME=OFF -DTEST_SUITE_SUBDIRS=SYCL  -DTEST_SUITE_LIT=<PATH_TO_llvm-lit> -DSYCL_BE=PI_LEVEL0 -DSYCL_TARGET_DEVICES="gpu" -C../cmake/caches/clang_fsycl.cmake  ..
+```
+
+# LIT features which can be used to configure test execution:
+ - **windows**, **linux** - host OS;
+ - **cpu**, **gpu**, **host**, **acc** - target devices;
+ - **cuda**, **opencl**, **level0** - target backend;
+ - **sycl-ls** - sycl-ls tool is available;
+ - **dump_ir**: is set to true if compiler supports dumiping IR. Can be set by setting DUMP_IR_SUPPORTED in cmake. Default is false.
+

SYCL/Basic/aot/Inputs/aot.cpp

@@ -0,0 +1,76 @@
+//==----- 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;
+}

SYCL/Basic/aot/accelerator.cpp

@@ -0,0 +1,13 @@
+//==----- accelerator.cpp - AOT compilation for fpga devices using aoc  ------==//
+//
+// 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: aoc, accelerator
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga-unknown-unknown-sycldevice %S/Inputs/aot.cpp -o %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+// XFAIL: *

SYCL/Basic/aot/cpu.cpp

@@ -0,0 +1,12 @@
+//==----- 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-unknown-unknown-sycldevice %S/Inputs/aot.cpp -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out

SYCL/Basic/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-unknown-unknown-sycldevice -Xsycl-target-backend=spir64_gen-unknown-unknown-sycldevice "-device skl" %S/Inputs/aot.cpp -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out

SYCL/Basic/aot/spec_const_aot.cpp

@@ -0,0 +1,66 @@
+// REQUIRES: opencl-aot, cpu
+//
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64-unknown-unknown-sycldevice %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+//
+// The test checks that the specialization constant feature works with ahead
+// of time compilation.
+
+#include <CL/sycl.hpp>
+
+#include <iostream>
+#include <vector>
+
+class MyInt32Const;
+
+using namespace sycl;
+
+class Kernel;
+
+int main(int argc, char **argv) {
+  cl::sycl::queue q(default_selector{}, [](exception_list l) {
+    for (auto ep : l) {
+      try {
+        std::rethrow_exception(ep);
+      } catch (cl::sycl::exception &e0) {
+        std::cout << e0.what();
+      } catch (std::exception &e1) {
+        std::cout << e1.what();
+      } catch (...) {
+        std::cout << "*** catch (...)\n";
+      }
+    }
+  });
+
+  std::cout << "Running on " << q.get_device().get_info<info::device::name>() << "\n";
+  cl::sycl::program prog(q.get_context());
+
+  cl::sycl::experimental::spec_constant<int32_t, MyInt32Const> i32 =
+      prog.set_spec_constant<MyInt32Const>(10);
+
+  prog.build_with_kernel_type<Kernel>();
+
+  std::vector<int> vec(1);
+  {
+    cl::sycl::buffer<int, 1> buf(vec.data(), vec.size());
+
+    q.submit([&](cl::sycl::handler &cgh) {
+      auto acc = buf.get_access<cl::sycl::access::mode::write>(cgh);
+      cgh.single_task<Kernel>(
+          prog.get_kernel<Kernel>(),
+          [=]() {
+            acc[0] = i32.get();
+          });
+    });
+  }
+  bool passed = true;
+  int val = vec[0];
+  int gold = 0; // with AOT, spec constant is set to C++ default for the type
+
+  if (val != gold) {
+    std::cout << "*** ERROR: " << val << " != " << gold << "(gold)\n";
+    passed = false;
+  }
+  std::cout << (passed ? "passed\n" : "FAILED\n");
+  return passed ? 0 : 1;
+}

SYCL/Basic/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

SYCL/Basic/bit_cast/bit_cast.cpp

@@ -0,0 +1,84 @@
+// 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
+
+#include <CL/sycl.hpp>
+
+#include <iostream>
+#include <math.h>
+#include <type_traits>
+
+constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write;
+
+template <typename To, typename From>
+class BitCastKernel;
+
+template <typename To, typename From>
+To doBitCast(const From &ValueToConvert) {
+  std::vector<To> Vec(1);
+  {
+    sycl::buffer<To, 1> Buf(Vec.data(), 1);
+    sycl::queue Queue;
+    Queue.submit([&](sycl::handler &cgh) {
+      auto acc = Buf.template get_access<sycl_write>(cgh);
+      cgh.single_task<class BitCastKernel<To, From>>([=]() {
+        // TODO: change to sycl::bit_cast in the future
+        acc[0] = sycl::detail::bit_cast<To>(ValueToConvert);
+      });
+    });
+  }
+  return Vec[0];
+}
+
+template <typename To, typename From>
+int test(const From &Value) {
+  auto ValueConvertedTwoTimes = doBitCast<From>(doBitCast<To>(Value));
+  bool isOriginalValueEqualsToConvertedTwoTimes = false;
+  if (std::is_integral<From>::value) {
+    isOriginalValueEqualsToConvertedTwoTimes = Value == ValueConvertedTwoTimes;
+  } else if ((std::is_floating_point<From>::value) || std::is_same<From, cl::sycl::half>::value) {
+    static const float Epsilon = 0.0000001f;
+    isOriginalValueEqualsToConvertedTwoTimes = fabs(Value - ValueConvertedTwoTimes) < Epsilon;
+  } else {
+    std::cerr << "Type " << typeid(From).name() << " neither integral nor floating point nor cl::sycl::half\n";
+    return 1;
+  }
+  if (!isOriginalValueEqualsToConvertedTwoTimes) {
+    std::cerr << "FAIL: Original value which is " << Value << " != value converted two times which is " << ValueConvertedTwoTimes << "\n";
+    return 1;
+  }
+  std::cout << "PASS\n";
+  return 0;
+}
+
+int main() {
+  int ReturnCode = 0;
+
+  std::cout << "cl::sycl::half to unsigned short ...\n";
+  ReturnCode += test<unsigned short>(cl::sycl::half(1.0f));
+
+  std::cout << "unsigned short to cl::sycl::half ...\n";
+  ReturnCode += test<cl::sycl::half>(static_cast<unsigned short>(16384));
+
+  std::cout << "cl::sycl::half to short ...\n";
+  ReturnCode += test<short>(cl::sycl::half(1.0f));
+
+  std::cout << "short to cl::sycl::half ...\n";
+  ReturnCode += test<cl::sycl::half>(static_cast<short>(16384));
+
+  std::cout << "int to float ...\n";
+  ReturnCode += test<float>(static_cast<int>(2));
+
+  std::cout << "float to int ...\n";
+  ReturnCode += test<int>(static_cast<float>(-2.4f));
+
+  std::cout << "unsigned int to float ...\n";
+  ReturnCode += test<float>(static_cast<unsigned int>(6));
+
+  std::cout << "float to unsigned int ...\n";
+  ReturnCode += test<unsigned int>(static_cast<float>(-2.4f));
+
+  return ReturnCode;
+}