[SYCL] next portion of the tests (#33)

* Add/move tests covering different kernel param types
* Add test covering loading kernel images from file
* Added kernel/program compile/build tests
* Add multi_ptr test
* add function pointer tests
* add functor tests
* test for context with multiple devices
* XFAIL kernel and program tests failing on CUDA

Author: vladimirlaz

SYCL/Basic/config/allowlist.cpp renamed to SYCL/Basic/Config/allowlist.cpp

@@ -20,8 +20,9 @@ using namespace cl;
 
 static void replaceSpecialCharacters(std::string &Str) {
   // Replace common special symbols with '.' which matches to any character
-  std::replace_if(Str.begin(), Str.end(),
-                  [](const char Sym) { return '(' == Sym || ')' == Sym; }, '.');
+  std::replace_if(
+      Str.begin(), Str.end(),
+      [](const char Sym) { return '(' == Sym || ')' == Sym; }, '.');
 }
 
 int main() {

SYCL/Basic/config/config.cpp renamed to SYCL/Basic/Config/config.cpp

@@ -0,0 +1,44 @@
+// UNSUPPORTED: cuda
+// CUDA does not support SPIR-V.
+
+// RUN: %clangxx -fsycl-device-only -fno-sycl-use-bitcode -Xclang -fsycl-int-header=%t.h -c %s -o %t.spv -I %sycl_include -Xclang -verify-ignore-unexpected=note,warning -Wno-sycl-strict
+// RUN: %clangxx -include %t.h %s -o %t.out -lsycl -I %sycl_include -Xclang -verify-ignore-unexpected=note,warning
+// RUN: env SYCL_BE=%sycl_be SYCL_USE_KERNEL_SPV=%t.spv %t.out | FileCheck %s
+// CHECK: Passed
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+using namespace cl::sycl;
+
+int main(int argc, char **argv) {
+  int data = 5;
+
+  try {
+    queue myQueue;
+    buffer<int, 1> buf(&data, range<1>(1));
+
+    event e = myQueue.submit([&](handler &cgh) {
+      auto ptr = buf.get_access<access::mode::read_write>(cgh);
+
+      cgh.single_task<class my_kernel>([=]() { ptr[0]++; });
+    });
+    e.wait_and_throw();
+
+  } catch (cl::sycl::exception const &e) {
+    std::cerr << "SYCL exception caught:\n";
+    std::cerr << e.what() << "\n";
+    return 2;
+  } catch (...) {
+    std::cerr << "unknown exception caught\n";
+    return 1;
+  }
+
+  if (data == 6) {
+    std::cout << "Passed\n";
+    return 0;
+  } else {
+    std::cout << "Failed: " << data << "!= 6(gold)\n";
+    return 1;
+  }
+}

SYCL/Basic/Config/kernel_from_file.cpp

@@ -0,0 +1,79 @@
+// UNSUPPORTED: windows
+// UNSUPPORTED: cuda || level_zero
+// CUDA does not support the function pointer as kernel argument extension.
+// Hangs on level zero
+
+// RUN: %clangxx -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// FIXME: This test should use runtime early exit once correct check for
+// corresponding extension is implemented
+
+#include <CL/sycl.hpp>
+
+#include <algorithm>
+#include <iostream>
+#include <vector>
+
+[[intel::device_indirectly_callable]] extern "C" int add(int A, int B) {
+  return A + B;
+}
+
+int main() {
+  const int Size = 10;
+  std::vector<long> A(Size, 1);
+  std::vector<long> B(Size, 2);
+
+  cl::sycl::queue Q;
+  cl::sycl::device D = Q.get_device();
+  cl::sycl::context C = Q.get_context();
+  cl::sycl::program P(C);
+
+  P.build_with_kernel_type<class K>();
+  cl::sycl::kernel KE = P.get_kernel<class K>();
+
+  auto FptrStorage = cl::sycl::ONEAPI::get_device_func_ptr(&add, "add", P, D);
+  if (!D.is_host()) {
+    // FIXME: update this check with query to supported extension
+    // For now, we don't have runtimes that report required OpenCL extension and
+    // it is hard to understand should this functionality be supported or not.
+    // So, let's skip this test if FptrStorage is 0, which means that by some
+    // reason we failed to obtain device function pointer. Just to avoid false
+    // alarms
+    if (0 == FptrStorage) {
+      std::cout << "Test PASSED. (it was actually skipped)" << std::endl;
+      return 0;
+    }
+  }
+
+  cl::sycl::buffer<long> BufA(A.data(), cl::sycl::range<1>(Size));
+  cl::sycl::buffer<long> BufB(B.data(), cl::sycl::range<1>(Size));
+
+  Q.submit([&](cl::sycl::handler &CGH) {
+    auto AccA =
+        BufA.template get_access<cl::sycl::access::mode::read_write>(CGH);
+    auto AccB = BufB.template get_access<cl::sycl::access::mode::read>(CGH);
+    CGH.parallel_for<class K>(
+        KE, cl::sycl::range<1>(Size), [=](cl::sycl::id<1> Index) {
+          auto Fptr =
+              cl::sycl::ONEAPI::to_device_func_ptr<decltype(add)>(FptrStorage);
+          AccA[Index] = Fptr(AccA[Index], AccB[Index]);
+        });
+  });
+
+  auto HostAcc = BufA.get_access<cl::sycl::access::mode::read>();
+  auto *Data = HostAcc.get_pointer();
+
+  if (std::all_of(Data, Data + Size, [](long V) { return V == 3; })) {
+    std::cout << "Test PASSED." << std::endl;
+  } else {
+    std::cout << "Test FAILED." << std::endl;
+    for (int I = 0; I < Size; ++I) {
+      std::cout << HostAcc[I] << " ";
+    }
+    std::cout << std::endl;
+  }
+
+  return 0;
+}

SYCL/Basic/FunctionPointers/fp-as-kernel-arg.cpp

@@ -0,0 +1,98 @@
+// UNSUPPORTED: windows
+// UNSUPPORTED: cuda || level_zero
+// CUDA does not support the function pointer as kernel argument extension.
+// Hangs on level zero
+
+// RUN: %clangxx -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// FIXME: This test should use runtime early exit once correct check for
+// corresponding extension is implemented
+
+#include <CL/sycl.hpp>
+
+#include <algorithm>
+#include <iostream>
+#include <vector>
+
+[[intel::device_indirectly_callable]] extern "C" int add(int A, int B) {
+  return A + B;
+}
+
+[[intel::device_indirectly_callable]] extern "C" int sub(int A, int B) {
+  return A - B;
+}
+
+int main() {
+  const int Size = 10;
+
+  cl::sycl::queue Q;
+  cl::sycl::device D = Q.get_device();
+  cl::sycl::context C = Q.get_context();
+  cl::sycl::program P(C);
+
+  P.build_with_kernel_type<class K>();
+  cl::sycl::kernel KE = P.get_kernel<class K>();
+
+  cl::sycl::buffer<cl::sycl::ONEAPI::device_func_ptr_holder_t> DispatchTable(2);
+  {
+    auto DTAcc =
+        DispatchTable.get_access<cl::sycl::access::mode::discard_write>();
+    DTAcc[0] = cl::sycl::ONEAPI::get_device_func_ptr(&add, "add", P, D);
+    DTAcc[1] = cl::sycl::ONEAPI::get_device_func_ptr(&sub, "sub", P, D);
+    if (!D.is_host()) {
+      // FIXME: update this check with query to supported extension
+      // For now, we don't have runtimes that report required OpenCL extension
+      // and it is hard to understand should this functionality be supported or
+      // not. So, let's skip this test if DTAcc[i] is 0, which means that by
+      // some reason we failed to obtain device function pointer. Just to avoid
+      // false alarms
+      if (0 == DTAcc[0] || 0 == DTAcc[1]) {
+        std::cout << "Test PASSED. (it was actually skipped)" << std::endl;
+        return 0;
+      }
+    }
+  }
+
+  for (int Mode = 0; Mode < 2; ++Mode) {
+    std::vector<int> A(Size, 1);
+    std::vector<int> B(Size, 2);
+
+    cl::sycl::buffer<int> bufA(A.data(), cl::sycl::range<1>(Size));
+    cl::sycl::buffer<int> bufB(B.data(), cl::sycl::range<1>(Size));
+
+    Q.submit([&](cl::sycl::handler &CGH) {
+      auto AccA =
+          bufA.template get_access<cl::sycl::access::mode::read_write>(CGH);
+      auto AccB = bufB.template get_access<cl::sycl::access::mode::read>(CGH);
+      auto AccDT =
+          DispatchTable.template get_access<cl::sycl::access::mode::read>(CGH);
+      CGH.parallel_for<class K>(
+          KE, cl::sycl::range<1>(Size), [=](cl::sycl::id<1> Index) {
+            auto FP = cl::sycl::ONEAPI::to_device_func_ptr<int(int, int)>(
+                AccDT[Mode]);
+
+            AccA[Index] = FP(AccA[Index], AccB[Index]);
+          });
+    });
+
+    auto HostAcc = bufA.get_access<cl::sycl::access::mode::read>();
+
+    int Reference = Mode == 0 ? 3 : -1;
+    auto *Data = HostAcc.get_pointer();
+
+    if (std::all_of(Data, Data + Size,
+                    [=](long V) { return V == Reference; })) {
+      std::cout << "Test " << Mode << " PASSED." << std::endl;
+    } else {
+      std::cout << "Test " << Mode << " FAILED." << std::endl;
+      for (int I = 0; I < Size; ++I) {
+        std::cout << HostAcc[I] << " ";
+      }
+      std::cout << std::endl;
+    }
+  }
+
+  return 0;
+}

SYCL/Basic/FunctionPointers/pass-fp-through-buffer.cpp

@@ -0,0 +1,60 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -o %t.out %s
+// 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 <CL/sycl.hpp>
+#include <cassert>
+
+constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
+constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;
+
+struct SecondBase {
+  SecondBase(int _E) : E(_E) {}
+  int E;
+};
+
+struct InnerFieldBase {
+  InnerFieldBase(int _D) : D(_D) {}
+  int D;
+};
+
+struct InnerField : public InnerFieldBase {
+  InnerField(int _C, int _D) : C(_C), InnerFieldBase(_D) {}
+  int C;
+};
+
+struct Base {
+  Base(int _B, int _C, int _D) : B(_B), InnerObj(_C, _D) {}
+  int B;
+  InnerField InnerObj;
+};
+
+struct Derived : public Base, public SecondBase {
+  Derived(int _A, int _B, int _C, int _D, int _E,
+          cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &_Acc)
+      : A(_A), Acc(_Acc), /*Out(_Out),*/ Base(_B, _C, _D), SecondBase(_E) {}
+  void operator()() const {
+    Acc[0] = this->A + this->B + this->InnerObj.C + this->InnerObj.D + this->E;
+  }
+
+  int A;
+  cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> Acc;
+};
+
+int main() {
+  int A[] = {10};
+  {
+    cl::sycl::queue Q;
+    cl::sycl::buffer<int, 1> Buf(A, 1);
+
+    Q.submit([&](cl::sycl::handler &cgh) {
+      auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
+      Derived F = {1, 2, 3, 4, 5, Acc /*, Out*/};
+      cgh.single_task(F);
+    });
+  }
+  assert(A[0] == 15);
+  return 0;
+}