Skip to content

Commit 51568f7

Browse files
elizabethandrewselizabethandrews
authored
[SYCL][OpenMP] Correctly handle __bf16 intrinsics using deferred diagnostics (#8116)
Commits 1 and 2 are cherry-picks from llvm.org - https://reviews.llvm.org/D141375 and https://reviews.llvm.org/rGc10615e4a94fc8dec65a48a6eb8f7efccc3fb1fc. Commit 3 reverts header changes made in intel/llvm.
1 parent ad86982 commit 51568f7

File tree

10 files changed

+39
-18
lines changed

10 files changed

+39
-18
lines changed

‎clang/lib/AST/ASTContext.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2146,6 +2146,11 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const {
21462146
if (Target->hasBFloat16Type()) {
21472147
Width = Target->getBFloat16Width();
21482148
Align = Target->getBFloat16Align();
2149+
} else if ((getLangOpts().SYCLIsDevice ||
2150+
(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice)) &&
2151+
AuxTarget->hasBFloat16Type()) {
2152+
Width = AuxTarget->getBFloat16Width();
2153+
Align = AuxTarget->getBFloat16Align();
21492154
}
21502155
break;
21512156
case BuiltinType::Float16:

‎clang/lib/AST/ItaniumMangle.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3050,7 +3050,11 @@ void CXXNameMangler::mangleType(const BuiltinType *T) {
30503050
break;
30513051
}
30523052
case BuiltinType::BFloat16: {
3053-
const TargetInfo *TI = &getASTContext().getTargetInfo();
3053+
const TargetInfo *TI = ((getASTContext().getLangOpts().OpenMP &&
3054+
getASTContext().getLangOpts().OpenMPIsDevice) ||
3055+
getASTContext().getLangOpts().SYCLIsDevice)
3056+
? getASTContext().getAuxTargetInfo()
3057+
: &getASTContext().getTargetInfo();
30543058
Out << TI->getBFloat16Mangling();
30553059
break;
30563060
}

‎clang/lib/Headers/avxintrin.h

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -45,11 +45,9 @@ typedef _Float16 __v16hf __attribute__((__vector_size__(32), __aligned__(32)));
4545
typedef _Float16 __m256h __attribute__((__vector_size__(32), __aligned__(32)));
4646
typedef _Float16 __m256h_u __attribute__((__vector_size__(32), __aligned__(1)));
4747

48-
#ifndef __SYCL_DEVICE_ONLY__
4948
typedef __bf16 __v16bf __attribute__((__vector_size__(32), __aligned__(32)));
5049
typedef __bf16 __m256bh __attribute__((__vector_size__(32), __aligned__(32)));
5150
#endif
52-
#endif
5351

5452
/* Define the default attributes for the functions in this file. */
5553
#define __DEFAULT_FN_ATTRS __attribute__((__always_inline__, __nodebug__, __target__("avx"), __min_vector_width__(256)))

‎clang/lib/Headers/avxneconvertintrin.h

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,6 @@
2525
__attribute__((__always_inline__, __nodebug__, __target__("avxneconvert"), \
2626
__min_vector_width__(256)))
2727

28-
#ifndef __SYCL_DEVICE_ONLY__
2928
/// Convert scalar BF16 (16-bit) floating-point element
3029
/// stored at memory locations starting at location \a __A to a
3130
/// single-precision (32-bit) floating-point, broadcast it to packed
@@ -91,7 +90,6 @@ static __inline__ __m256 __DEFAULT_FN_ATTRS256
9190
_mm256_bcstnebf16_ps(const void *__A) {
9291
return (__m256)__builtin_ia32_vbcstnebf162ps256((const __bf16 *)__A);
9392
}
94-
#endif
9593

9694
/// Convert scalar half-precision (16-bit) floating-point element
9795
/// stored at memory locations starting at location \a __A to a
@@ -159,7 +157,6 @@ _mm256_bcstnesh_ps(const void *__A) {
159157
return (__m256)__builtin_ia32_vbcstnesh2ps256((const _Float16 *)__A);
160158
}
161159

162-
#ifndef __SYCL_DEVICE_ONLY__
163160
/// Convert packed BF16 (16-bit) floating-point even-indexed elements
164161
/// stored at memory locations starting at location \a __A to packed
165162
/// single-precision (32-bit) floating-point elements, and store the results in
@@ -225,7 +222,6 @@ static __inline__ __m256 __DEFAULT_FN_ATTRS256
225222
_mm256_cvtneebf16_ps(const __m256bh *__A) {
226223
return (__m256)__builtin_ia32_vcvtneebf162ps256((const __v16bf *)__A);
227224
}
228-
#endif
229225

230226
/// Convert packed half-precision (16-bit) floating-point even-indexed elements
231227
/// stored at memory locations starting at location \a __A to packed
@@ -293,7 +289,6 @@ _mm256_cvtneeph_ps(const __m256h *__A) {
293289
return (__m256)__builtin_ia32_vcvtneeph2ps256((const __v16hf *)__A);
294290
}
295291

296-
#ifndef __SYCL_DEVICE_ONLY__
297292
/// Convert packed BF16 (16-bit) floating-point odd-indexed elements
298293
/// stored at memory locations starting at location \a __A to packed
299294
/// single-precision (32-bit) floating-point elements, and store the results in
@@ -359,7 +354,6 @@ static __inline__ __m256 __DEFAULT_FN_ATTRS256
359354
_mm256_cvtneobf16_ps(const __m256bh *__A) {
360355
return (__m256)__builtin_ia32_vcvtneobf162ps256((const __v16bf *)__A);
361356
}
362-
#endif
363357

364358
/// Convert packed half-precision (16-bit) floating-point odd-indexed elements
365359
/// stored at memory locations starting at location \a __A to packed
@@ -427,7 +421,6 @@ _mm256_cvtneoph_ps(const __m256h *__A) {
427421
return (__m256)__builtin_ia32_vcvtneoph2ps256((const __v16hf *)__A);
428422
}
429423

430-
#ifndef __SYCL_DEVICE_ONLY__
431424
/// Convert packed single-precision (32-bit) floating-point elements in \a __A
432425
/// to packed BF16 (16-bit) floating-point elements, and store the results in \a
433426
/// dst.
@@ -483,7 +476,6 @@ static __inline__ __m128bh __DEFAULT_FN_ATTRS256
483476
_mm256_cvtneps_avx_pbh(__m256 __A) {
484477
return (__m128bh)__builtin_ia32_vcvtneps2bf16256((__v8sf)__A);
485478
}
486-
#endif
487479

488480
#undef __DEFAULT_FN_ATTRS128
489481
#undef __DEFAULT_FN_ATTRS256

‎clang/lib/Headers/emmintrin.h

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -44,11 +44,9 @@ typedef _Float16 __v8hf __attribute__((__vector_size__(16), __aligned__(16)));
4444
typedef _Float16 __m128h __attribute__((__vector_size__(16), __aligned__(16)));
4545
typedef _Float16 __m128h_u __attribute__((__vector_size__(16), __aligned__(1)));
4646

47-
#ifndef __SYCL_DEVICE_ONLY__
4847
typedef __bf16 __v8bf __attribute__((__vector_size__(16), __aligned__(16)));
4948
typedef __bf16 __m128bh __attribute__((__vector_size__(16), __aligned__(16)));
5049
#endif
51-
#endif
5250

5351
/* Define the default attributes for the functions in this file. */
5452
#define __DEFAULT_FN_ATTRS \

‎clang/lib/Headers/immintrin.h

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -229,7 +229,6 @@
229229
#include <avx512vlfp16intrin.h>
230230
#endif
231231

232-
#ifndef __SYCL_DEVICE_ONLY__
233232
#if !(defined(_MSC_VER) || defined(__SCE__)) || __has_feature(modules) || \
234233
defined(__AVX512BF16__)
235234
#include <avx512bf16intrin.h>
@@ -239,7 +238,6 @@
239238
(defined(__AVX512VL__) && defined(__AVX512BF16__))
240239
#include <avx512vlbf16intrin.h>
241240
#endif
242-
#endif
243241

244242
#if !(defined(_MSC_VER) || defined(__SCE__)) || __has_feature(modules) || \
245243
defined(__PKU__)

‎clang/lib/Sema/Sema.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2073,6 +2073,8 @@ void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) {
20732073
(Ty->isIbm128Type() && !Context.getTargetInfo().hasIbm128Type()) ||
20742074
(Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
20752075
!Context.getTargetInfo().hasInt128Type()) ||
2076+
(Ty->isBFloat16Type() && !Context.getTargetInfo().hasBFloat16Type() &&
2077+
!LangOpts.CUDAIsDevice) ||
20762078
LongDoubleMismatched) {
20772079
PartialDiagnostic PD = PDiag(diag::err_target_unsupported_type);
20782080
if (D)

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -434,6 +434,7 @@ static void checkSYCLType(Sema &S, QualType Ty, SourceRange Loc,
434434
if (Ty->isSpecificBuiltinType(BuiltinType::Int128) ||
435435
Ty->isSpecificBuiltinType(BuiltinType::UInt128) ||
436436
Ty->isSpecificBuiltinType(BuiltinType::LongDouble) ||
437+
Ty->isSpecificBuiltinType(BuiltinType::BFloat16) ||
437438
(Ty->isSpecificBuiltinType(BuiltinType::Float128) &&
438439
!S.Context.getTargetInfo().hasFloat128Type())) {
439440
S.SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported)

‎clang/lib/Sema/SemaType.cpp

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1527,9 +1527,10 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
15271527
break;
15281528
case DeclSpec::TST_half: Result = Context.HalfTy; break;
15291529
case DeclSpec::TST_BFloat16:
1530-
if (!S.Context.getTargetInfo().hasBFloat16Type())
1531-
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
1532-
<< "__bf16";
1530+
if (!S.Context.getTargetInfo().hasBFloat16Type() &&
1531+
!(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice) &&
1532+
!S.getLangOpts().SYCLIsDevice)
1533+
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported) << "__bf16";
15331534
Result = Context.BFloat16Ty;
15341535
break;
15351536
case DeclSpec::TST_float: Result = Context.FloatTy; break;

‎clang/test/SemaSYCL/bf16.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
// RUN: %clang_cc1 -triple spir64 -aux-triple x86_64-unknown-linux-gnu -fsycl-is-device -verify -fsyntax-only %s
2+
3+
template <typename Name, typename Func>
4+
__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
5+
kernelFunc(); // expected-note {{called by 'kernel}}
6+
}
7+
8+
void host_ok(void) {
9+
__bf16 A;
10+
}
11+
12+
int main()
13+
{ host_ok();
14+
__bf16 var; // expected-note {{'var' defined here}}
15+
kernel<class variables>([=]() {
16+
(void)var; // expected-error {{'var' requires 16 bit size '__bf16' type support, but target 'spir64' does not support it}}
17+
int B = sizeof(__bf16);
18+
});
19+
20+
return 0;
21+
}
22+

0 commit comments

Comments
 (0)