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.

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] Move bfloat support from experimental to supported. #6524

Merged
merged 99 commits into from
Nov 28, 2022
Merged
Show file tree
Hide file tree
Changes from 45 commits
Commits
Show all changes
99 commits
Select commit Hold shift + click to select a range
6014cef
[SYCL] Move bfloat support from experimental to supported.
rdeodhar Aug 3, 2022
bdd88e5
Corrections to tests.
rdeodhar Aug 3, 2022
73ed541
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Aug 24, 2022
0fe1884
Moved another file out of experimental space.
rdeodhar Aug 24, 2022
feb9d5f
Responses to review comments.
rdeodhar Aug 25, 2022
129f53f
Removed unneeded sycl::half conversion and updated doc.
rdeodhar Aug 26, 2022
2115f09
Added conversion from sycl::half to bfloat16.
rdeodhar Aug 29, 2022
3c2eb80
Cleanup of documentation.
rdeodhar Aug 31, 2022
74aa175
Hooked up bfloat16 aspect within OpenCL plugin.
rdeodhar Sep 2, 2022
bd05711
Support for bfloat16 aspect, and native or fallback support.
rdeodhar Sep 8, 2022
f8e894c
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 8, 2022
2ad68f6
Formatting changes.
rdeodhar Sep 8, 2022
4b78c03
Formatting changes.
rdeodhar Sep 8, 2022
0fce16d
Update to documentation.
rdeodhar Sep 8, 2022
4bcb383
Deprecate bfloat16 aspect.
rdeodhar Sep 8, 2022
35308f8
Fixes for ESIMD.
rdeodhar Sep 9, 2022
fa045e2
Reinstated to_float and from_float, used by NVidia, updated doc.
rdeodhar Sep 9, 2022
3322d6a
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 12, 2022
b12fd94
Update to doc.
rdeodhar Sep 12, 2022
87b0f09
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 14, 2022
f217eb4
Corrections to headers.
rdeodhar Sep 14, 2022
a908b11
Formatting change.
rdeodhar Sep 14, 2022
aab4c78
bfloat16 class supports all sm_xx devices.
Sep 15, 2022
a2568ba
Merge pull request #1 from JackAKirk/bfloat16-cuda-allarch
rdeodhar Sep 15, 2022
4d7a22b
Changes to keep bfloat math functions experimental for now.
rdeodhar Sep 16, 2022
38e5ad4
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 16, 2022
b9accad
Cleanup of bfloat16_math extension.
rdeodhar Sep 16, 2022
ca7880a
Document updates and minor changes.
rdeodhar Sep 19, 2022
dc3b2b5
Fixes for long lines in doc, a different way to check for NaN.
rdeodhar Sep 19, 2022
c955d36
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 20, 2022
1aa6ad3
Broke long lines into multiple lines.
rdeodhar Sep 20, 2022
ff04ce1
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 21, 2022
802f502
Changed library order on Windows.
rdeodhar Sep 21, 2022
8d7f46a
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 22, 2022
190f2a3
Fix for AOT compilation and correction to new headers.
rdeodhar Sep 22, 2022
84c50f3
Noted AOT limitation in doc.
rdeodhar Sep 23, 2022
df058ba
Adjustment for AOT compilation.
rdeodhar Sep 24, 2022
fed4d1d
Fixes for AOT builds.
rdeodhar Sep 26, 2022
28259d0
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 26, 2022
c11115b
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 26, 2022
6b05a2a
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 27, 2022
a82d73a
Fixes for AOT multiple devices.
rdeodhar Sep 27, 2022
3fc8885
Updated documentation.
rdeodhar Sep 27, 2022
1ec6838
Added back missing Status section in documentation.
rdeodhar Sep 27, 2022
105094b
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 27, 2022
432e775
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Sep 29, 2022
c135643
Added tests, corrected aspect check.
rdeodhar Oct 1, 2022
4eca414
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 1, 2022
8876ac8
Added missing newlines.
rdeodhar Oct 3, 2022
f0f2727
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 3, 2022
17673bf
Corrections to tests and macros, added host code emulation.
rdeodhar Oct 4, 2022
1094b8c
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 4, 2022
8d40228
Small corrections.
rdeodhar Oct 4, 2022
c5a85cf
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 4, 2022
cf8f6e0
Fixes for AOT.
rdeodhar Oct 4, 2022
5e50646
Formatting change.
rdeodhar Oct 4, 2022
45d3e70
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 5, 2022
a7be718
Renamed the bfloat aspects.
rdeodhar Oct 5, 2022
cac1c18
Fixes for generic JIT compilation.
rdeodhar Oct 6, 2022
208c09a
Changes for AOT sycl-targets switch.
rdeodhar Oct 6, 2022
46f406d
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 6, 2022
6830857
Corrected aspects queries.
rdeodhar Oct 6, 2022
46e5278
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 6, 2022
10fc9a3
Change in the way fallback/native libs are selected.
rdeodhar Oct 8, 2022
6195545
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 8, 2022
437e34a
Changed type of string.
rdeodhar Oct 10, 2022
09dc4c5
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 12, 2022
386353e
Replaced bfloat16 aspect with bfloat16_math_functions aspect.
rdeodhar Oct 12, 2022
0f93586
Improved devices check in clang driver.
rdeodhar Oct 13, 2022
48f3cac
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 13, 2022
d33cb10
Enhanced test for improved bfloat16 target detection.
rdeodhar Oct 13, 2022
28992c2
Updated bfloat16 driver test for windows.
rdeodhar Oct 13, 2022
ec28c8b
Use STL for parsing devices.
rdeodhar Oct 13, 2022
b958fc7
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 24, 2022
ec70b20
Allow spir64 target to be JIT even when combined with AOT targets.
rdeodhar Oct 24, 2022
1b86012
Updated documentation.
rdeodhar Oct 24, 2022
3e1e681
Modifications for mixed JIT and AOT compilations, added tests.
rdeodhar Oct 25, 2022
8c633d3
Corrections to comments.
rdeodhar Oct 25, 2022
1a59e03
Update to documentation.
rdeodhar Oct 25, 2022
b2fd6cc
Updated doc.
rdeodhar Oct 25, 2022
fab2e54
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 26, 2022
35b8910
Adjustments to tests.
rdeodhar Oct 27, 2022
a05c872
Test cleanup.
rdeodhar Oct 27, 2022
ac5f603
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Oct 27, 2022
6d45ed1
Adjustments to more tests.
rdeodhar Oct 27, 2022
077d0fe
Change to tests to ensure AOT components are available.
rdeodhar Oct 28, 2022
2ff6a9d
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Nov 7, 2022
d7c80ee
Adjustment to test for new bfloat16 header.
rdeodhar Nov 7, 2022
20d13df
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Nov 8, 2022
cd1d0a2
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Nov 15, 2022
4bf60b9
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Nov 18, 2022
45c32f7
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Nov 21, 2022
5de1bf7
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Nov 22, 2022
6ec2bb9
Changes for indirect accesses.
rdeodhar Nov 22, 2022
49e9cd1
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Nov 22, 2022
2065060
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Nov 23, 2022
e24e57b
Fixed conflicts.
rdeodhar Nov 23, 2022
41098ab
Merge branch 'sycl' of https://github.com/intel/llvm into bfloat16
rdeodhar Nov 25, 2022
37b05f0
Correction to library list.
rdeodhar Nov 25, 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
2 changes: 1 addition & 1 deletion clang/lib/Basic/Targets/NVPTX.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -181,7 +181,7 @@ void NVPTXTargetInfo::getTargetDefines(const LangOptions &Opts,
MacroBuilder &Builder) const {
Builder.defineMacro("__PTX__");
Builder.defineMacro("__NVPTX__");
if (Opts.CUDAIsDevice || Opts.OpenMPIsDevice) {
if (Opts.CUDAIsDevice || Opts.OpenMPIsDevice || Opts.SYCLIsDevice) {
// Set __CUDA_ARCH__ for the GPU specified.
std::string CUDAArchCode = [this] {
switch (GPU) {
Expand Down
75 changes: 74 additions & 1 deletion clang/lib/Driver/Driver.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -5081,6 +5081,62 @@ class OffloadingActionBuilder final {
}
}

// Return whether to use native bfloat16 library.
bool useNativeBfloat(const ToolChain *TC, bool &isAOT) {
isAOT = false;
if (!TC->getTriple().isSPIR())
return false;

const OptTable &Opts = C.getDriver().getOpts();
const char *TargetOpt = nullptr;
const char *DeviceOpt = nullptr;
for (auto *A : Args) {
llvm::Triple *TargetBE = nullptr;

auto GetTripleIt = [&, this](llvm::StringRef Triple) {
llvm::Triple TargetTriple{Triple};
auto TripleIt = llvm::find_if(SYCLTripleList, [&](auto &SYCLTriple) {
return SYCLTriple == TargetTriple;
});
return TripleIt != SYCLTripleList.end() ? &*TripleIt : nullptr;
};

if (A->getOption().matches(options::OPT_fsycl_targets_EQ)) {
// Passing arg: -fsycl-targets=<targets>.
isAOT = true;
Copy link
Contributor

Choose a reason for hiding this comment

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

In this case -fsycl-targets=spir64 will also set isAOT to true. Is this the intention? The spir64 target isn't considered AOT.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I've renamed the bool variable to be more meaningful. It tells whether addition of bfloat libs is needed.

TargetBE = GetTripleIt(A->getValue(0));
if (TargetBE)
TargetOpt = A->getValue(0);
else
continue;
} else if (A->getOption().matches(options::OPT_Xsycl_backend_EQ)) {
// Passing device args: -Xsycl-target-backend=<triple> -opt=val.
TargetBE = GetTripleIt(A->getValue(0));
if (TargetBE)
DeviceOpt = A->getValue(1);
else
continue;
} else if (A->getOption().matches(options::OPT_Xsycl_backend)) {
// Passing device args: -Xsycl-target-backend -opt=val.
TargetBE = &SYCLTripleList.front();
DeviceOpt = A->getValue(0);
} else {
mdtoguchi marked this conversation as resolved.
Show resolved Hide resolved
continue;
};
}
if (TC->getTriple().getSubArch() != llvm::Triple::SPIRSubArch_gen)
return false;

if (TargetOpt && DeviceOpt) {
// Currently we support only single AOT target for bfloat16.
if (!(strstr(TargetOpt, "*") || strstr(TargetOpt, ",")))
return strstr(DeviceOpt, "pvc") || strstr(DeviceOpt, "ats");
else
return false;
}
return false;
}

bool addSYCLDeviceLibs(const ToolChain *TC, ActionList &DeviceLinkObjects,
bool isSpirvAOT, bool isMSVCEnv) {
struct DeviceLibOptInfo {
Expand All @@ -5094,7 +5150,8 @@ class OffloadingActionBuilder final {
// of "internal" libraries cannot be affected via -fno-sycl-device-lib.
llvm::StringMap<bool> devicelib_link_info = {
{"libc", true}, {"libm-fp32", true}, {"libm-fp64", true},
{"libimf-fp32", true}, {"libimf-fp64", true}, {"internal", true}};
{"libimf-fp32", true}, {"libimf-fp64", true}, {"libm-bfloat16", true},
{"internal", true}};
if (Arg *A = Args.getLastArg(options::OPT_fsycl_device_lib_EQ,
options::OPT_fno_sycl_device_lib_EQ)) {
if (A->getValues().size() == 0)
Expand Down Expand Up @@ -5153,6 +5210,10 @@ class OffloadingActionBuilder final {
{"libsycl-fallback-cmath-fp64", "libm-fp64"},
{"libsycl-fallback-imf", "libimf-fp32"},
{"libsycl-fallback-imf-fp64", "libimf-fp64"}};
const SYCLDeviceLibsList sycl_device_bfloat16_fallback_lib = {
{"libsycl-fallback-bfloat16", "libm-bfloat16"}};
const SYCLDeviceLibsList sycl_device_bfloat16_native_lib = {
{"libsycl-native-bfloat16", "libm-bfloat16"}};
// ITT annotation libraries are linked in separately whenever the device
// code instrumentation is enabled.
const SYCLDeviceLibsList sycl_device_annotation_libs = {
Expand Down Expand Up @@ -5202,6 +5263,18 @@ class OffloadingActionBuilder final {
addInputs(sycl_device_wrapper_libs);
if (isSpirvAOT || TC->getTriple().isNVPTX())
addInputs(sycl_device_fallback_libs);

bool isAOT;
bool useNativeBfloatLib = useNativeBfloat(TC, isAOT);
if (isAOT &&
TC->getTriple().getSubArch() != llvm::Triple::SPIRSubArch_fpga) {
Copy link
Contributor

Choose a reason for hiding this comment

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

The fpga check is unnecessary given your explicit spir64_gen check in selectBfloatLibs

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Yes, that was an oversight.

// Add native or fallback bfloat16 library.
if (useNativeBfloatLib)
addInputs(sycl_device_bfloat16_native_lib);
else
addInputs(sycl_device_bfloat16_fallback_lib);
}
mdtoguchi marked this conversation as resolved.
Show resolved Hide resolved

if (Args.hasFlag(options::OPT_fsycl_instrument_device_code,
options::OPT_fno_sycl_instrument_device_code, true))
addInputs(sycl_device_annotation_libs);
Expand Down
26 changes: 26 additions & 0 deletions libdevice/bfloat16_wrapper.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,26 @@
//==--- bfloat16_wrapper.cpp - wrappers for bfloat16 library functions ----==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===---------------------------------------------------------------------===//

#include "device.h"

#ifdef __SPIR__

#include <CL/__spirv/spirv_ops.hpp>
#include <cstdint>

DEVICE_EXTERN_C_INLINE
uint16_t __devicelib_ConvertFToBF16INTEL(const float &x) {
return __spirv_ConvertFToBF16INTEL(x);
}

DEVICE_EXTERN_C_INLINE
float __devicelib_ConvertBF16ToFINTEL(const uint16_t &x) {
return __spirv_ConvertBF16ToFINTEL(x);
}

#endif // __SPIR__
4 changes: 4 additions & 0 deletions libdevice/cmake/modules/SYCLLibdevice.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -101,6 +101,7 @@ set(complex_obj_deps device_complex.h device.h sycl-compiler)
set(cmath_obj_deps device_math.h device.h sycl-compiler)
set(imf_obj_deps device_imf.hpp imf_half.hpp device.h sycl-compiler)
set(itt_obj_deps device_itt.h spirv_vars.h device.h sycl-compiler)
set(bfloat16_obj_deps sycl-compiler)

add_devicelib_obj(libsycl-itt-stubs SRC itt_stubs.cpp DEP ${itt_obj_deps})
add_devicelib_obj(libsycl-itt-compiler-wrappers SRC itt_compiler_wrappers.cpp DEP ${itt_obj_deps})
Expand All @@ -113,6 +114,7 @@ add_devicelib_obj(libsycl-cmath SRC cmath_wrapper.cpp DEP ${cmath_obj_deps})
add_devicelib_obj(libsycl-cmath-fp64 SRC cmath_wrapper_fp64.cpp DEP ${cmath_obj_deps} )
add_devicelib_obj(libsycl-imf SRC imf_wrapper.cpp DEP ${imf_obj_deps})
add_devicelib_obj(libsycl-imf-fp64 SRC imf_wrapper_fp64.cpp DEP ${imf_obj_deps})
add_devicelib_obj(libsycl-bfloat16 SRC bfloat16_wrapper.cpp DEP ${cmath_obj_deps} )
if(WIN32)
add_devicelib_obj(libsycl-msvc-math SRC msvc_math.cpp DEP ${cmath_obj_deps})
endif()
Expand All @@ -123,6 +125,8 @@ add_fallback_devicelib(libsycl-fallback-complex SRC fallback-complex.cpp DEP ${c
add_fallback_devicelib(libsycl-fallback-complex-fp64 SRC fallback-complex-fp64.cpp DEP ${complex_obj_deps} )
add_fallback_devicelib(libsycl-fallback-cmath SRC fallback-cmath.cpp DEP ${cmath_obj_deps})
add_fallback_devicelib(libsycl-fallback-cmath-fp64 SRC fallback-cmath-fp64.cpp DEP ${cmath_obj_deps})
add_fallback_devicelib(libsycl-fallback-bfloat16 SRC fallback-bfloat16.cpp DEP ${bfloat16_obj_deps})
add_fallback_devicelib(libsycl-native-bfloat16 SRC bfloat16_wrapper.cpp DEP ${bfloat16_obj_deps})

file(MAKE_DIRECTORY ${obj_binary_dir}/libdevice)
set(imf_fallback_src_dir ${obj_binary_dir}/libdevice)
Expand Down
46 changes: 46 additions & 0 deletions libdevice/fallback-bfloat16.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,46 @@
//==------- fallback-bfloat16.cpp - bfloat16 conversions in software -------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===---------------------------------------------------------------------===//

#include "device.h"

#ifdef __SPIR__

#include <cstdint>

// To support fallback device libraries on-demand loading, please update the
// DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add
// or remove any item in this file.
// TODO: generate the DeviceLibFuncMap in sycl-post-link.cpp automatically
// during the build based on libdevice to avoid manually sync.

DEVICE_EXTERN_C_INLINE uint16_t
__devicelib_ConvertFToBF16INTEL(const float &a) {
// In case float value is nan - propagate bfloat16's qnan
Copy link
Contributor

Choose a reason for hiding this comment

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

Why is this code commented off?

Copy link
Contributor

Choose a reason for hiding this comment

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

I think it is because std::isnan is not available in device code currently.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Yes

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Regarding necessary changes to llvm-test-suite. They are in the corresponding PR intel/llvm-test-suite#1129.
Cross-testing #5424 and intel/llvm-test-suite#1129 has been successful.

Copy link
Contributor

Choose a reason for hiding this comment

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

Doesn't that mean that the emulation is broken, then? Is there another way to test for NaN?

Copy link
Contributor

@MrSidims MrSidims Sep 19, 2022

Choose a reason for hiding this comment

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

Can we use here and later a classic if (a !=a ) here to test if a value is NaN (leaving std::isnan only for the host code)?

if (__spirv_IsNan(a))
return 0xffc1;
union {
uint32_t intStorage;
float floatValue;
};
floatValue = a;
// Do RNE and truncate
uint32_t roundingBias = ((intStorage >> 16) & 0x1) + 0x00007FFF;
return static_cast<uint16_t>((intStorage + roundingBias) >> 16);
}

DEVICE_EXTERN_C_INLINE float
__devicelib_ConvertBF16ToFINTEL(const uint16_t &a) {
union {
uint32_t intStorage;
float floatValue;
};
intStorage = a << 16;
return floatValue;
}

#endif // __SPIR__
5 changes: 5 additions & 0 deletions llvm/tools/sycl-post-link/SYCLDeviceLibReqMask.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -415,6 +415,10 @@ SYCLDeviceLibFuncMap SDLMap = {
DeviceLibExt::cl_intel_devicelib_imf_fp64},
{"__devicelib_imf_longlong_as_double",
DeviceLibExt::cl_intel_devicelib_imf_fp64},
{"__devicelib_ConvertFToBF16INTEL",
DeviceLibExt::cl_intel_devicelib_bfloat16},
{"__devicelib_ConvertBF16ToFINTEL",
DeviceLibExt::cl_intel_devicelib_bfloat16},
};

// Each fallback device library corresponds to one bit in "require mask" which
Expand All @@ -429,6 +433,7 @@ SYCLDeviceLibFuncMap SDLMap = {
// fallback-cstring: 0x20
// fallback-imf: 0x40
// fallback-imf-fp64: 0x80
// fallback-bfloat16: 0x100
uint32_t getDeviceLibBits(const std::string &FuncName) {
auto DeviceLibFuncIter = SDLMap.find(FuncName);
return ((DeviceLibFuncIter == SDLMap.end())
Expand Down
1 change: 1 addition & 0 deletions llvm/tools/sycl-post-link/SYCLDeviceLibReqMask.h
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,7 @@ enum class DeviceLibExt : std::uint32_t {
cl_intel_devicelib_cstring,
cl_intel_devicelib_imf,
cl_intel_devicelib_imf_fp64,
cl_intel_devicelib_bfloat16,
};

uint32_t getSYCLDeviceLibReqMask(const Module &M);
Expand Down
Loading