Amx sys matrix

.github/CODEOWNERS

@@ -0,0 +1,37 @@
+* @pvchupin
+
+# Use runtime team as the umbrella for most of the tests
+/SYCL/ @intel/llvm-reviewers-runtime
+
+# SYCL sub-directory matchers are grouped by code owner first, followed by
+# alphabetical order within the group. Please, keep this ordering.
+
+# Group algorithms
+/SYCL/GroupAlgorithm/ @Pennycook @intel/llvm-reviewers-runtime
+/SYCL/GroupLocalMemory/ @Pennycook @intel/llvm-reviewers-runtime
+/SYCL/SubGroup/ @Pennycook @intel/llvm-reviewers-runtime
+/SYCL/SubGroupMask/ @Pennycook @intel/llvm-reviewers-runtime
+
+# Plugin interface for Level Zero
+/SYCL/Plugin/*level-zero* @intel/dpcpp-l0-pi-reviewers
+/SYCL/Plugin/*level_zero* @intel/dpcpp-l0-pi-reviewers
+
+# Explicit SIMD
+/SYCL/ESIMD/ @intel/dpcpp-esimd-reviewers
+
+# BFloat16 conversion
+/SYCL/BFloat16/ @intel/dpcpp-tools-reviewers
+
+# Compiler tests
+/SYCL/AOT/ @intel/dpcpp-tools-reviewers
+/SYCL/DeviceCodeSplit/ @intel/dpcpp-tools-reviewers
+/SYCL/SeparateCompile/ @intel/dpcpp-tools-reviewers
+
+# Printf
+/SYCL/Printf/ @intel/dpcpp-tools-reviewers
+
+# Specialization constant
+/SYCL/SpecConstants/ @intel/dpcpp-tools-reviewers
+
+# invoke_simd
+/SYCL/InvokeSimd/ @rolandschulz @kbobrovs

.github/workflows/clang-format.yml

@@ -0,0 +1,34 @@
+name: clang-format-check
+
+on:
+  pull_request:
+    branches:
+    - intel
+
+jobs:
+  build:
+    runs-on: ubuntu-latest
+
+    container:
+      image: ghcr.io/intel/llvm/sycl_ubuntu2004_nightly:no-drivers
+
+    steps:
+    - uses: actions/checkout@v2
+      with:
+        fetch-depth: 2
+
+    - name: Run clang-format for the patch
+      shell: bash {0}
+      run: |
+        git config --global --add safe.directory /__w/llvm-test-suite/llvm-test-suite
+        git clang-format ${GITHUB_SHA}^1
+        git diff > ./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"

.gitignore

@@ -2,3 +2,4 @@
 # External/*
 /test-suite-externals
 *.pyc
+/build*

CONTRIBUTING.md

@@ -0,0 +1,88 @@
+# 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
+
+### Testing
+
+- See [SYCL/README.md](SYCL/README.md) for instructions.
+
+### 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/).
+  - PR description should follow same rules as commit message. It is used as
+    commit message on the final merge.
+- Changes addressing comments made during code review should be added as a
+  separate commits to the same PR.
+- CI will run checks which are prerequisites for submitting PR:
+  - clang-format-check/build checks that the patch matches coding style
+    (see [clang-format](https://clang.llvm.org/docs/ClangFormat.html));
+  - Jenkins/pre-ci-cuda - runs all related tests on CUDA backend for GPU device
+    on Ubuntu 18.04;
+  - Jenkins/pre-ci-linux - runs all related tests on Ubuntu 18.04 machine with
+  Level_Zero backend (GPU device) and OpenCL backend (CPU, GPU and FPGA
+  emulator devices);
+  - Jenkins/pre-ci-windows - runs all related tests on Windows Server 2019 with
+  Level_Zero backend (GPU device) and OpenCL backend (CPU, GPU and FPGA
+  emulator devices).
+
+The last three checks are done for the latest available nightly build for DPC++
+compiler and runtime from [intel/llvm](https://github.com/intel/llvm). The
+build happens around 18:00 UTC if there are new commits since previous build.
+
+Once the PR is approved and all checks have passed, the pull request is
+ready for merge.
+
+### Merge
+
+Project maintainers merge pull requests by "Squash and merge". PR description
+is used as final commit message.
+
+\*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[A-Za-z0-9_-]*) *:|expected-"

SYCL/AOT/Inputs/aot.cpp

@@ -0,0 +1,74 @@
+//==--- 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<sycl::cl_int, array_size> A = {{1, 2, 3, 4}}, B = {{1, 2, 3, 4}},
+                                       C;
+  std::array<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 - 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/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/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/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/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