Merge branch 'intel:sycl' into sycl-hostpipe-runtime

Author: zibaiwan

CONTRIBUTING.md

@@ -69,7 +69,7 @@ To contribute:
     can be reviewed, tested and merged independently.
 - For changes which require modification in tests outside of the current repository
   the commit message should contain the link to corresponding test PR.
-  For example: intel/llvm-test-suite#88 or KhronosGroup/SYCL-CTS#65. (see
+  For example: KhronosGroup/SYCL-CTS#65. (see
   [Autolinked references and URLs](https://docs.github.com/en/get-started/writing-on-github/working-with-advanced-formatting/autolinked-references-and-urls)
   for more details). The same message should be present both in commit
   message and in PR description.

llvm/test/Transforms/SPIRITTAnnotations/itt_atomic_load.ll

@@ -1,7 +1,7 @@
 ;; The test serves a purpose to check if Atomic load instruction is being
 ;; annotated by SPIRITTAnnotations pass
 ;;
-;; Compiled from https://github.com/intel/llvm-test-suite/blob/intel/SYCL/AtomicRef/load.cpp
+;; Compiled from https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/AtomicRef/load.cpp
 ;; with following commands:
 ;; clang++ -fsycl -fsycl-device-only load.cpp -o load.bc
 

llvm/test/Transforms/SPIRITTAnnotations/itt_atomic_store.ll

@@ -1,7 +1,7 @@
 ;; The test serves a purpose to check if Atomic store instruction is being
 ;; annotated by SPIRITTAnnotations pass
 ;;
-;; Compiled from https://github.com/intel/llvm-test-suite/blob/intel/SYCL/AtomicRef/load.cpp
+;; Compiled from https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/AtomicRef/load.cpp
 ;; with following commands:
 ;; clang++ -fsycl -fsycl-device-only load.cpp -o load.bc
 

llvm/test/Transforms/SPIRITTAnnotations/itt_barrier.ll

@@ -1,7 +1,8 @@
 ;; The test serves a purpose to check if barrier instruction is being annotated
 ;; by SPIRITTAnnotations pass
 ;;
-;; Compiled from https://github.com/intel/llvm-test-suite/blob/intel/SYCL/KernelAndProgram/kernel-and-program.cpp
+;; Compiled from the SYCL End-to-End test KernelAndProgram/kernel-and-program.cpp last available in
+;; https://github.com/intel/llvm-test-suite/blob/3be0c4dae326f187664cdaa1b9730b446ec4ae29/SYCL/KernelAndProgram/kernel-and-program.cpp
 ;; with following commands:
 ;; clang++ -fsycl -fsycl-device-only kernel-and-program.cpp -o kernel_and_program_optimized.bc
 

sycl/doc/GetStartedGuide.md

@@ -533,10 +533,10 @@ architecture. It is possible to change it by adding
 `-Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=<target>` to the CMake
 variable `SYCL_CLANG_EXTRA_FLAGS`.
 
-#### Run DPC++ E2E test suite
+#### Run DPC++ E2E tests
 
 Follow instructions from the link below to build and run tests:
-[README](https://github.com/intel/llvm-test-suite/tree/intel/SYCL#execution)
+[README](https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/README.md#build-and-run-tests)
 
 #### Run Khronos\* SYCL\* conformance test suite (optional)
 

sycl/doc/design/SharedLibraries.md

@@ -607,4 +607,4 @@ same name as one of the explicitly linked libraries.
 ## Related links
 
 1. Test plan for this feature
-https://github.com/intel/llvm-test-suite/blob/intel/SYCL/TestPlans/DynamicLinkingTestPlan.md
+https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/TestPlans/DynamicLinkingTestPlan.md

sycl/doc/developer/ContributeToDPCPP.md

@@ -49,13 +49,13 @@ Every product change should be accompanied with corresponding test modification
 (adding new test(s), extending, removing or modifying existing test(s)).
 
 There are 3 types of tests which are used for DPC++ toolchain validation:
-* DPC++ in-tree tests
+* DPC++ device-independent tests
 * DPC++ end-to-end (E2E) tests
 * SYCL Conformance Test Suite (CTS)
 
-### DPC++ in-tree tests
+### DPC++ device-independent tests
 
-DPC++ in-tree tests are hosted in this repository. They can be run by
+DPC++ device-independent tests are hosted in this repository. They can be run by
 [check-llvm](/llvm/test), [check-clang](/clang/test),
 [check-llvm-spirv](/llvm-spirv/test) and [check-sycl](/sycl/test) targets.
 These tests are expected not to have hardware (e.g. GPU, FPGA, etc.) or
@@ -144,12 +144,13 @@ When adding new test to `check-sycl`, please consider the following:
 
 ### DPC++ end-to-end (E2E) tests
 
-These tests are extension to
-[LLVM test suite](https://github.com/intel/llvm-test-suite/tree/intel/SYCL)
-and hosted at separate repository.
+These tests are located in [/sycl/test-e2e](/sycl/test-e2e) directory and are not
+configured to be run by default. See
+[End-to-End tests documentation](/sycl/test-e2e/README.md)
+for instructions on how to run them.
+
 A test which requires full stack including backend runtimes (e.g. OpenCL,
-Level Zero or CUDA) should be added to DPC++ E2E test suite following
-[CONTRIBUTING](https://github.com/intel/llvm-test-suite/blob/intel/CONTRIBUTING.md).
+Level Zero or CUDA) should be added to DPC++ E2E tests.
 
 ### SYCL Conformance Test Suite (CTS)
 

sycl/doc/extensions/deprecated/sycl_ext_oneapi_spec_constants.md

@@ -55,10 +55,10 @@ Here the values of specialization constants `SC0` and `SC1` are changed on
 every loop iteration. All what's needed is re-creating a `program` class
 instance, setting new values and rebuilding it via
 `program::build_with_kernel_type`. JIT compiler will effectively replace
-`sc0.get()` and  `sc1.get()` within thhe device code with the corresponding
-constant values (`sc_vals[i][0]` and `sc_vals[i][1]`). Full runnable example
+`sc0.get()` and  `sc1.get()` within the device code with the corresponding
+constant values (`sc_vals[i][0]` and `sc_vals[i][1]`). Full runnable examples
 can be found on
-[github](https://github.com/intel/llvm-test-suite/blob/intel/SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp).
+[github](https://github.com/intel/llvm/tree/sycl/sycl/test-e2e/SpecConstants/2020/).
 
 Specialization constants can be used in programs compiled Ahead-Of-Time, in this
 case a specialization constant takes default value for its type (as specified by

sycl/doc/extensions/experimental/sycl_ext_intel_esimd/README.md

@@ -6,7 +6,7 @@ yet allows to write explicitly vectorized device code. This helps programmer to
 have more control over the generated code and depend less on compiler
 optimizations. The [specification](sycl_ext_intel_esimd.md),
 [API reference](https://intel.github.io/llvm-docs/doxygen/group__sycl__esimd.html), and
-[working code examples](https://github.com/intel/llvm-test-suite/tree/intel/SYCL/ESIMD) are available on the Intel DPC++ project's github.
+[working code examples](https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/ESIMD/) are available on the Intel DPC++ project's github.
 
 **_NOTE:_** _Some parts of this extension is under active development and APIs in the
 `sycl::ext::intel::experimental::esimd` package are subject to change. There are
@@ -45,7 +45,7 @@ a special attribute - `[[intel::sycl_explicit_simd]]`. This tells the compiler t
 this kernel is a ESIMD one and ESIMD APIs can be used inside it. Here the `simd`
 objects and `copy_to` intrinsics are used which are avaiable only in the ESIMD extension.
 Full runnable code sample can be found on the
-[github repo](https://github.com/intel/llvm-test-suite/blob/intel/SYCL/ESIMD/vadd_usm.cpp).
+[github repo](https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/ESIMD/vadd_usm.cp).
 
 #### Compiling and running ESIMD code
 

sycl/doc/extensions/experimental/sycl_ext_intel_esimd/esimd_emulator.md

@@ -73,7 +73,7 @@ List of target platforms supported by CM_EMU is as follows
 Compilation step for ESIMD kernels prepared for ESIMD_EMULATOR backend
 is same as for OpenCL and Level Zero backends. Full runnable code
 sample used below can be found on the [github
-repo](https://github.com/intel/llvm-test-suite/blob/intel/SYCL/ESIMD/vadd_usm.cpp).
+repo](/sycl/test-e2e/ESIMD/vadd_usm.cpp).
 
 To compile using the open-source Intel DPC++ compiler:
 > `$ clang++ -fsycl vadd_usm.cpp`
@@ -92,39 +92,35 @@ so they can't be used simultaneously by a SYCL offload application process. On t
 hand, it is OK to mix the emulator with non-Intel GPU devices or CPU device in
 `ONEAPI_DEVICE_SELECTOR`.
 
-## Running ESIMD examples from [ESIMD test suite](https://github.com/intel/llvm-test-suite/tree/intel/SYCL/ESIMD) on github with ESIMD_EMULATOR backend
+## Running ESIMD examples from [ESIMD test suite](/sycl/test-e2e/ESIMD/) on github with ESIMD_EMULATOR backend
 
 ```
 # Get sources
-git clone https://github.com/intel/llvm-test-suite
-cd llvm-test-suite
+git clone https://github.com/intel/llvm
+cd llvm/sycl/test-e2e
 mkdir build && cd build
 
 # Configure for make utility with compiler tools available in $PATH
 cmake \
- -DCMAKE_CXX_COMPILER=clang++ \
- -DTEST_SUITE_SUBDIRS=SYCL \
- -DSYCL_BE="ext_intel_esimd_emulator" \
- -DSYCL_TARGET_DEVICES="gpu" \
+ -DSYCL_CXX_COMPILER=clang++ \
+ -DSYCL_TEST_E2E_TARGETS="ext_intel_esimd_emulator:gpu" \
  ..
 
 # Build and Run
-make check
+make check-sycl-e2e
 
 # Or, for Ninja utility
-cmake -G Ninja \
- -DCMAKE_CXX_COMPILER=clang++ \
- -DTEST_SUITE_SUBDIRS=SYCL \
- -DSYCL_BE="ext_intel_esimd_emulator" \
- -DSYCL_TARGET_DEVICES="gpu" \
+cmake  -G Ninja \
+ -DSYCL_CXX_COMPILER=clang++ \
+ -DSYCL_TEST_E2E_TARGETS="ext_intel_esimd_emulator:gpu" \
  ..
 
 # Build and Run
-ninja check
+ninja check-sycl-e2e
 
 ```
 
-Note that only [ESIMD Kernels](https://github.com/intel/llvm-test-suite/tree/intel/SYCL/ESIMD) are
+Note that only [ESIMD Kernels](/sycl/test-e2e/ESIMD/) are
 tested with above command examples due to ESIMD_EMULATOR's limitations
 below. And, if 'CM_RT_PLATFORM' is not set, 'skl' is set by default.
 

sycl/doc/extensions/experimental/sycl_ext_intel_esimd/sycl_ext_intel_esimd.md

@@ -668,7 +668,7 @@ looks as below if packed as 8x8 matrix with `uint32_t` elements:
 
 and is passed to `xmx::dpas()` as simd<uint32_t, 8*8> {(b0<<16)|a0, | (b1<<16)|a1, ..., (b7<<16)|a7, (d0<<16)|c0, ..., ..., p7<<16\|o7}.
 
-Some more examples can be found in test for `xmx::dpas()` in [LIT tests](https://github.com/intel/llvm-test-suite/tree/intel/SYCL/ESIMD/dpas).
+Some more examples can be found in test for `xmx::dpas()` in [LIT tests](https://github.com/intel/llvm/tree/sycl/sycl/test-e2e/ESIMD/dpas).
 
 ### Other APIs
 
@@ -907,4 +907,4 @@ int main(void) {
 }
 ```
 more examples can be found in the
-[ESIMD test suite](https://github.com/intel/llvm-test-suite/tree/intel/SYCL/ESIMD) on github.
+[ESIMD test suite](https://github.com/intel/llvm/tree/sycl/sycl/test-e2e/ESIMD) on github.

sycl/include/sycl/group_algorithm.hpp

@@ -315,29 +315,6 @@ reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
 }
 
 // ---- joint_reduce
-template <typename Group, typename Ptr, class BinaryOperation>
-detail::enable_if_t<
-    (is_group_v<std::decay_t<Group>> && detail::is_pointer<Ptr>::value &&
-     detail::is_arithmetic_or_complex<
-         typename detail::remove_pointer<Ptr>::type>::value &&
-     detail::is_plus_or_multiplies_if_complex<
-         typename detail::remove_pointer<Ptr>::type, BinaryOperation>::value),
-    typename detail::remove_pointer<Ptr>::type>
-joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
-#ifdef __SYCL_DEVICE_ONLY__
-  using T = typename detail::remove_pointer<Ptr>::type;
-  T init = detail::identity_for_ga_op<T, BinaryOperation>();
-  return joint_reduce(g, first, last, init, binary_op);
-#else
-  (void)g;
-  (void)first;
-  (void)last;
-  (void)binary_op;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
-#endif
-}
-
 template <typename Group, typename Ptr, typename T, class BinaryOperation>
 detail::enable_if_t<
     (is_group_v<std::decay_t<Group>> && detail::is_pointer<Ptr>::value &&
@@ -373,6 +350,29 @@ joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
 #endif
 }
 
+template <typename Group, typename Ptr, class BinaryOperation>
+detail::enable_if_t<
+    (is_group_v<std::decay_t<Group>> && detail::is_pointer<Ptr>::value &&
+     detail::is_arithmetic_or_complex<
+         typename detail::remove_pointer<Ptr>::type>::value &&
+     detail::is_plus_or_multiplies_if_complex<
+         typename detail::remove_pointer<Ptr>::type, BinaryOperation>::value),
+    typename detail::remove_pointer<Ptr>::type>
+joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  using T = typename detail::remove_pointer<Ptr>::type;
+  T init = detail::identity_for_ga_op<T, BinaryOperation>();
+  return joint_reduce(g, first, last, init, binary_op);
+#else
+  (void)g;
+  (void)first;
+  (void)last;
+  (void)binary_op;
+  throw runtime_error("Group algorithms are not supported on host.",
+                      PI_ERROR_INVALID_DEVICE);
+#endif
+}
+
 // ---- any_of_group
 template <typename Group>
 detail::enable_if_t<is_group_v<std::decay_t<Group>>, bool>

sycl/test-e2e/External/README.md

@@ -56,8 +56,8 @@ is specified binaries will be used directly ignoring other parameters.
 Get sources
 
 ```
-git clone https://github.com/intel/llvm-test-suite
-cd llvm-test-suite
+git clone https://github.com/intel/llvm
+cd llvm/sycl/test-e2e
 mkdir build
 cd build
 ```
@@ -68,13 +68,12 @@ With compiler tools available in the PATH:
 # Configure
 cmake \
  -DCMAKE_CXX_COMPILER=clang++ \
- -DTEST_SUITE_SUBDIRS=SYCL \
- -DCHECK_SYCL_ALL="level_zero:gpu" \
+ -DSYCL_TEST_E2E_TARGETS="level_zero:gpu" \
  -DSYCL_EXTERNAL_TESTS="RSBench" \
  ..
 
 # Build and Run
-make check-sycl-all
+make check-sycl-e2e
 
 ```
 

sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp

@@ -34,13 +34,18 @@ void test(queue q, InputContainer input, OutputContainer output,
       cgh.parallel_for<SpecializationKernelName>(
           nd_range<1>(G, G), [=](nd_item<1> it) {
             group<1> g = it.get_group();
+            auto sg = it.get_sub_group();
             int lid = it.get_local_id(0);
             out[0] = reduce_over_group(g, in[lid], binary_op);
             out[1] = reduce_over_group(g, in[lid], init, binary_op);
             out[2] = joint_reduce(g, in.get_pointer(), in.get_pointer() + N,
                                   binary_op);
             out[3] = joint_reduce(g, in.get_pointer(), in.get_pointer() + N,
                                   init, binary_op);
+            out[4] = joint_reduce(sg, in.get_pointer(), in.get_pointer() + N,
+                                  binary_op);
+            out[5] = joint_reduce(sg, in.get_pointer(), in.get_pointer() + N,
+                                  init, binary_op);
           });
     });
   }
@@ -54,6 +59,10 @@ void test(queue q, InputContainer input, OutputContainer output,
          std::accumulate(input.begin(), input.end(), identity, binary_op));
   assert(output[3] ==
          std::accumulate(input.begin(), input.end(), init, binary_op));
+  assert(output[4] ==
+         std::accumulate(input.begin(), input.end(), identity, binary_op));
+  assert(output[5] ==
+         std::accumulate(input.begin(), input.end(), init, binary_op));
 }
 
 int main() {
@@ -65,7 +74,7 @@ int main() {
 
   constexpr int N = 128;
   std::array<int, N> input;
-  std::array<int, 4> output;
+  std::array<int, 6> output;
   std::iota(input.begin(), input.end(), 0);
   std::fill(output.begin(), output.end(), 0);
 
@@ -93,7 +102,7 @@ int main() {
   // sycl::plus binary operation.
 #ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
   std::array<std::complex<float>, N> input_cf;
-  std::array<std::complex<float>, 4> output_cf;
+  std::array<std::complex<float>, 6> output_cf;
   std::iota(input_cf.begin(), input_cf.end(), 0);
   std::fill(output_cf.begin(), output_cf.end(), 0);
   test<class KernelNamePlusComplexF>(q, input_cf, output_cf,

sycl/test/extensions/inline_asm.cpp

@@ -1,5 +1,5 @@
 // This is a basic acceptance test for inline ASM feature. More tests can be
-// found in https://github.com/intel/llvm-test-suite/tree/intel/SYCL/InlineAsm
+// found in https://github.com/intel/llvm/tree/sycl/sycl/test-e2e/InlineAsm
 // RUN: %clangxx -fsycl -fsyntax-only %s
 
 #include <cmath>

sycl/unittests/CMakeLists.txt

@@ -24,7 +24,7 @@ include(AddSYCLUnitTest)
 
 add_custom_target(check-sycl-unittests)
 
-# TODO PI tests require real hardware and must be moved to intel/llvm-test-suite
+# TODO PI tests require real hardware and must be moved to sycl/test-e2e.
 option(SYCL_PI_TESTS "Enable PI-specific unit tests" ON)
 
 if (SYCL_PI_TESTS)