DPC++ Runtime
Runtime libraries for oneAPI DPC++
kernel_handler.hpp
Go to the documentation of this file.
1 //==------ kernel_handler.hpp -- SYCL standard header file -----*- C++ -*---==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <sycl/detail/defines.hpp> // for __SYCL_TYPE
12 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
13 #include <sycl/detail/pi.h> // for PI_ERROR_INVALID_OPERATION
14 #include <sycl/exception.hpp> // for feature_not_supported
15 
16 #ifdef __SYCL_DEVICE_ONLY__
17 #include <CL/__spirv/spirv_ops.hpp>
18 #endif
19 
20 #include <type_traits> // for remove_reference_t
21 
22 #ifdef __SYCL_DEVICE_ONLY__
23 // Get the value of the specialization constant with given symbolic ID.
24 // `SymbolicID` is a unique string ID of a specialization constant.
25 // `DefaultValue` contains a pointer to a global variable with the initializer,
26 // which should be used as the default value of the specialization constants.
27 // `RTBuffer` is a pointer to a runtime buffer, which holds values of all
28 // specialization constant and should be used if native specialization constants
29 // are not available.
30 template <typename T>
31 __DPCPP_SYCL_EXTERNAL T __sycl_getScalar2020SpecConstantValue(
32  const char *SymbolicID, const void *DefaultValue, const void *RTBuffer);
33 
34 template <typename T>
35 __DPCPP_SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue(
36  const char *SymbolicID, const void *DefaultValue, const void *RTBuffer);
37 #endif
38 
39 namespace sycl {
40 inline namespace _V1 {
44 class __SYCL_TYPE(kernel_handler) kernel_handler {
45 public:
46  template <auto &S>
47  __SYCL_ALWAYS_INLINE typename std::remove_reference_t<decltype(S)>::value_type
48  get_specialization_constant() {
49 #ifdef __SYCL_DEVICE_ONLY__
50  return getSpecializationConstantOnDevice<S>();
51 #else
52  // TODO: add support of host device
53  throw sycl::feature_not_supported(
54  "kernel_handler::get_specialization_constant() is not yet supported by "
55  "host device.",
56  PI_ERROR_INVALID_OPERATION);
57 #endif // __SYCL_DEVICE_ONLY__
58  }
59 
60 private:
61  void __init_specialization_constants_buffer(
62  char *SpecializationConstantsBuffer = nullptr) {
63  MSpecializationConstantsBuffer = SpecializationConstantsBuffer;
64  }
65 
66 #ifdef __SYCL_DEVICE_ONLY__
67  template <
68  auto &S,
69  typename T = typename std::remove_reference_t<decltype(S)>::value_type,
70  std::enable_if_t<std::is_fundamental_v<T>> * = nullptr>
71  __SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() {
72  const char *SymbolicID = __builtin_sycl_unique_stable_id(S);
73  return __sycl_getScalar2020SpecConstantValue<T>(
74  SymbolicID, &S, MSpecializationConstantsBuffer);
75  }
76  template <
77  auto &S,
78  typename T = typename std::remove_reference_t<decltype(S)>::value_type,
79  std::enable_if_t<std::is_compound_v<T>> * = nullptr>
80  __SYCL_ALWAYS_INLINE T getSpecializationConstantOnDevice() {
81  const char *SymbolicID = __builtin_sycl_unique_stable_id(S);
82  return __sycl_getComposite2020SpecConstantValue<T>(
83  SymbolicID, &S, MSpecializationConstantsBuffer);
84  }
85 #endif // __SYCL_DEVICE_ONLY__
86 
87  char *MSpecializationConstantsBuffer = nullptr;
88 };
89 
90 } // namespace _V1
91 } // namespace sycl
#define __SYCL_ALWAYS_INLINE
#define __DPCPP_SYCL_EXTERNAL
Definition: access.hpp:18