Merge branch 'intel' of https://github.com/AidanBeltonS/llvm-test-suite into local_accessor

Author: aidan.belton

SYCL/Assert/assert_in_simultaneous_kernels.cpp

@@ -1,4 +1,6 @@
 // REQUIRES: linux
+// FIXME: Flaky on HIP
+// UNSUPPORTED: hip
 // 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,5 +1,6 @@
-// FIXME flaky output on Level Zero
-// UNSUPPORTED: level_zero
+// FIXME flaky fail on CUDA and HIP
+// UNSUPPORTED: cuda || hip
+//
 // 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/AtomicRef/assignment_atomic64.cpp

@@ -4,9 +4,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// XFAIL: hip
-// Expected failure because hip does not have atomic64 check implementation
-
 #include "assignment.h"
 #include <iostream>
 using namespace sycl;

SYCL/Basic/aspects.cpp

@@ -4,9 +4,6 @@
 // Hip is missing some of the parameters tested here so it fails with NVIDIA
 // XFAIL: hip_nvidia
 
-// XFAIL: hip
-// Expected failure because hip does not have atomic64 check implementation
-
 //==--------------- aspects.cpp - SYCL device test ------------------------==//
 //
 // Returns the various aspects of a device  and platform.

SYCL/Basic/buffer/buffer_allocator.cpp

@@ -1,7 +1,3 @@
-// RUN: %clangxx -fsycl -DSYCL2020_CONFORMANT_APIS -fsycl-targets=%sycl_triple %s -o %t.out
-// RUN: env SYCL_DEVICE_FILTER=opencl %CPU_RUN_PLACEHOLDER %t.out
-// RUN: env SYCL_DEVICE_FILTER=opencl %GPU_RUN_PLACEHOLDER %t.out
-// RUN: env SYCL_DEVICE_FILTER=opencl %ACC_RUN_PLACEHOLDER %t.out
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: env SYCL_DEVICE_FILTER=opencl %CPU_RUN_PLACEHOLDER %t.out
 // RUN: env SYCL_DEVICE_FILTER=opencl %GPU_RUN_PLACEHOLDER %t.out
@@ -15,9 +11,9 @@
 //
 //===----------------------------------------------------------------------===//
 
-// Tests that the default (and explicit with SYCL 2020) buffer_allocator used by
-// buffers are as defined by the spec and will allocate the right amount of
-// memory on the device.
+// Tests that the default and explicit buffer_allocator used by buffers are as
+// defined by the spec and will allocate the right amount of memory on the
+// device.
 
 #include <iostream>
 #include <sycl/sycl.hpp>
@@ -45,40 +41,27 @@ bool checkResult(sycl::buffer<T, 1, AllocT> &Buf, size_t N, std::string &TName,
 template <typename T> bool runTest(sycl::queue &Q, std::string TName) {
   sycl::buffer<T, 1> DefaultBuf{NumElems};
 
-#ifdef SYCL2020_CONFORMANT_APIS
   static_assert(std::is_same_v<decltype(DefaultBuf),
                                sycl::buffer<T, 1, sycl::buffer_allocator<T>>>);
 
   sycl::buffer<T, 1, sycl::buffer_allocator<char>> CharAllocBuf{NumElems};
   sycl::buffer<T, 1, sycl::buffer_allocator<long>> LongAllocBuf{NumElems};
-#else
-  static_assert(std::is_same_v<decltype(DefaultBuf),
-                               sycl::buffer<T, 1, sycl::buffer_allocator>>);
-#endif
 
   Q.submit([&](sycl::handler &CGH) {
      auto DefaultAcc = DefaultBuf.get_access(CGH);
-#ifdef SYCL2020_CONFORMANT_APIS
      auto CharAllocAcc = CharAllocBuf.get_access(CGH);
      auto LongAllocAcc = LongAllocBuf.get_access(CGH);
-#endif
      CGH.parallel_for<TestKernel<T>>(NumElems, [=](sycl::item<1> It) {
        DefaultAcc[It] = static_cast<T>(It[0]);
-#ifdef SYCL2020_CONFORMANT_APIS
        CharAllocAcc[It] = static_cast<T>(It[0]);
        LongAllocAcc[It] = static_cast<T>(It[0]);
-#endif
      });
    }).wait();
 
-#ifdef SYCL2020_CONFORMANT_APIS
   return checkResult(DefaultBuf, NumElems, TName,
                      "buffer_allocator<" + TName + ">") &&
          checkResult(CharAllocBuf, NumElems, TName, "buffer_allocator<char>") &&
          checkResult(LongAllocBuf, NumElems, TName, "buffer_allocator<long>");
-#else
-  return checkResult(DefaultBuf, NumElems, TName, "buffer_allocator");
-#endif
 }
 
 #define RUN_TEST(TYPE, Q) runTest<TYPE>(Q, #TYPE)

SYCL/Basic/buffer/buffer_dev_to_dev.cpp

@@ -1,3 +1,5 @@
+// FIXME flaky fail on HIP
+// UNSUPPORTED: hip
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out

SYCL/Basic/code_location_e2e.cpp

@@ -0,0 +1,156 @@
+// UNSUPPORTED: hip || cuda
+
+// RUN: %clangxx -fsycl -DNDEBUG %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
+// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
+// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
+
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
+// RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
+// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
+
+/*
+    clang++ -fsycl -DNDEBUG -o smyl.bin code_location_e2e.cpp   // <<--- NDEBUG
+   will suppress fileLocation
+
+    -DNDEBUG  <-- not debugging, meaning fileName should not be retrievable (for
+   data privacy reasons)
+
+   Between having to test against NDEBUG and test from kernel code,
+   using assert/abort type of logic is not available. This test just
+   outputs strings that it checks with FileCheck.
+
+   This test uses hardcoded file code location positions. If modified, those
+   may change. pass in any value to signal "report only" and not run tests.
+
+*/
+
+#include <sycl/sycl.hpp>
+using namespace sycl;
+
+// llvm/sycl/doc/design/DeviceLibExtensions.rst
+// Our devicelib support for <cstring> only includes three memory
+// operations, none of the string ones. So we need to provide
+// our own string comparison for kernel calls.
+bool stringsAreSameP(const char *a, const char *b) {
+  // If both are nullptr, then they are the same,
+  if ((a == nullptr) && (b == nullptr))
+    return true;
+  // but if only one, they are not.
+  if ((a == nullptr) || (b == nullptr))
+    return false;
+
+  int index = 0;
+  while (true) {
+    if (a[index] != b[index]) {
+      return false;
+    }
+    if (a[index] == '\0') {
+      return true;
+    } // If we are on this line we know a[i]==b[i].
+    index++;
+  }
+  // We will never arrive here.
+  return true;
+}
+
+template <typename OS> void report(OS &out, detail::code_location code_loc) {
+  out << "function {line:col} => " << code_loc.functionName() << " {"
+      << code_loc.lineNumber() << ":" << code_loc.columnNumber() << "}"
+      << "\n";
+
+  auto fileName = code_loc.fileName();
+  if (fileName == nullptr)
+    out << "nullptr for fileName"
+        << "\n";
+  else
+    out << "fileName: " << code_loc.fileName() << "\n";
+}
+
+template <typename OS>
+void test(OS &out, detail::code_location &code_loc, const char *fileName,
+          const char *funcName, int line, int col) {
+
+  // functionName
+  auto funcNameStr = code_loc.functionName();
+  auto fNameResult =
+      ((funcNameStr != nullptr) && stringsAreSameP(funcNameStr, funcName))
+          ? "OK"
+          : "WRONG";
+  out << "code_location.functionName: " << fNameResult << "\n";
+
+  // lineNumber
+  auto lineNumberResult = (code_loc.lineNumber() == line) ? "OK" : "WRONG";
+  out << "code_location.lineNumber: " << lineNumberResult << "\n";
+
+  // columnNumber
+  auto colNumberResult = (code_loc.columnNumber() == col) ? "OK" : "WRONG";
+  out << "code_location.columnNumber: " << colNumberResult << "\n";
+
+  // fileName
+  auto fileNameStr = code_loc.fileName();
+#ifdef NDEBUG
+  // NDEBUG == not debugging == no fileName (for security).
+  auto fileNameResult =
+      (fileNameStr == nullptr)
+          ? "OK"
+          : "WRONG - fileName should not be present when NDEBUG defined";
+#else
+  auto fileNameResult = stringsAreSameP(fileName, fileNameStr) ? "OK" : "WRONG";
+#endif
+  out << "code_location.fileName: " << fileNameResult << "\n";
+}
+
+int main(int argc, char **argv) {
+  bool testing =
+      (argc <= 1); // Passing in ANY argument means report only, do not test.
+  auto code_loc = sycl::detail::code_location::current(); // <--
+  int EXPECTED_LINE = __LINE__ - 1;
+#ifdef NDEBUG
+  std::cout << "NDEBUG, therefore no fileName" << std::endl;
+#else
+  std::cout << "NDEBUG NOT DEFINED, make sure file name isn't a full path "
+            << std::endl;
+  std::cout << "file name: " << code_loc.fileName() << std::endl;
+#endif
+  std::cout << "------- host test -------" << std::endl;
+  report(std::cout, code_loc);
+
+  if (testing)
+    test(std::cout, code_loc, "code_location_e2e.cpp", "main", EXPECTED_LINE,
+         19);
+
+  std::cout << "------- kernel test -------" << std::endl;
+
+  queue q;
+  q.submit([testing](handler &cgh) {
+    sycl::stream out(2024, 400, cgh);
+    cgh.single_task<class KHALID>([=]() {
+      auto kernel_loc = sycl::detail::code_location::current(); // <--
+      int EXPECTED_LINE = __LINE__ - 1;
+      report(out, kernel_loc);
+
+      if (testing)
+        test(out, kernel_loc, "code_location_e2e.cpp", "operator()",
+             EXPECTED_LINE, 25);
+
+      out << "-------" << sycl::endl;
+    });
+  });
+  q.wait();
+
+  return 0;
+}
+
+// CHECK: ------- host test -------
+// CHECK:      code_location.functionName: OK
+// CHECK-NEXT: code_location.lineNumber: OK
+// CHECK-NEXT: code_location.columnNumber: OK
+// XHECK-NEXT: code_location.fileName: OK
+
+// CHECK: ------- kernel test -------
+// CHECK:      code_location.functionName: OK
+// CHECK-NEXT: code_location.lineNumber: OK
+// CHECK-NEXT: code_location.columnNumber: OK
+// XHECK-NEXT: code_location.fileName: OK

SYCL/Basic/context-with-multiple-devices.cpp

@@ -1,6 +1,3 @@
-// UNSUPPORTED: windows || linux
-//   temporarily disabled
-
 // REQUIRES: accelerator, opencl-aot
 
 // RUN: %clangxx -fsycl -fintelfpga -fsycl-unnamed-lambda %s -o %t2.out

SYCL/Basic/default_device.cpp

@@ -0,0 +1,21 @@
+//==--------- default_device.cpp - SYCL device default ctor test -----------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// 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 <sycl/sycl.hpp>
+
+int main() {
+  sycl::device Dev;
+  assert(Dev == sycl::device{sycl::default_selector_v});
+  return 0;
+}

SYCL/Basic/default_platform.cpp

@@ -0,0 +1,21 @@
+//==------- default_platform.cpp - SYCL platform default ctor test ---------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// 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 <sycl/sycl.hpp>
+
+int main() {
+  sycl::platform Plt;
+  assert(Plt == sycl::platform{sycl::default_selector_v});
+  return 0;
+}