Merge branch 'intel' into ptx_atomic_ldst

Author: t4c1

.github/CODEOWNERS

@@ -14,13 +14,14 @@
 /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
+/SYCL/Plugin/*level_zero* @intel/dpcpp-l0-pi-reviewers
 
 # Explicit SIMD
 /SYCL/ESIMD/ @intel/dpcpp-esimd-reviewers
 
 # BFloat16 conversion
-/SYCL/BFloat16/ @AlexeySotkin @intel/dpcpp-tools-reviewers
+/SYCL/BFloat16/ @intel/dpcpp-tools-reviewers
 
 # Compiler tests
 /SYCL/AOT/ @intel/dpcpp-tools-reviewers

.gitignore

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

SYCL/AOT/accelerator.cpp

@@ -1,12 +1,12 @@
-//==--- accelerator.cpp - AOT compilation for fpga devices using aoc ------==//
+//=-- 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: aoc, accelerator
+// 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/multiple-devices.cpp

@@ -6,7 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 
-// REQUIRES: opencl-aot, ocloc, aoc, cpu, gpu, accelerator
+// REQUIRES: opencl-aot, ocloc, cpu, gpu, accelerator
 // UNSUPPORTED: cuda
 // CUDA is not compatible with SPIR.
 

SYCL/Assert/assert_in_kernels.cpp

@@ -1,7 +1,7 @@
 // REQUIRES: linux
 // FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 // RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true

SYCL/Assert/assert_in_kernels_win.cpp

@@ -1,6 +1,6 @@
 // REQUIRES: windows
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 // RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true

SYCL/Assert/assert_in_multiple_tus.cpp

@@ -1,7 +1,7 @@
 // REQUIRES: linux
 // FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 // RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true

SYCL/Assert/assert_in_multiple_tus_one_ndebug.cpp

@@ -1,7 +1,7 @@
 // REQUIRES: linux
 // FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 // RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true

SYCL/Assert/assert_in_multiple_tus_one_ndebug_win.cpp

@@ -1,6 +1,6 @@
 // REQUIRES: windows
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 // RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true

SYCL/Assert/assert_in_multiple_tus_win.cpp

@@ -1,6 +1,6 @@
 // REQUIRES: windows
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 // RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true

SYCL/Assert/assert_in_one_kernel.cpp

@@ -1,7 +1,7 @@
 // REQUIRES: linux
 // FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 // RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true

SYCL/Assert/assert_in_one_kernel_win.cpp

@@ -1,6 +1,6 @@
 // REQUIRES: windows
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 // RUN: %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true

SYCL/Assert/assert_in_simultaneous_kernels.cpp

@@ -1,7 +1,7 @@
 // REQUIRES: linux
 // FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out %threads_lib
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out %threads_lib
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 //

SYCL/Assert/assert_in_simultaneous_kernels_win.cpp

@@ -1,6 +1,6 @@
 // REQUIRES: windows
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out %threads_lib
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple %s -o %t.out %threads_lib
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 //

SYCL/Assert/assert_in_simultaneously_multiple_tus.cpp

@@ -1,6 +1,6 @@
 // FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out %threads_lib
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out %threads_lib
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 //

SYCL/Assert/assert_in_simultaneously_multiple_tus_one_ndebug.cpp

@@ -1,6 +1,6 @@
 // FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -DSYCL_ENABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_simultaneously_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out %threads_lib
+// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_simultaneously_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out %threads_lib
 // RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
 // RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
 //

SYCL/AtomicRef/add.cpp

@@ -1,44 +1,12 @@
-// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out \
-// RUN: -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_60
+// See https://github.com/intel/llvm-test-suite/issues/867 for detailed status
+// UNSUPPORTED: hip
+
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-device-code-split=per_kernel -fsycl-targets=%sycl_triple %s -o %t.out -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
 #include "add.h"
-#include <iostream>
-using namespace sycl;
-
-// Floating-point types do not support pre- or post-increment
-template <> void add_test<float>(queue q, size_t N) {
-  add_fetch_test<::sycl::ext::oneapi::atomic_ref,
-                 access::address_space::global_space, float>(q, N);
-  add_fetch_test<::sycl::atomic_ref, access::address_space::global_space,
-                 float>(q, N);
-  add_plus_equal_test<::sycl::ext::oneapi::atomic_ref,
-                      access::address_space::global_space, float>(q, N);
-  add_plus_equal_test<::sycl::atomic_ref, access::address_space::global_space,
-                      float>(q, N);
-}
-
-int main() {
-  queue q;
-
-  constexpr int N = 32;
-  add_test<int>(q, N);
-  add_test<unsigned int>(q, N);
-  add_test<float>(q, N);
-
-  // Include long tests if they are 32 bits wide
-  if constexpr (sizeof(long) == 4) {
-    add_test<long>(q, N);
-    add_test<unsigned long>(q, N);
-  }
-
-  // Include pointer tests if they are 32 bits wide
-  if constexpr (sizeof(char *) == 4) {
-    add_test<char *, ptrdiff_t>(q, N);
-  }
 
-  std::cout << "Test passed." << std::endl;
-}
+int main() { add_test_all<access::address_space::global_space>(); }