Skip to content

Latest commit

 

History

History
563 lines (448 loc) · 18.9 KB

sycl_ext_oneapi_annotated_arg.asciidoc

File metadata and controls

563 lines (448 loc) · 18.9 KB
Raw

sycl_ext_oneapi_annotated_arg

Notice

Copyright (c) 2022-2023 Intel Corporation. All rights reserved.

Khronos® is a registered trademark and SYCL™ and SPIR™ are trademarks of The Khronos Group Inc. OpenCL™ is a trademark of Apple Inc. used by permission by Khronos.

Contact

To report problems with this extension, please open a new issue at: https://github.com/intel/llvm/issues

Contributors

Joe Garvey, Intel
Abhishek Tiwari, Intel

Dependencies

This extension is written against the SYCL 2020 specification, Revision 5 and the following extensions:

Status

This is an experimental extension specification, intended to provide early access to features and gather community feedback. Interfaces defined in this specification are implemented in DPC++, but they are not finalized and may change incompatibly in future versions of DPC++ without prior notice. Shipping software products should not rely on APIs defined in this specification.

Overview

This extension introduces a wrapper class sycl::ext::oneapi::experimental::annotated_arg that provides a mechanism to attach compile-time constant information to kernel arguments in a reliable manner.

Note

In this document, we use the shortened form annotated_arg to refer to the proposed sycl::ext::oneapi::experimental::annotated_arg class.

The purpose of this document is to clearly describe and specify annotated_arg and related concepts, types, and mechanisms, and to give examples and context for their usage.

Specification

Feature Test Macro

This extension provides a feature-test macro as described in the core SYCL specification section 6.3.3 "Feature test macros". Therefore, an implementation supporting this extension must predefine the macro SYCL_EXT_ONEAPI_ANNOTATED_ARG to one of the values defined in the table below. Applications can test for the existence of this macro to determine if the implementation supports this feature, or applications can test the macro’s value to determine which of the extension’s APIs the implementation supports.

Value Description

1

Initial extension version. Base features are supported.

Introduction

The annotated_arg class enables users to attach compile-time constant information to kernel arguments with properties.

The example below shows a use case with some properties.

Toy Example
using sycl::ext::oneapi::experimental;
struct MyKernel {
  annotated_arg<int, properties<PropA>> a;
  annotated_arg<int, properties<PropB>> b;
  annotated_arg<int*, properties<PropC>> c;
  MyKernel(int a_, int b_, int* c_) : a(a_), b(b_), c(c_) { }
  void operator()() const {
    *c = a + b;
  }
};

int main () {
  sycl::queue q;
  auto result_ptr = ...
  q.single_task(MyKernel{2, 5, result_ptr}).wait();
  ...
  sycl::free(result_ptr, q);
}

PropA, PropB and PropC are expected to be preserved on the kernel arguments in a reliable manner.

Note
 PropC will only influence the kernel argument and not any pointer dereference sites. Use the sycl::ext::oneapi::experimental::annotated_ptr class template to apply properties that must be preserved at pointer dereference sites.

The example below shows a type of use-case which is not meant to be supported by annotated_arg:

Unsupported Usage Example
using sycl::ext::oneapi::experimental;
struct MyType {
  annotated_arg<int, properties<PropC>> a;
  annotated_arg<int, properties<PropD>> b;
};

struct MyKernel {
  MyType arg_a;
  MyType arg_b;
  ...
  void operator()() const {
    ...
  }
};

It is illegal to apply annotated_arg to members of kernel arguments. In the above example, encapsulating annotated_arg within MyType is illegal.

Representation of annotated_arg

namespace sycl::ext::oneapi::experimental {
template < typename T, typename PropertyListT = empty_properties_t>
class annotated_arg {
  ...

annotated_arg is a class template, parameterized by the type of the underlying allocation T and a list of associated properties specified by PropertyListT.

T can be any type except the following types or a structure containing one of the following types:

  • sycl::accessor

  • sycl::stream

  • sycl::local_accessor

  • sycl::unsampled_image_accessor

  • sycl::sampled_image_accessor

  • sycl::half

The properties supported with annotated_arg may be defined in separate extensions. Please note that there cannot be duplicated property in a properties list. Otherwise, a compiler time error is triggered.

If the type T is trivially copyable, then annotated_arg<T, …​> is also trivially copyable.

If the type T is device copyable, then annotated_arg<T, …​> is also device copyable and the implementation sets the is_device_copyable_v trait to true for this type.

The section below describes the constructors and member functions for annotated_arg.

The template parameter T in the definition of annotated_arg template below must be a legal parameter type as defined by the SYCL specification. Given annotated_arg<T, …​>, T must be a device copy-able type.

namespace sycl::ext::oneapi::experimental {
template <typename T, typename PropertyListT = empty_properties_t>
class annotated_arg {
  public:
    annotated_arg() noexcept;
    annotated_arg(const T& v_, const PropertyListT &P = properties{}) noexcept;
    template<typename... PropertyValueTs>
    annotated_arg(const T& v_, PropertyValueTs... props) noexcept;

    template <typename T2, typename P> explicit annotated_arg(
      const annotated_arg<T2, P>&) noexcept;
    template <typename T2, typename PropertyListU, typename PropertyListV>
    explicit annotated_arg(const annotated_arg<T2, PropertyListU>&,
      properties<PropertyListV>) noexcept;

    annotated_arg(const annotated_arg&) = default;
    annotated_arg& operator=(annotated_arg&) = default;

    // Conversion operator to convert to the underlying type
    operator T() noexcept;
    operator T() const noexcept;

    // Available if type T is pointer type
    T operator->() const noexcept;

    // Available if the operator[] is valid for objects of type T, return
    // type will match the return type of T::operator[](std::ptrdiff_t)
    /* ... */ operator[](std::ptrdiff_t idx) const noexcept;

    template<typename propertyT>
    static constexpr bool has_property();

    // The return type is an unspecified internal class used to represent
    // instances of propertyT
    template<typename propertyT>
    static constexpr /*unspecified*/ get_property();

    // Overloaded arithmetic operators
    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() + std::declval<T2>())>
    R operator+(const annotated_arg<T2, PropertyList2> &other) const;

    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() - std::declval<T2>())>
    R operator-(const annotated_arg<T2, PropertyList2> &other) const;

    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() * std::declval<T2>())>
    R operator*(const annotated_arg<T2, PropertyList2> &other) const;

    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() / std::declval<T2>())>
    R operator/(const annotated_arg<T2, PropertyList2> &other) const;

    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() % std::declval<T2>())>
    R operator%(const annotated_arg<T2, PropertyList2> &other) const;

    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() & std::declval<T2>())>
    R operator&(const annotated_arg<T2, PropertyList2> &other) const;

    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() | std::declval<T2>())>
    R operator|(const annotated_arg<T2, PropertyList2> &other) const;

    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() ^ std::declval<T2>())>
    R operator^(const annotated_arg<T2, PropertyList2> &other) const;

    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() >> std::declval<T2>())>
    R operator>>(const annotated_arg<T2, PropertyList2> &other) const;

    template <typename T2, typename PropertyList2,
              typename R = decltype(std::declval<T>() << std::declval<T2>())>
    R operator<<(const annotated_arg<T2, PropertyList2> &other) const;
};

//Deduction guides
template <typename T, typename... PropertyValueTs>
annotated_arg(T, PropertyValueTs... values) ->
    annotated_arg<T, decltype(properties{PropertyValueTs...})>;

template <typename T, typename PropertiesA, typename PropertiesB>
annotated_arg(annotated_arg<T, PropertiesA>, PropertiesB>) ->
    annotated_arg<T, /* a type that combines the properties of PropertiesA and PropertiesB */>;

// Overloaded operators for raw type
template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() + std::declval<T2>())>
R operator+(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() - std::declval<T2>())>
R operator-(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() * std::declval<T2>())>
R operator*(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() / std::declval<T2>())>
R operator/(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() % std::declval<T2>())>
R operator%(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() & std::declval<T2>())>
R operator&(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() | std::declval<T2>())>
R operator|(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() ^ std::declval<T2>())>
R operator^(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() >> std::declval<T2>())>
R operator>>(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() << std::declval<T2>())>
R operator<<(const annotated_arg<T, PropertyList> &a, const T2 &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() + std::declval<T2>())>
R operator+(const T &a, const annotated_arg<T2, PropertyList2> &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() - std::declval<T2>())>
R operator-(const T &a, const annotated_arg<T2, PropertyList2> &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() * std::declval<T2>())>
R operator*(const T &a, const annotated_arg<T2, PropertyList2> &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() / std::declval<T2>())>
R operator/(const T &a, const annotated_arg<T2, PropertyList2> &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() % std::declval<T2>())>
R operator%(const T &a, const annotated_arg<T2, PropertyList2> &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() & std::declval<T2>())>
R operator&(const T &a, const annotated_arg<T2, PropertyList2> &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() | std::declval<T2>())>
R operator|(const T &a, const annotated_arg<T2, PropertyList2> &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() ^ std::declval<T2>())>
R operator^(const T &a, const annotated_arg<T2, PropertyList2> &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() >> std::declval<T2>())>
R operator>>(const T &a, const annotated_arg<T2, PropertyList2> &b);

template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() << std::declval<T2>())>
R operator<<(const T &a, const annotated_arg<T2, PropertyList2> &b);

} // namespace sycl::ext::oneapi::experimental

The following table describes the member functions of the annotated_arg class:

Functions Description 
annotated_arg() noexcept;

Not available in device code. Constructs an annotated_arg object which is default initialized.

 
annotated_arg(const T& v_, const PropertyListT &P = properties{}) noexcept;

Not available in device code. Constructs an annotated_arg object from the input object v_.

The new property set PropertyListT contains all properties in P.

 
template<typename... PropertyValueTs>
annotated_arg(const T& v_, const PropertyValueTs... props) noexcept;

Not available in device code. Constructs an annotated_arg object from the input object v_.

The new property set PropertyListT contains all properties listed in props.

 
template <typename T2, typename P> explicit annotated_arg(
  const annotated_arg<T2, P> &ConvertFrom);

Not available in device code. Constructs the annotated_arg object from the ConvertFrom object if the list of properties in PropertyListT is a superset of the list of properties in P.

T2 must be implicitly convertible to T.

 
template <typename T2, typename PropertyListU, typename PropertyListV>
explicit annotated_arg(const annotated_arg<T2, PropertyListU>& v_,
  properties<PropertyListV> P) noexcept;

Not available in device code. Constructs the annotated_arg object from the input object v_.

The new PropertyListT is the union of all properties contained within PropertyListU and PropertyListV. If there are any common properties in the two lists with different values, a compile-time error is triggered. Common properties with the same value (or no value) are allowed.

T2 must be implicitly convertible to T.

 
annotated_arg(const annotated_arg&) = default;

Not available in device code. Compiler generated copy constructor.

 
annotated_arg& operator=(annotated_arg&) = default;

Compiler generated assignment operator.

 
operator T() noexcept;
operator const T() const noexcept;

Implicit conversion to a reference to the underlying type T.

 
T operator->() const noexcept;

Member access of type T if T is a pointer type.

 
/* ... */ operator[](std::ptrdiff_t idx) const noexcept;

Available if the operator[] is valid for objects of type T. This function will call the subscript operator defined for T.

 
template<typename propertyT>
static constexpr bool has_property();

Returns true if PropertyListT contains the property with property key class propertyT. Returns false if it does not.

Available only when propertyT is a property key class.

 
template<typename propertyT>
static constexpr /* unspecified */ get_property();

Returns a copy of the property value contained in PropertyListT. Must produce a compile error if PropertyListT does not contain a property with the propertyT key.

Available only if propertyT is the property key class of a compile-time constant property.

 
template <typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() OP std::declval<T2>())>
R operatorOP(const annotated_arg<T2, PropertyList2> &other) const;

Defines the overloaded operators for types annotated_arg<T, …​> and annotated_arg<T2, …​>. Let operatorOP denote the operator used. The overloaded operator operatorOP utilizes operatorOP(T, T2) and is available only if operatorOP(T, T2) is well formed. The value and result is the same as the result of operatorOP(T, T2) applied to the objects of type T and T2.

operatorOP is any of the following: operator+, operator-, operator*, operator/, operator%, operator&, operator|, operator^, operator>>, and operator<<.

Free functions of the following form are defined for the overloaded operators for types annotated_arg<T, …​> and raw type T2. Let operatorOP denote the operator used. The overloaded operator operatorOP utilizes operatorOP(T, T2) and is available only if operatorOP(T, T2) is well formed. The value and result is the same as the result of operatorOP(T, T2) applied to the underlying objects of type T and T2.

template <typename T, typename PropertyList, typename T2,
          typename R = decltype(std::declval<T>() OP std::declval<T2>())>
R operatorOP(const annotated_arg<T, PropertyList> &a, const T2 &b);
template <typename T, typename T2, typename PropertyList2,
          typename R = decltype(std::declval<T>() OP std::declval<T2>())>
R operatorOP(const T &a, const annotated_arg<T2, PropertyList2> &b);

operatorOP is any of the following: operator+, operator-, operator*, operator/, operator%, operator&, operator\|, operator^, operator>>, and operator<<.

Issues

None.

Revision History

Rev Date Author Changes

1

2022-03-09

Abhishek Tiwari

Initial working draft