Skip to content
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.

Sign up for GitHub

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

[SYCL] Add vec<bfloat16> support to math builtins #14002

Merged
merged 16 commits into from
Jun 7, 2024
Merged

[SYCL] Add vec<bfloat16> support to math builtins #14002

Show file tree
Hide file tree
Changes from 7 commits
Commits
Show all changes
16 commits
Select commit Hold shift + click to select a range
0aa7a9a
Add copy constructor
uditagarwal97 Apr 18, 2024
f2a1dc2
Merge branch 'sycl' of https://github.com/uditagarwal97/llvm into sycl
uditagarwal97 May 28, 2024
361eea7
Merge branch 'sycl' of https://github.com/uditagarwal97/llvm into sycl
uditagarwal97 May 31, 2024
d90deb9
Add basic support for vec<bfloat16> to math builtins.
uditagarwal97 Jun 1, 2024
e3d3de8
Update doc; Fix formatting
uditagarwal97 Jun 1, 2024
19b52ee
Add device code test for vec<bfloat16> math builtins.
uditagarwal97 Jun 1, 2024
72cb653
Fix formatting
uditagarwal97 Jun 1, 2024
0e0a740
Update sycl_ext_oneapi_bfloat16_math_functions.asciidoc
uditagarwal97 Jun 4, 2024
7df752a
Update sycl_ext_oneapi_bfloat16_math_functions.asciidoc
uditagarwal97 Jun 4, 2024
6b43d13
Implement math builtins overloads for swizzles
uditagarwal97 Jun 5, 2024
b59e6d5
Update device code test
uditagarwal97 Jun 5, 2024
8e0379f
Check that vec or swizzle are of type bfloat16
uditagarwal97 Jun 5, 2024
3a98092
Remove redundant func
uditagarwal97 Jun 5, 2024
5f47ea7
Update sycl_ext_oneapi_bfloat16_math_functions.asciidoc
uditagarwal97 Jun 5, 2024
21b6ba8
Remove redundant statement from the extension doc
uditagarwal97 Jun 6, 2024
9b478af
Fix test failures
uditagarwal97 Jun 6, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
46 changes: 22 additions & 24 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_bfloat16_math_functions.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -105,7 +105,8 @@ then it supports the `bfloat16` math functions described in the next section.

The following functions are only available when `T` is `bfloat16` or
`sycl::marray<bfloat16, {N}>`, where `{N}` means any positive value of
`size_t` type.
`size_t` type, or `sycl::vec<bfloat16, {N}>`, where `{N}` is a valid
`sycl::vec` size.

==== isnan

Expand All @@ -116,12 +117,15 @@ bool isnan(bfloat16 x);

template <size_t N>
sycl::marray<bool, N> isnan(sycl::marray<bfloat16, N> x);

template <size_t N>
sycl::vec<bool, N> isnan(sycl::vec<bfloat16, N> x);
} // namespace sycl::ext::oneapi::experimental
```

===== Description

Returns true if x is NAN value, otherwise returns false.
Returns true if x is or contains a NAN value, otherwise returns false.

==== fma

Expand Down Expand Up @@ -187,7 +191,7 @@ T fabs(T x);

===== Description

Compute absolute value of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute absolute value of a `bfloat16` value.

==== ceil

Expand All @@ -200,7 +204,7 @@ T ceil(T x);

===== Description

Returns `x` rounded to an integral value using the round to positive infinity rounding mode
Returns `x` rounded to an integral value using the round to positive infinity rounding mode.

==== floor

Expand All @@ -214,7 +218,7 @@ T floor(T x);
===== Description

Returns `x` rounded to an integral value using the round to negative infinity rounding mode
for a `bfloat16` value or `sycl::marray<bfloat16, N>`.
for a `bfloat16` value.

==== cos

Expand All @@ -227,7 +231,7 @@ T cos(T x);

===== Description

Compute cosine of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute cosine of a `bfloat16` value.

==== sin

Expand All @@ -240,7 +244,7 @@ T sin(T x);

===== Description

Compute sine of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute sine of a `bfloat16` value.


==== exp
Expand All @@ -254,7 +258,7 @@ T exp(T x);

===== Description

Compute the base-e exponential of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute the base-e exponential of a `bfloat16` value.

==== exp2

Expand All @@ -267,7 +271,7 @@ T exp2(T x);

===== Description

Compute the base-2 exponential of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute the base-2 exponential of a `bfloat16` value.

==== exp10

Expand All @@ -280,7 +284,7 @@ T exp10(T x);

===== Description

Compute the base-10 exponential of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute the base-10 exponential of a `bfloat16` value.

==== log

Expand All @@ -293,7 +297,7 @@ T log(T x);

===== Description

Compute natural logarithm of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute natural logarithm of a `bfloat16` value.

==== log2

Expand All @@ -306,7 +310,7 @@ T log2(T x);

===== Description

Compute base-2 logarithm of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute base-2 logarithm of a `bfloat16` value.

==== log10

Expand All @@ -319,7 +323,7 @@ T log10(T x);

===== Description

Compute base-10 logarithm of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute base-10 logarithm of a `bfloat16` value.


==== rint
Expand All @@ -334,7 +338,7 @@ T rint(T x);
===== Description

Returns `x` rounded to an integral value using the round to nearest even rounding mode
for a `bfloat16` value or `sycl::marray<bfloat16, N>`.
for a `bfloat16` value.

==== sqrt

Expand All @@ -347,7 +351,7 @@ T sqrt(T x);

===== Description

Compute square root of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute square root of a `bfloat16` value.

==== rsqrt

Expand All @@ -360,7 +364,7 @@ T rsqrt(T x);

===== Description

Compute inverse square root of a `bfloat16` value or `sycl::marray<bfloat16, N>`.
Compute inverse square root of a `bfloat16` value.

==== trunc

Expand All @@ -374,15 +378,9 @@ T trunc(T x);
===== Description

Returns `x` rounded to an integral value using the round to zero rounding mode
for a `bfloat16` value or `sycl::marray<bfloat16, N>`.
for a `bfloat16` value.

== Issues

1. The CUDA backend does not have a use case that would necessitate support
of the `vec` class in bfloat16 math functions, and `marray` would always be
preferred over `vec` if `vec` support were to be added in the CUDA backend.
For portability reasons, support for the `vec` class can be easily added if
other backends require it.

2. We should decide on a roadmap to extend support of `bfloat16` to other
1. We should decide on a roadmap to extend support of `bfloat16` to other
SYCL 2020 math functions.
83 changes: 83 additions & 0 deletions sycl/include/sycl/ext/oneapi/experimental/bfloat16_math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,14 @@ template <size_t N> sycl::marray<bool, N> isnan(sycl::marray<bfloat16, N> x) {
return res;
}

template <int N> sycl::vec<bool, N> isnan(sycl::vec<bfloat16, N> x) {
sycl::vec<bool, N> res;
for (size_t i = 0; i < N; i++) {
res[i] = isnan(x[i]);
}
return res;
}
gmlueck marked this conversation as resolved.
Show resolved Hide resolved

template <typename T>
std::enable_if_t<std::is_same_v<T, bfloat16>, T> fabs(T x) {
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__) && \
Expand Down Expand Up @@ -89,6 +97,14 @@ sycl::marray<bfloat16, N> fabs(sycl::marray<bfloat16, N> x) {
return res;
}

template <int N> sycl::vec<bfloat16, N> fabs(sycl::vec<bfloat16, N> x) {
sycl::vec<bfloat16, N> res;
for (size_t i = 0; i < N; i++) {
res[i] = fabs(x[i]);
}
return res;
}

template <typename T>
std::enable_if_t<std::is_same_v<T, bfloat16>, T> fmin(T x, T y) {
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__) && \
Expand Down Expand Up @@ -146,6 +162,16 @@ sycl::marray<bfloat16, N> fmin(sycl::marray<bfloat16, N> x,
return res;
}

template <int N>
sycl::vec<bfloat16, N> fmin(sycl::vec<bfloat16, N> x,
sycl::vec<bfloat16, N> y) {
sycl::vec<bfloat16, N> res;
for (size_t i = 0; i < N; i++) {
res[i] = fmin(x[i], y[i]);
}
return res;
}

template <typename T>
std::enable_if_t<std::is_same_v<T, bfloat16>, T> fmax(T x, T y) {
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__) && \
Expand Down Expand Up @@ -202,6 +228,16 @@ sycl::marray<bfloat16, N> fmax(sycl::marray<bfloat16, N> x,
return res;
}

template <int N>
sycl::vec<bfloat16, N> fmax(sycl::vec<bfloat16, N> x,
sycl::vec<bfloat16, N> y) {
sycl::vec<bfloat16, N> res;
for (size_t i = 0; i < N; i++) {
res[i] = fmax(x[i], y[i]);
}
return res;
}

template <typename T>
std::enable_if_t<std::is_same_v<T, bfloat16>, T> fma(T x, T y, T z) {
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__) && \
Expand Down Expand Up @@ -248,6 +284,16 @@ sycl::marray<bfloat16, N> fma(sycl::marray<bfloat16, N> x,
return res;
}

template <int N>
sycl::vec<bfloat16, N> fma(sycl::vec<bfloat16, N> x, sycl::vec<bfloat16, N> y,
sycl::vec<bfloat16, N> z) {
sycl::vec<bfloat16, N> res;
for (size_t i = 0; i < N; i++) {
res[i] = fma(x[i], y[i], z[i]);
}
return res;
}

#define BFLOAT16_MATH_FP32_WRAPPERS(op) \
template <typename T> \
std::enable_if_t<std::is_same<T, bfloat16>::value, T> op(T x) { \
Expand All @@ -264,37 +310,74 @@ sycl::marray<bfloat16, N> fma(sycl::marray<bfloat16, N> x,
return res; \
}

#define BFLOAT16_MATH_FP32_WRAPPERS_VEC(op) \
template <int N> sycl::vec<bfloat16, N> op(sycl::vec<bfloat16, N> x) { \
sycl::vec<bfloat16, N> res; \
for (size_t i = 0; i < N; i++) { \
res[i] = op(x[i]); \
} \
return res; \
}

BFLOAT16_MATH_FP32_WRAPPERS(ceil)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(ceil)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(ceil)

BFLOAT16_MATH_FP32_WRAPPERS(cos)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(cos)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(cos)

BFLOAT16_MATH_FP32_WRAPPERS(exp)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(exp)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(exp)

BFLOAT16_MATH_FP32_WRAPPERS(exp10)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(exp10)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(exp10)

BFLOAT16_MATH_FP32_WRAPPERS(exp2)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(exp2)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(exp2)

BFLOAT16_MATH_FP32_WRAPPERS(floor)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(floor)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(floor)

BFLOAT16_MATH_FP32_WRAPPERS(log)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(log)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(log)

BFLOAT16_MATH_FP32_WRAPPERS(log2)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(log2)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(log2)

BFLOAT16_MATH_FP32_WRAPPERS(log10)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(log10)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(log10)

BFLOAT16_MATH_FP32_WRAPPERS(rint)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(rint)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(rint)

BFLOAT16_MATH_FP32_WRAPPERS(rsqrt)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(rsqrt)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(rsqrt)

BFLOAT16_MATH_FP32_WRAPPERS(sin)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(sin)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(sin)

BFLOAT16_MATH_FP32_WRAPPERS(sqrt)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(sqrt)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(sqrt)

BFLOAT16_MATH_FP32_WRAPPERS(trunc)
BFLOAT16_MATH_FP32_WRAPPERS_MARRAY(trunc)
BFLOAT16_MATH_FP32_WRAPPERS_VEC(trunc)

#undef BFLOAT16_MATH_FP32_WRAPPERS
#undef BFLOAT16_MATH_FP32_WRAPPERS_MARRAY
#undef BFLOAT16_MATH_FP32_WRAPPERS_VEC
} // namespace ext::oneapi::experimental
} // namespace _V1
} // namespace sycl
Loading
Loading