Skip to content

Commit e2def2e

Browse files
ronliebronlieb
authored
merge main into amd-staging (llvm#1468)
2 parents 0c05887 + e4bb07a commit e2def2e

File tree

175 files changed

+9428
-8169
lines changed

Some content is hidden

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

175 files changed

+9428
-8169
lines changed

‎clang/docs/ReleaseNotes.rst

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -167,6 +167,7 @@ Non-comprehensive list of changes in this release
167167

168168
- Support parsing the `cc` operand modifier and alias it to the `c` modifier (#GH127719).
169169
- Added `__builtin_elementwise_exp10`.
170+
- For AMDPGU targets, added `__builtin_v_cvt_off_f32_i4` that maps to the `v_cvt_off_f32_i4` instruction.
170171

171172
New Compiler Flags
172173
------------------

‎clang/include/clang/AST/StmtOpenACC.h

Lines changed: 26 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -829,24 +829,42 @@ class OpenACCUpdateConstruct final
829829

830830
// This class represents the 'atomic' construct, which has an associated
831831
// statement, but no clauses.
832-
class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct {
832+
class OpenACCAtomicConstruct final
833+
: public OpenACCAssociatedStmtConstruct,
834+
private llvm::TrailingObjects<OpenACCAtomicConstruct,
835+
const OpenACCClause *> {
833836

834837
friend class ASTStmtReader;
838+
friend TrailingObjects;
835839
OpenACCAtomicKind AtomicKind = OpenACCAtomicKind::None;
836840

837-
OpenACCAtomicConstruct(EmptyShell)
841+
OpenACCAtomicConstruct(unsigned NumClauses)
838842
: OpenACCAssociatedStmtConstruct(
839843
OpenACCAtomicConstructClass, OpenACCDirectiveKind::Atomic,
840844
SourceLocation{}, SourceLocation{}, SourceLocation{},
841-
/*AssociatedStmt=*/nullptr) {}
845+
/*AssociatedStmt=*/nullptr) {
846+
std::uninitialized_value_construct(
847+
getTrailingObjects<const OpenACCClause *>(),
848+
getTrailingObjects<const OpenACCClause *>() + NumClauses);
849+
setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
850+
NumClauses));
851+
}
842852

843853
OpenACCAtomicConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
844854
OpenACCAtomicKind AtKind, SourceLocation End,
855+
ArrayRef<const OpenACCClause *> Clauses,
845856
Stmt *AssociatedStmt)
846857
: OpenACCAssociatedStmtConstruct(OpenACCAtomicConstructClass,
847858
OpenACCDirectiveKind::Atomic, Start,
848859
DirectiveLoc, End, AssociatedStmt),
849-
AtomicKind(AtKind) {}
860+
AtomicKind(AtKind) {
861+
// Initialize the trailing storage.
862+
std::uninitialized_copy(Clauses.begin(), Clauses.end(),
863+
getTrailingObjects<const OpenACCClause *>());
864+
865+
setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
866+
Clauses.size()));
867+
}
850868

851869
void setAssociatedStmt(Stmt *S) {
852870
OpenACCAssociatedStmtConstruct::setAssociatedStmt(S);
@@ -857,10 +875,12 @@ class OpenACCAtomicConstruct final : public OpenACCAssociatedStmtConstruct {
857875
return T->getStmtClass() == OpenACCAtomicConstructClass;
858876
}
859877

860-
static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C);
878+
static OpenACCAtomicConstruct *CreateEmpty(const ASTContext &C,
879+
unsigned NumClauses);
861880
static OpenACCAtomicConstruct *
862881
Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
863-
OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt);
882+
OpenACCAtomicKind AtKind, SourceLocation End,
883+
ArrayRef<const OpenACCClause *> Clauses, Stmt *AssociatedStmt);
864884

865885
OpenACCAtomicKind getAtomicKind() const { return AtomicKind; }
866886
const Stmt *getAssociatedStmt() const {

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -140,6 +140,7 @@ BUILTIN(__builtin_amdgcn_cvt_pknorm_u16, "E2Usff", "nc")
140140
BUILTIN(__builtin_amdgcn_cvt_pk_i16, "E2sii", "nc")
141141
BUILTIN(__builtin_amdgcn_cvt_pk_u16, "E2UsUiUi", "nc")
142142
BUILTIN(__builtin_amdgcn_cvt_pk_u8_f32, "UifUiUi", "nc")
143+
BUILTIN(__builtin_amdgcn_cvt_off_f32_i4, "fi", "nc")
143144
BUILTIN(__builtin_amdgcn_sad_u8, "UiUiUiUi", "nc")
144145
BUILTIN(__builtin_amdgcn_msad_u8, "UiUiUiUi", "nc")
145146
BUILTIN(__builtin_amdgcn_sad_hi_u8, "UiUiUiUi", "nc")
@@ -254,7 +255,7 @@ TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s", "t", "at
254255
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t", "atomic-global-pk-add-bf16-inst")
255256
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "atomic-ds-pk-add-16-insts")
256257
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h", "t", "atomic-ds-pk-add-16-insts")
257-
TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "gfx940-insts")
258+
TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi", "t", "vmem-to-lds-load-insts")
258259

259260
//===----------------------------------------------------------------------===//
260261
// Deep learning builtins.

‎clang/include/clang/Basic/DiagnosticGroups.td

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -377,12 +377,13 @@ def CXX11WarnSuggestOverride : DiagGroup<"suggest-override">;
377377
def WarnUnnecessaryVirtualSpecifier : DiagGroup<"unnecessary-virtual-specifier"> {
378378
code Documentation = [{
379379
Warns when a ``final`` class contains a virtual method (including virtual
380-
destructors) that does not override anything. Since ``final`` classes cannot be
381-
subclassed, their methods cannot be overridden, so there is no point to
382-
introducing new ``virtual`` methods.
380+
destructors). Since ``final`` classes cannot be subclassed, their methods
381+
cannot be overridden, and hence the ``virtual`` specifier is useless.
383382

384383
The warning also detects virtual methods in classes whose destructor is
385384
``final``, for the same reason.
385+
386+
The warning does not fire on virtual methods which are also marked ``override``.
386387
}];
387388
}
388389

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2733,7 +2733,7 @@ def note_final_dtor_non_final_class_silence : Note<
27332733
"mark %0 as '%select{final|sealed}1' to silence this warning">;
27342734
def warn_unnecessary_virtual_specifier : Warning<
27352735
"virtual method %0 is inside a 'final' class and can never be overridden">,
2736-
InGroup<WarnUnnecessaryVirtualSpecifier>;
2736+
InGroup<WarnUnnecessaryVirtualSpecifier>, DefaultIgnore;
27372737

27382738
// C++11 attributes
27392739
def err_repeat_attribute : Error<"%0 attribute cannot be repeated">;

‎clang/lib/AST/StmtOpenACC.cpp

Lines changed: 13 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -307,20 +307,26 @@ OpenACCUpdateConstruct::Create(const ASTContext &C, SourceLocation Start,
307307
}
308308

309309
OpenACCAtomicConstruct *
310-
OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C) {
311-
void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct));
312-
auto *Inst = new (Mem) OpenACCAtomicConstruct(EmptyShell{});
310+
OpenACCAtomicConstruct::CreateEmpty(const ASTContext &C, unsigned NumClauses) {
311+
void *Mem = C.Allocate(
312+
OpenACCAtomicConstruct::totalSizeToAlloc<const OpenACCClause *>(
313+
NumClauses));
314+
auto *Inst = new (Mem) OpenACCAtomicConstruct(NumClauses);
313315
return Inst;
314316
}
315317

316318
OpenACCAtomicConstruct *OpenACCAtomicConstruct::Create(
317319
const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
318-
OpenACCAtomicKind AtKind, SourceLocation End, Stmt *AssociatedStmt) {
319-
void *Mem = C.Allocate(sizeof(OpenACCAtomicConstruct));
320-
auto *Inst = new (Mem)
321-
OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind, End, AssociatedStmt);
320+
OpenACCAtomicKind AtKind, SourceLocation End,
321+
ArrayRef<const OpenACCClause *> Clauses, Stmt *AssociatedStmt) {
322+
void *Mem = C.Allocate(
323+
OpenACCAtomicConstruct::totalSizeToAlloc<const OpenACCClause *>(
324+
Clauses.size()));
325+
auto *Inst = new (Mem) OpenACCAtomicConstruct(Start, DirectiveLoc, AtKind,
326+
End, Clauses, AssociatedStmt);
322327
return Inst;
323328
}
329+
324330
OpenACCCacheConstruct *OpenACCCacheConstruct::CreateEmpty(const ASTContext &C,
325331
unsigned NumVars) {
326332
void *Mem =

‎clang/lib/AST/StmtPrinter.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1258,6 +1258,7 @@ void StmtPrinter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
12581258
if (S->getAtomicKind() != OpenACCAtomicKind::None)
12591259
OS << " " << S->getAtomicKind();
12601260

1261+
PrintOpenACCClauseList(S);
12611262
OS << '\n';
12621263
PrintStmt(S->getAssociatedStmt());
12631264
}

‎clang/lib/AST/StmtProfile.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2839,6 +2839,8 @@ void StmtProfiler::VisitOpenACCUpdateConstruct(
28392839
void StmtProfiler::VisitOpenACCAtomicConstruct(
28402840
const OpenACCAtomicConstruct *S) {
28412841
VisitStmt(S);
2842+
OpenACCClauseProfiler P{*this};
2843+
P.VisitOpenACCClauseList(S->clauses());
28422844
}
28432845

28442846
void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) {

‎clang/lib/Basic/Targets/AMDGPU.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -260,7 +260,7 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple,
260260

261261
MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
262262
CUMode = !(GPUFeatures & llvm::AMDGPU::FEATURE_WGP);
263-
for (auto F : {"image-insts", "gws"})
263+
for (auto F : {"image-insts", "gws", "vmem-to-lds-load-insts"})
264264
ReadOnlyFeatures.insert(F);
265265
HalfArgsAndReturns = true;
266266
}

‎clang/lib/Sema/SemaOpenACC.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1832,9 +1832,8 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(
18321832
EndLoc, Clauses);
18331833
}
18341834
case OpenACCDirectiveKind::Atomic: {
1835-
assert(Clauses.empty() && "Atomic doesn't allow clauses");
18361835
return OpenACCAtomicConstruct::Create(
1837-
getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc,
1836+
getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc, Clauses,
18381837
AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
18391838
}
18401839
case OpenACCDirectiveKind::Cache: {

‎clang/lib/Sema/SemaOpenACCClause.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -55,6 +55,8 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
5555
case OpenACCDirectiveKind::ParallelLoop:
5656
case OpenACCDirectiveKind::SerialLoop:
5757
case OpenACCDirectiveKind::KernelsLoop:
58+
// OpenACC 3.4(prerelease) PR #511 adds 'if' to atomic.
59+
case OpenACCDirectiveKind::Atomic:
5860
return true;
5961
default:
6062
return false;

‎clang/lib/Sema/TreeTransform.h

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4222,10 +4222,11 @@ class TreeTransform {
42224222
SourceLocation DirLoc,
42234223
OpenACCAtomicKind AtKind,
42244224
SourceLocation EndLoc,
4225+
ArrayRef<OpenACCClause *> Clauses,
42254226
StmtResult AssociatedStmt) {
42264227
return getSema().OpenACC().ActOnEndStmtDirective(
42274228
OpenACCDirectiveKind::Atomic, BeginLoc, DirLoc, SourceLocation{},
4228-
SourceLocation{}, {}, AtKind, SourceLocation{}, EndLoc, {},
4229+
SourceLocation{}, {}, AtKind, SourceLocation{}, EndLoc, Clauses,
42294230
AssociatedStmt);
42304231
}
42314232

@@ -12744,6 +12745,10 @@ StmtResult TreeTransform<Derived>::TransformOpenACCAtomicConstruct(
1274412745
OpenACCAtomicConstruct *C) {
1274512746
getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
1274612747

12748+
llvm::SmallVector<OpenACCClause *> TransformedClauses =
12749+
getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
12750+
C->clauses());
12751+
1274712752
if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
1274812753
C->getBeginLoc(), {}))
1274912754
return StmtError();
@@ -12759,7 +12764,7 @@ StmtResult TreeTransform<Derived>::TransformOpenACCAtomicConstruct(
1275912764

1276012765
return getDerived().RebuildOpenACCAtomicConstruct(
1276112766
C->getBeginLoc(), C->getDirectiveLoc(), C->getAtomicKind(),
12762-
C->getEndLoc(), AssocStmt);
12767+
C->getEndLoc(), TransformedClauses, AssocStmt);
1276312768
}
1276412769

1276512770
template <typename Derived>

‎clang/lib/Serialization/ASTReaderStmt.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2940,9 +2940,7 @@ void ASTStmtReader::VisitOpenACCCacheConstruct(OpenACCCacheConstruct *S) {
29402940

29412941
void ASTStmtReader::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
29422942
VisitStmt(S);
2943-
S->Kind = Record.readEnum<OpenACCDirectiveKind>();
2944-
S->Range = Record.readSourceRange();
2945-
S->DirectiveLoc = Record.readSourceLocation();
2943+
VisitOpenACCConstructStmt(S);
29462944
S->AtomicKind = Record.readEnum<OpenACCAtomicKind>();
29472945
S->setAssociatedStmt(Record.readSubStmt());
29482946
}
@@ -4490,7 +4488,8 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
44904488
break;
44914489
}
44924490
case STMT_OPENACC_ATOMIC_CONSTRUCT: {
4493-
S = OpenACCAtomicConstruct::CreateEmpty(Context);
4491+
unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
4492+
S = OpenACCAtomicConstruct::CreateEmpty(Context, NumClauses);
44944493
break;
44954494
}
44964495
case EXPR_REQUIRES: {

‎clang/lib/Serialization/ASTWriterStmt.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -3014,9 +3014,7 @@ void ASTStmtWriter::VisitOpenACCWaitConstruct(OpenACCWaitConstruct *S) {
30143014

30153015
void ASTStmtWriter::VisitOpenACCAtomicConstruct(OpenACCAtomicConstruct *S) {
30163016
VisitStmt(S);
3017-
Record.writeEnum(S->Kind);
3018-
Record.AddSourceRange(S->Range);
3019-
Record.AddSourceLocation(S->DirectiveLoc);
3017+
VisitOpenACCConstructStmt(S);
30203018
Record.writeEnum(S->getAtomicKind());
30213019
Record.AddStmt(S->getAssociatedStmt());
30223020

‎clang/test/AST/ast-print-openacc-atomic-construct.cpp

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,3 +31,35 @@ void foo(int v, int x) {
3131
{ x--; v = x; }
3232

3333
}
34+
35+
void foo2(int v, int x) {
36+
// CHECK: #pragma acc atomic read if(v)
37+
// CHECK-NEXT: v = x;
38+
#pragma acc atomic read if (v)
39+
v = x;
40+
// CHECK-NEXT: pragma acc atomic write if(x)
41+
// CHECK-NEXT: v = x + 1;
42+
#pragma acc atomic write if (x)
43+
v = x + 1;
44+
// CHECK-NEXT: pragma acc atomic update if(true)
45+
// CHECK-NEXT: x++;
46+
#pragma acc atomic update if (true)
47+
x++;
48+
// CHECK-NEXT: pragma acc atomic if(false)
49+
// CHECK-NEXT: x--;
50+
#pragma acc atomic if (false)
51+
x--;
52+
// CHECK-NEXT: pragma acc atomic capture if(v < x)
53+
// CHECK-NEXT: v = x++;
54+
#pragma acc atomic capture if (v < x)
55+
v = x++;
56+
57+
// CHECK-NEXT: #pragma acc atomic capture if(x > v)
58+
// CHECK-NEXT: {
59+
// CHECK-NEXT: x--;
60+
// CHECK-NEXT: v = x;
61+
// CHECK-NEXT: }
62+
#pragma acc atomic capture if (x > v)
63+
{ x--; v = x; }
64+
65+
}

‎clang/test/Analysis/Checkers/WebKit/ref-cntbl-crtp-base-no-virtual-dtor.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_analyze_cc1 -analyzer-checker=webkit.RefCntblBaseVirtualDtor -verify %s -Wno-unnecessary-virtual-specifier
1+
// RUN: %clang_analyze_cc1 -analyzer-checker=webkit.RefCntblBaseVirtualDtor -verify %s
22

33
#include "mock-types.h"
44

‎clang/test/CXX/class/p2-0x.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -28,7 +28,7 @@ struct C : A<int> { }; // expected-error {{base 'A' is marked 'final'}}
2828

2929
namespace Test4 {
3030

31-
struct A final { virtual void func() = 0; }; // expected-warning {{abstract class is marked 'final'}} expected-note {{unimplemented pure virtual method 'func' in 'A'}} expected-warning {{virtual method 'func' is inside a 'final' class}}}
31+
struct A final { virtual void func() = 0; }; // expected-warning {{abstract class is marked 'final'}} expected-note {{unimplemented pure virtual method 'func' in 'A'}}
3232
struct B { virtual void func() = 0; }; // expected-note {{unimplemented pure virtual method 'func' in 'C'}}
3333

3434
struct C final : B { }; // expected-warning {{abstract class is marked 'final'}}

‎clang/test/CodeGen/attr-btf_type_tag-similar-type.c

Lines changed: 15 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,21 @@
11
// RUN: %clang_cc1 -triple %itanium_abi_triple -debug-info-kind=limited -emit-llvm -o - %s | FileCheck %s
2+
// RUN: %clang_cc1 -triple %itanium_abi_triple -DDOUBLE_BRACKET_ATTRS=1 -debug-info-kind=limited -emit-llvm -o - %s | FileCheck %s
3+
4+
#if DOUBLE_BRACKET_ATTRS
5+
#define __tag1 [[clang::btf_type_tag("tag1")]]
6+
#define __tag2 [[clang::btf_type_tag("tag2")]]
7+
#define __tag3 [[clang::btf_type_tag("tag3")]]
8+
#define __tag4 [[clang::btf_type_tag("tag4")]]
9+
#else
10+
#define __tag1 __attribute__((btf_type_tag("tag1")))
11+
#define __tag2 __attribute__((btf_type_tag("tag2")))
12+
#define __tag3 __attribute__((btf_type_tag("tag3")))
13+
#define __tag4 __attribute__((btf_type_tag("tag4")))
14+
#endif
215

316
struct map_value {
4-
int __attribute__((btf_type_tag("tag1"))) __attribute__((btf_type_tag("tag3"))) *a;
5-
int __attribute__((btf_type_tag("tag2"))) __attribute__((btf_type_tag("tag4"))) *b;
17+
int __tag1 __tag3 *a;
18+
int __tag2 __tag4 *b;
619
};
720

821
struct map_value *func(void);

‎clang/test/CodeGen/attr-btf_type_tag-typedef-field.c

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,13 @@
11
// RUN: %clang_cc1 -triple %itanium_abi_triple -debug-info-kind=limited -emit-llvm -o - %s | FileCheck %s
2+
// RUN: %clang_cc1 -triple %itanium_abi_triple -DDOUBLE_BRACKET_ATTRS=1 -debug-info-kind=limited -emit-llvm -o - %s | FileCheck %s
23

4+
#if DOUBLE_BRACKET_ATTRS
5+
#define __tag1 [[clang::btf_type_tag("tag1")]]
6+
#define __tag2 [[clang::btf_type_tag("tag2")]]
7+
#else
38
#define __tag1 __attribute__((btf_type_tag("tag1")))
49
#define __tag2 __attribute__((btf_type_tag("tag2")))
10+
#endif
511

612
typedef void __fn_t(int);
713
typedef __fn_t __tag1 __tag2 *__fn2_t;
@@ -31,5 +37,5 @@ int *foo1(struct t *a1) {
3137
// CHECK: ![[L28]] = !DISubroutineType(types: ![[L29:[0-9]+]])
3238
// CHECK: ![[L29]] = !{null, ![[L4]]}
3339
// CHECK: ![[L30]] = !{![[L21]], ![[L23]]}
34-
// CHECK: ![[L31]] = !DIDerivedType(tag: DW_TAG_member, name: "c", scope: ![[#]], file: ![[#]], line: [[#]]1, baseType: ![[L32:[0-9]+]]
40+
// CHECK: ![[L31]] = !DIDerivedType(tag: DW_TAG_member, name: "c", scope: ![[#]], file: ![[#]], line: [[#]], baseType: ![[L32:[0-9]+]]
3541
// CHECK: ![[L32]] = !DIBasicType(name: "long", size: [[#]], encoding: DW_ATE_signed)

‎clang/test/CodeGen/link-builtin-bitcode.c

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -44,6 +44,6 @@ int bar() { return no_attr() + attr_in_target() + attr_not_in_target() + attr_in
4444
// CHECK-SAME: () #[[ATTR_INCOMPATIBLE:[0-9]+]] {
4545

4646
// CHECK: attributes #[[ATTR_BAR]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
47-
// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
48-
// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+extended-image-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
49-
// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,-gfx9-insts" }
47+
// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+vmem-to-lds-load-insts,+wavefrontsize64" }
48+
// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+extended-image-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+vmem-to-lds-load-insts,+wavefrontsize64" }
49+
// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-cpu"="gfx90a" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" }

0 commit comments

Comments
 (0)