[SYCL] Added kernel/program compile/build tests

Author: vladimirlaz

SYCL/Basic/KernelAndProgram/basic-program.cpp

@@ -0,0 +1,93 @@
+// XFAIL: cuda
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple  %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+//==--- basic-program.cpp - Basic test of program and kernel APIs ----------==//
+//
+// 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 <CL/sycl.hpp>
+
+#include <iostream>
+
+int main() {
+
+  // Test program and kernel APIs when building a kernel.
+  {
+    cl::sycl::queue q;
+    int data = 0;
+    {
+      cl::sycl::buffer<int, 1> buf(&data, cl::sycl::range<1>(1));
+      cl::sycl::program prg(q.get_context());
+      assert(prg.get_state() == cl::sycl::program_state::none);
+      prg.build_with_kernel_type<class BuiltKernel>();
+      assert(prg.get_state() == cl::sycl::program_state::linked);
+      cl::sycl::vector_class<cl::sycl::vector_class<char>> binaries =
+          prg.get_binaries();
+      assert(prg.has_kernel<class BuiltKernel>());
+      cl::sycl::kernel krn = prg.get_kernel<class BuiltKernel>();
+      cl::sycl::string_class name =
+          krn.get_info<cl::sycl::info::kernel::function_name>();
+      assert(prg.has_kernel(name));
+
+      q.submit([&](cl::sycl::handler &cgh) {
+        auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+        cgh.single_task<class BuiltKernel>(krn, [=]() { acc[0] = acc[0] + 1; });
+      });
+    }
+    assert(data == 1);
+  }
+
+  // Test program and kernel APIs when compiling / linking a kernel.
+  {
+    cl::sycl::queue q;
+    int data = 0;
+    {
+      cl::sycl::buffer<int, 1> buf(&data, cl::sycl::range<1>(1));
+      cl::sycl::program prg(q.get_context());
+      assert(prg.get_state() == cl::sycl::program_state::none);
+      prg.compile_with_kernel_type<class CompiledKernel>();
+      assert(prg.get_state() == cl::sycl::program_state::compiled);
+      prg.link();
+      assert(prg.get_state() == cl::sycl::program_state::linked);
+      cl::sycl::vector_class<cl::sycl::vector_class<char>> binaries =
+          prg.get_binaries();
+      assert(prg.has_kernel<class CompiledKernel>());
+      cl::sycl::kernel krn = prg.get_kernel<class CompiledKernel>();
+      cl::sycl::string_class name =
+          krn.get_info<cl::sycl::info::kernel::function_name>();
+      assert(prg.has_kernel(name));
+
+      q.submit([&](cl::sycl::handler &cgh) {
+        auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+        cgh.single_task<class CompiledKernel>(krn,
+                                              [=]() { acc[0] = acc[0] + 1; });
+      });
+    }
+    assert(data == 1);
+  }
+
+  {
+    sycl::context context;
+    std::vector<sycl::device> devices = context.get_devices();
+
+    sycl::program prg1(context, sycl::property_list{});
+    sycl::program prg2(
+        context, devices,
+        sycl::property_list{sycl::property::buffer::use_host_ptr{}});
+    if (!prg2.has_property<sycl::property::buffer::use_host_ptr>()) {
+      std::cerr << "Line " << __LINE__ << ": Property was not found"
+                << std::endl;
+      return 1;
+    }
+
+    sycl::property::buffer::use_host_ptr Prop =
+        prg2.get_property<sycl::property::buffer::use_host_ptr>();
+  }
+}

SYCL/Basic/KernelAndProgram/build-log.cpp

@@ -0,0 +1,51 @@
+// XFAIL: cuda
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple  %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+//==--- build-log.cpp - Test log message from faild build ----------==//
+//
+// 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 <CL/sycl.hpp>
+
+SYCL_EXTERNAL
+void symbol_that_does_not_exist();
+
+void test() {
+  cl::sycl::queue Queue;
+
+  // Submitting this kernel should result in a compile_program_error exception
+  // with a message indicating that "symbol_that_does_not_exist" is undefined.
+  auto Kernel = []() {
+#ifdef __SYCL_DEVICE_ONLY__
+    symbol_that_does_not_exist();
+#endif
+  };
+
+  std::string Msg;
+  int Result;
+
+  try {
+    Queue.submit([&](cl::sycl::handler &CGH) {
+      CGH.single_task<class SingleTask>(Kernel);
+    });
+    assert(false && "There must be compilation error");
+  } catch (const cl::sycl::compile_program_error &e) {
+    std::string Msg(e.what());
+    assert(Msg.find("symbol_that_does_not_exist") != std::string::npos);
+  } catch (...) {
+    assert(false && "There must be cl::sycl::compile_program_error");
+  }
+}
+
+int main() {
+  test();
+
+  return 0;
+}

SYCL/Basic/KernelAndProgram/cache-build-result.cpp

@@ -0,0 +1,49 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+SYCL_EXTERNAL
+void undefined();
+
+void test() {
+  cl::sycl::queue Queue;
+
+  auto Kernel = []() {
+#ifdef __SYCL_DEVICE_ONLY__
+    undefined();
+#endif
+  };
+
+  std::string Msg;
+  int Result;
+
+  for (int Idx = 0; Idx < 2; ++Idx) {
+    try {
+      Queue.submit([&](cl::sycl::handler &CGH) {
+        CGH.single_task<class SingleTask>(Kernel);
+      });
+      assert(false && "There must be compilation error");
+    } catch (const cl::sycl::compile_program_error &e) {
+      fprintf(stderr, "Exception: %s, %d\n", e.what(), e.get_cl_code());
+      if (Idx == 0) {
+        Msg = e.what();
+        Result = e.get_cl_code();
+      } else {
+        // Exception constantly adds info on its error code in the message
+        assert(Msg.find_first_of(e.what()) == 0 && "Exception text differs");
+        assert(Result == e.get_cl_code() && "Exception code differs");
+      }
+    } catch (...) {
+      assert(false && "There must be cl::sycl::compile_program_error");
+    }
+  }
+}
+
+int main() {
+  test();
+
+  return 0;
+}

SYCL/Basic/KernelAndProgram/get-options.cpp

@@ -0,0 +1,47 @@
+// RUN: %clangxx -fsycl %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 <CL/sycl.hpp>
+
+#include <cassert>
+
+// Check program::get_compile/link/build_options functions
+
+class KernelName;
+void submitKernel() {
+  cl::sycl::queue q;
+  q.submit(
+      [&](cl::sycl::handler &cgh) { cgh.single_task<KernelName>([]() {}); });
+}
+
+int main() {
+  const cl::sycl::string_class CompileOpts{"-cl-opt-disable"};
+  const cl::sycl::string_class LinkOpts{"-cl-fast-relaxed-math"};
+  const cl::sycl::string_class BuildOpts{
+      "-cl-opt-disable -cl-fast-relaxed-math"};
+
+  cl::sycl::context Ctx;
+  cl::sycl::program PrgA{Ctx};
+  assert(PrgA.get_compile_options().empty());
+  assert(PrgA.get_link_options().empty());
+  assert(PrgA.get_build_options().empty());
+
+  PrgA.build_with_kernel_type<KernelName>(BuildOpts);
+  assert(PrgA.get_compile_options().empty());
+  assert(PrgA.get_link_options().empty());
+  assert(PrgA.get_build_options() == (PrgA.is_host() ? "" : BuildOpts));
+
+  cl::sycl::program PrgB{Ctx};
+  PrgB.compile_with_kernel_type<KernelName>(CompileOpts);
+  assert(PrgB.get_compile_options() == (PrgB.is_host() ? "" : CompileOpts));
+  assert(PrgB.get_link_options().empty());
+  assert(PrgB.get_build_options() == (PrgB.is_host() ? "" : CompileOpts));
+
+  PrgB.link(LinkOpts);
+  assert(PrgB.get_compile_options() == (PrgB.is_host() ? "" : CompileOpts));
+  assert(PrgB.get_link_options() == (PrgB.is_host() ? "" : LinkOpts));
+  assert(PrgB.get_build_options() == (PrgB.is_host() ? "" : LinkOpts));
+}