Skip to content

[SYCL] Introduce external arbitrary fixed point functions #2062

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Jul 10, 2020

Conversation

vmaksimo
Copy link
Contributor

@vmaksimo vmaksimo commented Jul 7, 2020

This change enables support of functions that operate on arbitrary precision
fixed point numbers. This datatype and its corresponding operations can be
useful on targets that can take advantage of narrower representation such as
FPGAs.

Arbitrary precision fixed point numbers are represented using arbitrary
precision integers (_ExtInt(N) wrapped in ap_int). Additional parameters are:

  • W - total width of the datatype. It's encoded in the width of the OpTypeInt.
  • I - determines the position of the decimal point.
  • S - determines if this is a signed or an unsigned number.

Usage example:
ap_int a;
auto ap_fixed_Sqrt = __spirv_FixedSqrtINTEL<W,rW>(a, S, I, rI);
// auto to be deducted to ap_int

This change enables support of functions that operate on arbitrary precision
fixed point numbers. This datatype and its corresponding operations can be
useful on targets that can take advantage of narrower representation such as
FPGAs.

Arbitrary precision fixed point numbers are represented using arbitrary
precision integers (_ExtInt(N) wrapped in ap_int<N>). Additional parameters are:
* W - total width of the datatype. It's encoded in the width of the OpTypeInt.
* I - determines the position of the decimal point.
* S - determines if this is a signed or an unsigned number.

Usage example:
  ap_int<W> a;
  auto ap_fixed_Sqrt = __spirv_FixedSqrtINTEL<W,rW>(a, S, I, rI);
  // auto to be deducted to ap_int<rW>
@vmaksimo vmaksimo requested review from MrSidims and mlychkov July 7, 2020 18:28
@vmaksimo vmaksimo requested a review from a team as a code owner July 7, 2020 18:28
@vmaksimo vmaksimo requested a review from rbegam July 7, 2020 18:28
@keryell
Copy link
Contributor

keryell commented Jul 8, 2020

That sounds like a good generic extension useful for others. :-)
Perhaps something to add as a KHR extension in SPIR-V?

@AlexeySachkov
Copy link
Contributor

That sounds like a good generic extension useful for others. :-)
Perhaps something to add as a KHR extension in SPIR-V?

@keryell, here you go: #1934. This is not a KHR one, but at least some starting point

@keryell
Copy link
Contributor

keryell commented Jul 9, 2020

Do you have a host (& by side effect a host device) implementation for this too?
Actually all the data-types should be available under sycl::... to be able to do something useful globally and keep the single-source approach.
There might be some alignment constraints to be compatible with C++ on the host and also have a kind of packed representation.

@MrSidims
Copy link
Contributor

MrSidims commented Jul 9, 2020

@keryell there will a wrapper headers around these SPIR-V builtins and this header will include host implementation.

Copy link
Contributor

@rbegam rbegam left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM.

@bader bader merged commit 62be913 into intel:sycl Jul 10, 2020
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

7 participants