[SYCL] Add support of function pointers API to SYCL RT (#490)

* [SYCL] Add support of function pointers API to SYCL RT

Signed-off-by: Alexey Sachkov <[email protected]>

Author: AlexeySachkov

sycl/include/CL/sycl.hpp

@@ -21,6 +21,7 @@
 #include <CL/sycl/handler.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/image.hpp>
+#include <CL/sycl/intel/function_pointer.hpp>
 #include <CL/sycl/intel/sub_group.hpp>
 #include <CL/sycl/item.hpp>
 #include <CL/sycl/kernel.hpp>

sycl/include/CL/sycl/detail/pi.def

@@ -24,6 +24,7 @@ _PI_API(piDevicePartition)
 _PI_API(piDeviceRetain)
 _PI_API(piDeviceRelease)
 _PI_API(piextDeviceSelectBinary)
+_PI_API(piextGetDeviceFunctionPointer)
 // Context
 _PI_API(piContextCreate)
 _PI_API(piContextGetInfo)

sycl/include/CL/sycl/detail/pi.h

@@ -409,6 +409,20 @@ pi_result piextDeviceSelectBinary(
   pi_uint32           num_binaries,
   pi_device_binary *  selected_binary);
 
+/// Retrieves a device function pointer to a user-defined function
+/// \arg \c function_name. \arg \c function_pointer_ret is set to 0 if query
+/// failed.
+///
+/// \arg \c program must be built before calling this API. \arg \c device
+/// must present in the list of devices returned by \c get_device method for
+/// \arg \c program.
+///
+pi_result piextGetDeviceFunctionPointer(
+  pi_device        device,
+  pi_program       program,
+  const char *     function_name,
+  pi_uint64 *      function_pointer_ret);
+
 //
 // Context
 //

sycl/include/CL/sycl/detail/program_impl.hpp

@@ -172,6 +172,9 @@ class program_impl {
     return pi::cast<cl_program>(Program);
   }
 
+  RT::PiProgram &getHandleRef() { return Program; }
+  const RT::PiProgram &getHandleRef() const { return Program; }
+
   bool is_host() const { return Context.is_host(); }
 
   template <typename KernelT>

sycl/include/CL/sycl/intel/function_pointer.hpp

@@ -0,0 +1,93 @@
+//==----------- function_pointer.hpp --- SYCL Function pointers ------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/device.hpp>
+#include <CL/sycl/program.hpp>
+#include <CL/sycl/stl.hpp>
+
+#include <type_traits>
+
+namespace cl {
+namespace sycl {
+namespace intel {
+
+// This is a preview extension implementation, intended to provide early access
+// to a feature for review and community feedback.
+//
+// Because the interfaces defined by this header file are not final and are
+// subject to change they are not intended to be used by shipping software
+// products. If you are interested in using this feature in your software
+// product, please let us know!
+
+using device_func_ptr_holder_t = cl::sycl::cl_ulong;
+
+/// \brief this function performs a cast from device_func_ptr_holder_t type
+/// to the provided function pointer type.
+template <
+    class FuncType,
+    typename FuncPtrType = typename std::add_pointer<FuncType>::type,
+    typename std::enable_if<std::is_function<FuncType>::value, int>::type = 0>
+inline FuncPtrType to_device_func_ptr(device_func_ptr_holder_t FptrHolder) {
+  return reinterpret_cast<FuncPtrType>(FptrHolder);
+}
+
+template <class FuncType>
+using enable_if_is_function_pointer_t = typename std::enable_if<
+    std::is_pointer<FuncType>::value &&
+        std::is_function<typename std::remove_pointer<FuncType>::type>::value,
+    int>::type;
+
+/// \brief this function can be used only on host side to obtain device function
+/// pointer for the specified function.
+///
+/// \param F - pointer to function to make it work for SYCL Host device
+/// \param FuncName - name of the function. Please note that by default names of
+/// functions are mangled since SYCL is a C++. To avoid the need ot specifying
+/// mangled name here, use `extern "C"`
+/// \param P - sycl::program object which will be used to extract device
+/// function pointer
+/// \param D - sycl::device object which will be used to extract device
+/// function pointer
+///
+/// \returns device_func_ptr_holder_t object which can be used inside a device
+/// code. This object must be converted back to a function pointer using
+/// `to_device_func_ptr` prior to actual usage.
+///
+/// Returned value is valid only within device code which was compiled for the
+/// specified program and device. Returned value invalidates whenever program
+/// is released or re-built
+template <class FuncType, enable_if_is_function_pointer_t<FuncType> = 0>
+device_func_ptr_holder_t get_device_func_ptr(FuncType F, const char *FuncName,
+                                             program &P, device &D) {
+  // TODO: drop function name argument and map host function pointer directly to
+  // a device function pointer
+  if (D.is_host()) {
+    return reinterpret_cast<device_func_ptr_holder_t>(F);
+  }
+
+  if (program_state::linked != P.get_state()) {
+    throw invalid_parameter_error(
+        "Program must be built before passing to get_device_func_ptr");
+  }
+
+  device_func_ptr_holder_t FPtr = 0;
+  // FIXME: return value must be checked here, but since we cannot yet check
+  // if corresponding extension is supported, let's silently ignore it here.
+  PI_CALL_RESULT(RT::piextGetDeviceFunctionPointer(
+      detail::pi::cast<pi_device>(detail::getSyclObjImpl(D)->getHandleRef()),
+      detail::pi::cast<pi_program>(detail::getSyclObjImpl(P)->getHandleRef()),
+      FuncName, &FPtr));
+
+  return FPtr;
+}
+
+} // namespace intel
+} // namespace sycl
+} // namespace cl

sycl/source/detail/pi_opencl.cpp

@@ -241,6 +241,39 @@ pi_result OCL(piSamplerCreate)(pi_context context,
   return error_code;
 }
 
+pi_result OCL(piextGetDeviceFunctionPointer)(pi_device device,
+                                             pi_program program,
+                                             const char *func_name,
+                                             pi_uint64 *function_pointer_ret) {
+  pi_platform platform;
+  PI_CALL(piDeviceGetInfo(device, PI_DEVICE_INFO_PLATFORM, sizeof(platform),
+                          &platform, nullptr));
+  using FuncT =
+      cl_int(CL_API_CALL *)(cl_device_id, cl_program, const char *, cl_ulong *);
+
+  // TODO: add check that device supports corresponding extension
+  FuncT func_ptr =
+      reinterpret_cast<FuncT>(clGetExtensionFunctionAddressForPlatform(
+          cast<cl_platform_id>(platform),
+          "clGetDeviceFunctionPointerINTEL"));
+  // TODO: once we have check that device supports corresponding extension,
+  // we can insert an assertion that func_ptr is not nullptr. For now, let's
+  // just return an error if failed to query such function
+  // PI_ASSERT(
+  //     func_ptr != nullptr,
+  //     "Failed to get address of clGetDeviceFunctionPointerINTEL function");
+
+  if (!func_ptr) {
+    if (function_pointer_ret)
+      *function_pointer_ret = 0;
+    return PI_INVALID_DEVICE;
+  }
+
+  return PI_CALL_RESULT(func_ptr(cast<cl_device_id>(device),
+                                 cast<cl_program>(program), func_name,
+                                 function_pointer_ret));
+}
+
 // Forward calls to OpenCL RT.
 #define _PI_CL(pi_api, ocl_api)                     \
 decltype(::pi_api) * pi_api##OclPtr =               \
@@ -256,6 +289,7 @@ _PI_CL(piDevicePartition,    clCreateSubDevices)
 _PI_CL(piDeviceRetain,       clRetainDevice)
 _PI_CL(piDeviceRelease,      clReleaseDevice)
 _PI_CL(piextDeviceSelectBinary, OCL(piextDeviceSelectBinary))
+_PI_CL(piextGetDeviceFunctionPointer, OCL(piextGetDeviceFunctionPointer))
   // Context
 _PI_CL(piContextCreate,     clCreateContext)
 _PI_CL(piContextGetInfo,    clGetContextInfo)

sycl/test/function-pointers/fp-as-kernel-arg.cpp

@@ -0,0 +1,74 @@
+// RUN: %clangxx -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out -lOpenCL
+// RUN: env SYCL_DEVICE_TYPE=HOST %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
+// UNSUPPORTED: windows
+
+#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::intel::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::intel::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/test/function-pointers/pass-fp-through-buffer.cpp

@@ -0,0 +1,93 @@
+// RUN: %clangxx -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out -lOpenCL
+// RUN: env SYCL_DEVICE_TYPE=HOST %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
+// UNSUPPORTED: windows
+
+#include <CL/sycl.hpp>
+
+#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::intel::device_func_ptr_holder_t> DispatchTable(2);
+  {
+    auto DTAcc =
+        DispatchTable.get_access<cl::sycl::access::mode::discard_write>();
+    DTAcc[0] = cl::sycl::intel::get_device_func_ptr(&add, "add", P, D);
+    DTAcc[1] = cl::sycl::intel::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::intel::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;
+}