Skip to content

[SYCL][CUDA] bfloat16 in oneapi namespace and also supporting CUDA #5393

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

Merged
merged 14 commits into from
Apr 5, 2022
Merged

[SYCL][CUDA] bfloat16 in oneapi namespace and also supporting CUDA #5393

Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
025cf7e
Added bfloat16 support for cuda backend.
JackAKirk Jan 25, 2022
66b4e33
deleted intel namespace bfloat16.
JackAKirk Jan 25, 2022
2d04406
Format.
JackAKirk Jan 25, 2022
9418f74
Changed extension macro name.
JackAKirk Jan 25, 2022
65fddfa
Merge branch 'sycl' into bf16-cvt-ext
JackAKirk Feb 17, 2022
4d99f3f
fixed test.
JackAKirk Feb 17, 2022
3982001
Used neg ptx7.0 builtin for unary minus
JackAKirk Mar 4, 2022
8d2d11f
Replaced SYCL_EXT_INTEL_BF16_CONVERSION.asciidoc with SYCL_EXT_ONEAPI…
JackAKirk Mar 7, 2022
d8bc53f
Merge branch 'sycl' into bf16-cvt-ext
JackAKirk Mar 8, 2022
2f9b7d7
Merge branch 'sycl' into bf16-cvt-ext
JackAKirk Mar 15, 2022
8a29c44
Renamed extension to cover all bfloat16 funct.
JackAKirk Mar 15, 2022
f1fba08
Updated macro name
JackAKirk Mar 16, 2022
461ddb8
Removed old extension doc
JackAKirk Mar 16, 2022
e433fbc
typo
JackAKirk Mar 31, 2022
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
3 changes: 2 additions & 1 deletion sycl/include/CL/sycl/feature_test.hpp.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,8 @@ namespace sycl {
#define SYCL_EXT_ONEAPI_USE_PINNED_HOST_MEMORY_PROPERTY 1
#define SYCL_EXT_ONEAPI_SRGB 1
#define SYCL_EXT_ONEAPI_SUB_GROUP 1
#define SYCL_EXT_INTEL_BF16_CONVERSION 1
#define SYCL_EXT_ONEAPI_BF16_CONVERSION 1
#define SYCL_EXT_INTEL_BITCAST 1
#define SYCL_EXT_INTEL_DATAFLOW_PIPES 1
#ifdef __clang__
#if __has_extension(sycl_extended_atomics)
Expand Down
15 changes: 13 additions & 2 deletions .../sycl/ext/intel/experimental/bfloat16.hpp → ...sycl/ext/oneapi/experimental/bfloat16.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,7 @@
__SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
namespace ext {
namespace intel {
namespace oneapi {
namespace experimental {

class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 {
Copy link
Contributor

Choose a reason for hiding this comment

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

This should be wrapped with SYCL_DEVICE_ONLY defined, for it to be built successfully by other compilers, right? (see #5594)
@AlexeySotkin, are you planning for adding this as a separate PR?

Expand All @@ -29,15 +29,26 @@ class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 {
// Explicit conversion functions
static storage_t from_float(const float &a) {
#if defined(__SYCL_DEVICE_ONLY__)
#if defined(__NVPTX__)
return __nvvm_f2bf16_rn(a);
#else
return __spirv_ConvertFToBF16INTEL(a);
#endif
#else
throw exception{errc::feature_not_supported,
"Bfloat16 conversion is not supported on host device"};
#endif
}
static float to_float(const storage_t &a) {
#if defined(__SYCL_DEVICE_ONLY__)
#if defined(__NVPTX__)
unsigned int y = a;
y = y << 16;
float *res = reinterpret_cast<float *>(&y);
return *res;
#else
return __spirv_ConvertBF16ToFINTEL(a);
#endif
#else
throw exception{errc::feature_not_supported,
"Bfloat16 conversion is not supported on host device"};
Expand Down Expand Up @@ -143,7 +154,7 @@ class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 {
};

} // namespace experimental
} // namespace intel
} // namespace oneapi
} // namespace ext

} // namespace sycl
Expand Down
4 changes: 2 additions & 2 deletions sycl/test/extensions/bfloat16.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,10 +2,10 @@

// UNSUPPORTED: cuda || hip_amd

#include <sycl/ext/intel/experimental/bfloat16.hpp>
#include <sycl/ext/oneapi/experimental/bfloat16.hpp>
#include <sycl/sycl.hpp>

using sycl::ext::intel::experimental::bfloat16;
using sycl::ext::oneapi::experimental::bfloat16;

SYCL_EXTERNAL uint16_t some_bf16_intrinsic(uint16_t x, uint16_t y);
SYCL_EXTERNAL void foo(long x, sycl::half y);
Expand Down