Skip to content

Commit 890471e

Browse files
vmaksimovmaksimo
committed
Merge from 'main' to 'sycl-web' (41 commits)
 CONFLICT (content): Merge conflict in libclc/cmake/modules/AddLibclc.cmake
2 parents b9a6aac + abe93d9 commit 890471e

File tree

196 files changed

+13682
-6982
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

196 files changed

+13682
-6982
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -710,6 +710,7 @@ TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8, "V8hV16iV16iIsV8hIbI
710710
TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8, "V8hV16iV16iIsV8hIbIb", "nc", "gfx1250-insts,wavefrontsize32")
711711
TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8, "V8hV16iV16iIsV8hIbIb", "nc", "gfx1250-insts,wavefrontsize32")
712712
TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8, "V8hV16iV16iIsV8hIbIb", "nc", "gfx1250-insts,wavefrontsize32")
713+
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4, "V8fIiV16iIiV16iIsV8f", "nc", "gfx1250-insts,wavefrontsize32")
713714
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8, "V8fV16iV16iIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
714715
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8, "V8fV16iV16iIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
715716
TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8, "V8fV16iV16iIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")

‎clang/include/clang/Driver/Options.td

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7551,6 +7551,8 @@ def fintrinsic_modules_path : Separate<["-"], "fintrinsic-modules-path">, Group
75517551
HelpText<"Specify where to find the compiled intrinsic modules">,
75527552
DocBrief<[{This option specifies the location of pre-compiled intrinsic modules,
75537553
if they are not in the default location expected by the compiler.}]>;
7554+
def fintrinsic_modules_path_EQ : Joined<["-"], "fintrinsic-modules-path=">,
7555+
Group<f_Group>, Alias<fintrinsic_modules_path>;
75547556

75557557
defm backslash : OptInFC1FFlag<"backslash", "Specify that backslash in string introduces an escape character">;
75567558
defm xor_operator : OptInFC1FFlag<"xor-operator", "Enable .XOR. as a synonym of .NEQV.">;

‎clang/lib/AST/ByteCode/ByteCodeEmitter.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -62,7 +62,7 @@ void ByteCodeEmitter::compileFunc(const FunctionDecl *FuncDecl,
6262
(Func->hasThisPointer() && !Func->isThisPointerExplicit());
6363
for (auto ParamOffset : llvm::drop_begin(Func->ParamOffsets, Drop)) {
6464
const ParmVarDecl *PD = FuncDecl->parameters()[ParamIndex];
65-
std::optional<PrimType> T = Ctx.classify(PD->getType());
65+
OptPrimType T = Ctx.classify(PD->getType());
6666
this->Params.insert({PD, {ParamOffset, T != std::nullopt}});
6767
++ParamIndex;
6868
}

‎clang/lib/AST/ByteCode/Compiler.cpp

Lines changed: 59 additions & 60 deletions
Large diffs are not rendered by default.

‎clang/lib/AST/ByteCode/Compiler.h

Lines changed: 4 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -254,12 +254,8 @@ class Compiler : public ConstStmtVisitor<Compiler<Emitter>, bool>,
254254
/// If the function does not exist yet, it is compiled.
255255
const Function *getFunction(const FunctionDecl *FD);
256256

257-
std::optional<PrimType> classify(const Expr *E) const {
258-
return Ctx.classify(E);
259-
}
260-
std::optional<PrimType> classify(QualType Ty) const {
261-
return Ctx.classify(Ty);
262-
}
257+
OptPrimType classify(const Expr *E) const { return Ctx.classify(E); }
258+
OptPrimType classify(QualType Ty) const { return Ctx.classify(Ty); }
263259

264260
/// Classifies a known primitive type.
265261
PrimType classifyPrim(QualType Ty) const {
@@ -306,7 +302,7 @@ class Compiler : public ConstStmtVisitor<Compiler<Emitter>, bool>,
306302
bool visitInitList(ArrayRef<const Expr *> Inits, const Expr *ArrayFiller,
307303
const Expr *E);
308304
bool visitArrayElemInit(unsigned ElemIndex, const Expr *Init,
309-
std::optional<PrimType> InitT);
305+
OptPrimType InitT);
310306
bool visitCallArgs(ArrayRef<const Expr *> Args, const FunctionDecl *FuncDecl,
311307
bool Activate);
312308

@@ -435,7 +431,7 @@ class Compiler : public ConstStmtVisitor<Compiler<Emitter>, bool>,
435431
bool InitStackActive = false;
436432

437433
/// Type of the expression returned by the function.
438-
std::optional<PrimType> ReturnType;
434+
OptPrimType ReturnType;
439435

440436
/// Switch case mapping.
441437
CaseMap CaseLabels;

‎clang/lib/AST/ByteCode/Context.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -306,7 +306,7 @@ static PrimType integralTypeToPrimTypeU(unsigned BitWidth) {
306306
llvm_unreachable("Unhandled BitWidth");
307307
}
308308

309-
std::optional<PrimType> Context::classify(QualType T) const {
309+
OptPrimType Context::classify(QualType T) const {
310310

311311
if (const auto *BT = dyn_cast<BuiltinType>(T.getCanonicalType())) {
312312
auto Kind = BT->getKind();
@@ -542,7 +542,7 @@ const Function *Context::getOrCreateFunction(const FunctionDecl *FuncDecl) {
542542
// Assign descriptors to all parameters.
543543
// Composite objects are lowered to pointers.
544544
for (const ParmVarDecl *PD : FuncDecl->parameters()) {
545-
std::optional<PrimType> T = classify(PD->getType());
545+
OptPrimType T = classify(PD->getType());
546546
PrimType PT = T.value_or(PT_Ptr);
547547
Descriptor *Desc = P->createDescriptor(PD, PT);
548548
ParamDescriptors.insert({ParamOffset, {PT, Desc}});
@@ -570,7 +570,7 @@ const Function *Context::getOrCreateObjCBlock(const BlockExpr *E) {
570570
// Assign descriptors to all parameters.
571571
// Composite objects are lowered to pointers.
572572
for (const ParmVarDecl *PD : BD->parameters()) {
573-
std::optional<PrimType> T = classify(PD->getType());
573+
OptPrimType T = classify(PD->getType());
574574
PrimType PT = T.value_or(PT_Ptr);
575575
Descriptor *Desc = P->createDescriptor(PD, PT);
576576
ParamDescriptors.insert({ParamOffset, {PT, Desc}});

‎clang/lib/AST/ByteCode/Context.h

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -82,10 +82,10 @@ class Context final {
8282
uint32_t getBitWidth(QualType T) const { return Ctx.getIntWidth(T); }
8383

8484
/// Classifies a type.
85-
std::optional<PrimType> classify(QualType T) const;
85+
OptPrimType classify(QualType T) const;
8686

8787
/// Classifies an expression.
88-
std::optional<PrimType> classify(const Expr *E) const {
88+
OptPrimType classify(const Expr *E) const {
8989
assert(E);
9090
if (E->isGLValue())
9191
return PT_Ptr;

‎clang/lib/AST/ByteCode/Descriptor.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -164,7 +164,7 @@ struct Descriptor final {
164164
/// The primitive type this descriptor was created for,
165165
/// or the primitive element type in case this is
166166
/// a primitive array.
167-
const std::optional<PrimType> PrimT = std::nullopt;
167+
const OptPrimType PrimT = std::nullopt;
168168
/// Flag indicating if the block is mutable.
169169
const bool IsConst = false;
170170
/// Flag indicating if a field is mutable.

‎clang/lib/AST/ByteCode/EvalEmitter.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -324,7 +324,7 @@ void EvalEmitter::updateGlobalTemporaries() {
324324
const Pointer &Ptr = P.getPtrGlobal(*GlobalIndex);
325325
APValue *Cached = Temp->getOrCreateValue(true);
326326

327-
if (std::optional<PrimType> T = Ctx.classify(E->getType())) {
327+
if (OptPrimType T = Ctx.classify(E->getType())) {
328328
TYPE_SWITCH(
329329
*T, { *Cached = Ptr.deref<T>().toAPValue(Ctx.getASTContext()); });
330330
} else {

‎clang/lib/AST/ByteCode/InterpBuiltin.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -53,7 +53,7 @@ static APSInt popToAPSInt(InterpStack &Stk, PrimType T) {
5353
static void pushInteger(InterpState &S, const APSInt &Val, QualType QT) {
5454
assert(QT->isSignedIntegerOrEnumerationType() ||
5555
QT->isUnsignedIntegerOrEnumerationType());
56-
std::optional<PrimType> T = S.getContext().classify(QT);
56+
OptPrimType T = S.getContext().classify(QT);
5757
assert(T);
5858

5959
unsigned BitWidth = S.getASTContext().getTypeSize(QT);
@@ -1530,7 +1530,7 @@ static bool interp__builtin_operator_new(InterpState &S, CodePtr OpPC,
15301530
return false;
15311531

15321532
bool IsArray = NumElems.ugt(1);
1533-
std::optional<PrimType> ElemT = S.getContext().classify(ElemType);
1533+
OptPrimType ElemT = S.getContext().classify(ElemType);
15341534
DynamicAllocator &Allocator = S.getAllocator();
15351535
if (ElemT) {
15361536
Block *B =
@@ -2879,7 +2879,7 @@ static bool copyRecord(InterpState &S, CodePtr OpPC, const Pointer &Src,
28792879

28802880
auto copyField = [&](const Record::Field &F, bool Activate) -> bool {
28812881
Pointer DestField = Dest.atField(F.Offset);
2882-
if (std::optional<PrimType> FT = S.Ctx.classify(F.Decl->getType())) {
2882+
if (OptPrimType FT = S.Ctx.classify(F.Decl->getType())) {
28832883
TYPE_SWITCH(*FT, {
28842884
DestField.deref<T>() = Src.atField(F.Offset).deref<T>();
28852885
if (Src.atField(F.Offset).isInitialized())

0 commit comments

Comments
 (0)