diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 7fbad9cea70c5..44697bcf0f58a 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -261,7 +261,9 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, const IdentifierInfo *Id = FDecl->getIdentifier(); if ((getEmissionReason(FDecl) == Sema::DeviceDiagnosticReason::Sycl) && Id && !Id->getName().startswith("__spirv_") && - !Id->getName().startswith("__sycl_")) { + !Id->getName().startswith("__sycl_") && + !Id->getName().startswith("__devicelib_ConvertBF16ToFINTEL") && + !Id->getName().startswith("__devicelib_ConvertFToBF16INTEL")) { SYCLDiagIfDeviceCode( *Locs.begin(), diag::err_sycl_device_function_is_called_from_esimd, Sema::DeviceDiagnosticReason::Esimd); diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index f96ee7b4fed85..ee3f9ecf3c5cf 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -54,7 +54,7 @@ static const char *LegalSYCLFunctions[] = { "^sycl::_V1::ext::oneapi::sub_group::.+", "^sycl::_V1::ext::oneapi::experimental::spec_constant<.+>::.+", "^sycl::_V1::ext::oneapi::experimental::this_sub_group", - "^sycl::_V1::ext::oneapi::experimental::bfloat16::.+", + "^sycl::_V1::ext::oneapi::bfloat16::.+", "^sycl::_V1::ext::oneapi::experimental::if_architecture_is"}; static const char *LegalSYCLFunctionsInStatelessMode[] = { diff --git a/sycl/test/esimd/fp16_converts.cpp b/sycl/test/esimd/fp16_converts.cpp new file mode 100644 index 0000000000000..9177702ea0d64 --- /dev/null +++ b/sycl/test/esimd/fp16_converts.cpp @@ -0,0 +1,50 @@ +// The test verifies support of bfloat16 <-> float conversions + +// Checks host+device compilation +// RUN: %clangxx -fsycl -fsyntax-only %s + +// Checks that lowerESIMD pass builds proper vc-intrinsics +// RUN: %clangxx -O2 -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o %t +// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table +// RUN: FileCheck %s -input-file=%t_esimd_0.ll + +#include +#include + +using namespace sycl::ext::intel::esimd; + +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void bf16_vector(); +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void bf16_scalar(); + +using bfloat16 = sycl::ext::oneapi::bfloat16; + +class EsimdFunctor { +public: + void operator()() __attribute__((sycl_explicit_simd)) { + bf16_vector(); + bf16_scalar(); + } +}; + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void bf16_vector() { + simd F32 = 0; + simd BF16 = F32; + // CHECK: call <8 x half> @llvm.genx.bf.cvt.v8f16.v8f32(<8 x float> {{[^)]+}}) + simd F32_conv = BF16; + // CHECK: call <8 x float> @llvm.genx.bf.cvt.v8f32.v8f16(<8 x half> {{[^)]+}}) +} + +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void bf16_scalar() { + // Note that this is the compilation test only. It checks that IR is correct. + // The actual support in GPU RT is on the way though. + float F32_scalar = 1; + bfloat16 BF16_scalar = F32_scalar; + // CHECK: call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{[^)]+}}) + float F32_scalar_conv = BF16_scalar; + // CHECK: call spir_func float @__devicelib_ConvertBF16ToFINTEL(i16 {{[^)]+}}) +}