[SYCL][ESIMD] Add ESIMD tests on TPM

.github/workflows/clang-format.yml

@@ -0,0 +1,34 @@
+name: clang-format-check
+
+on:
+  pull_request:
+    branches:
+    - intel
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+    steps:
+    - name: Get clang-format first
+      run: sudo apt-get install -yqq clang-format-9
+
+    - uses: actions/checkout@v2
+      with:
+        fetch-depth: 2
+
+    - name: Get clang-format-diff.py tool
+      run: |
+        wget https://raw.githubusercontent.com/intel/llvm/sycl/clang/tools/clang-format/clang-format-diff.py
+        chmod u+x clang-format-diff.py
+
+    - name: Run clang-format for the patch
+      run: |
+        git diff -U0 --no-color ${GITHUB_SHA}^1 ${GITHUB_SHA} -- | python clang-format-diff.py -p1 -binary clang-format-9 > ./clang-format.patch
+    # Add patch with formatting fixes to CI job artifacts
+    - uses: actions/upload-artifact@v1
+      with:
+        name: clang-format-patch
+        path: ./clang-format.patch
+
+    - name: Check if clang-format patch is empty
+      run: bash -c "if [ -s ./clang-format.patch ]; then cat ./clang-format.patch; exit 1; fi"

CONTRIBUTING.md

@@ -0,0 +1,72 @@
+# Contributing
+
+## License
+
+This project is licensed under the terms of the Apache License v2.0 with LLVM
+Exceptions license ([LICENSE.txt](LICENSE.TXT)) to ensure our ability to
+contribute this project to the LLVM test suite project under the same license.
+
+By contributing to this project, you agree to the Apache License v2.0 with LLVM
+Exceptions and copyright terms there in and release your contribution under
+these terms.
+
+## Contribution process
+
+### Development
+
+For any changes not related to DPC++, but rather to LLVM in general, it is
+strongly encouraged that you submit the patch to https://github.com/llvm/llvm-test-suite directly.
+See [LLVM contribution guidelines](https://llvm.org/docs/Contributing.html)
+for more information.
+
+- Create a personal fork of the project on GitHub
+  - For the DPC++ end-to-end test development, use **intel** branch as baseline
+    for your changes.
+- Prepare your patch
+  - follow [LLVM coding standards](https://llvm.org/docs/CodingStandards.html)
+  - [clang-format](https://clang.llvm.org/docs/ClangFormat.html) and
+    [clang-tidy](https://clang.llvm.org/extra/clang-tidy/) tools can be
+    integrated into your workflow to ensure formatting and stylistic
+    compliance of your changes.
+  - use
+
+    ```bash
+    wget https://raw.githubusercontent.com/intel/llvm/sycl/clang/tools/clang-format/git-clang-format
+    python git-clang-format `git merge-base origin/intel HEAD`
+    ```
+
+    to check the format of your current changes against the `origin/intel`
+    branch.
+    - `-f` to also correct unstaged changes
+    - `--diff` to only print the diff without applying
+- Instructions about running DPC++ tests
+  - see [README.md](SYCL/README.md).
+
+### Commit message
+
+- When writing your commit message, please make sure to follow
+  [LLVM developer policies](
+  https://llvm.org/docs/DeveloperPolicy.html#commit-messages) on the subject.
+- For any DPC++-related commit, the `[SYCL]` tag should be present in the
+  commit message title. To a reasonable extent, additional tags can be used
+  to signify the component changed, e.g.: `[LIT]`, `[NFC]`, `[Doc]`.
+
+### Review and acceptance testing
+
+- Create a pull request for your changes following [Creating a pull request
+instructions](https://help.github.com/articles/creating-a-pull-request/).
+- CI will run checks (e.g. clang-format-check) as soon as the PR is created.
+- Once the PR is approved and all checks have passed, the pull request is
+ready for merge.
+
+### Merge
+
+Project maintainers merge pull requests using one of the following options:
+
+- [Rebase and merge] The preferable choice for PRs containing a single commit
+- [Squash and merge] Used when there are multiple commits in the PR
+  - Squashing is done to make sure that the project is buildable on any commit
+- [Create a merge commit] Used for LLVM pull-down PRs to preserve hashes of the
+commits pulled from the LLVM community repository
+
+*Other names and brands may be claimed as the property of others.

README.md

@@ -0,0 +1,27 @@
+# LLVM* test suite repository
+
+Please see the LLVM testing infrastructure guide:
+
+  https://llvm.org/docs/TestSuiteGuide.html
+
+for more information on the contents of this repository.
+
+## Introduction
+
+Intel staging area for LLVM test suite contribution. Home for oneAPI Data
+Parallel C++ compiler tests extending LLVM test suite.
+
+## License
+
+See [LICENSE.txt](LICENSE.TXT) for details.
+
+## Contributing
+
+See [CONTRIBUTING.md](CONTRIBUTING.md) for details.
+
+## Related projects documentation
+
+* oneAPI Data Parallel C++ compiler - See
+  [DPC++ Documentation](https://intel.github.io/llvm-docs/)
+
+\*Other names and brands may be claimed as the property of others.

SYCL/.clang-format

@@ -0,0 +1,2 @@
+BasedOnStyle: LLVM
+CommentPragmas: "(RUN|FAIL|REQUIRES|UNSUPPORTED|CHECK) *:|expected-"

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;
+}

SYCL/AOT/accelerator.cpp

@@ -0,0 +1,12 @@
+//==--- 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

SYCL/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/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 *" %S/Inputs/aot.cpp -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out

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

SYCL/AOT/spec_const_aot.cpp

@@ -0,0 +1,64 @@
+// 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::ONEAPI::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;
+}