Skip to content

Commit 0e5a641

Browse files
authored
merge main into amd-staging (#792)
2 parents 8b12e8f + ec5efc4 commit 0e5a641

File tree

98 files changed

+1609
-728
lines changed

Some content is hidden

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

98 files changed

+1609
-728
lines changed

clang-tools-extra/clang-tidy/fuchsia/MultipleInheritanceCheck.cpp

Lines changed: 40 additions & 85 deletions
Original file line numberDiff line numberDiff line change
@@ -17,106 +17,61 @@ namespace clang::tidy::fuchsia {
1717

1818
namespace {
1919
AST_MATCHER(CXXRecordDecl, hasBases) {
20-
if (Node.hasDefinition())
21-
return Node.getNumBases() > 0;
22-
return false;
20+
return Node.hasDefinition() && Node.getNumBases() > 0;
2321
}
2422
} // namespace
2523

26-
// Adds a node to the interface map, if it was not present in the map
27-
// previously.
28-
void MultipleInheritanceCheck::addNodeToInterfaceMap(const CXXRecordDecl *Node,
29-
bool IsInterface) {
30-
InterfaceMap.try_emplace(Node, IsInterface);
31-
}
32-
33-
// Returns "true" if the boolean "isInterface" has been set to the
34-
// interface status of the current Node. Return "false" if the
35-
// interface status for the current node is not yet known.
36-
bool MultipleInheritanceCheck::getInterfaceStatus(const CXXRecordDecl *Node,
37-
bool &IsInterface) const {
38-
auto Pair = InterfaceMap.find(Node);
39-
if (Pair == InterfaceMap.end())
40-
return false;
41-
IsInterface = Pair->second;
42-
return true;
43-
}
24+
bool MultipleInheritanceCheck::isInterface(const CXXBaseSpecifier &Base) {
25+
const CXXRecordDecl *const Node = Base.getType()->getAsCXXRecordDecl();
26+
if (!Node)
27+
return true;
4428

45-
bool MultipleInheritanceCheck::isCurrentClassInterface(
46-
const CXXRecordDecl *Node) const {
47-
// Interfaces should have no fields.
48-
if (!Node->field_empty())
49-
return false;
50-
51-
// Interfaces should have exclusively pure methods.
52-
return llvm::none_of(Node->methods(), [](const CXXMethodDecl *M) {
53-
return M->isUserProvided() && !M->isPureVirtual() && !M->isStatic();
54-
});
55-
}
29+
assert(Node->isCompleteDefinition());
5630

57-
bool MultipleInheritanceCheck::isInterface(const CXXRecordDecl *Node) {
5831
// Short circuit the lookup if we have analyzed this record before.
59-
bool PreviousIsInterfaceResult = false;
60-
if (getInterfaceStatus(Node, PreviousIsInterfaceResult))
61-
return PreviousIsInterfaceResult;
62-
63-
// To be an interface, all base classes must be interfaces as well.
64-
for (const auto &I : Node->bases()) {
65-
if (I.isVirtual())
66-
continue;
67-
const auto *Base = I.getType()->getAsCXXRecordDecl();
68-
if (!Base)
69-
continue;
70-
assert(Base->isCompleteDefinition());
71-
if (!isInterface(Base)) {
72-
addNodeToInterfaceMap(Node, false);
73-
return false;
74-
}
75-
}
76-
77-
const bool CurrentClassIsInterface = isCurrentClassInterface(Node);
78-
addNodeToInterfaceMap(Node, CurrentClassIsInterface);
32+
if (const auto CachedValue = InterfaceMap.find(Node);
33+
CachedValue != InterfaceMap.end())
34+
return CachedValue->second;
35+
36+
// To be an interface, a class must have...
37+
const bool CurrentClassIsInterface =
38+
// ...no bases that aren't interfaces...
39+
llvm::none_of(Node->bases(),
40+
[&](const CXXBaseSpecifier &I) {
41+
return !I.isVirtual() && !isInterface(I);
42+
}) &&
43+
// ...no fields, and...
44+
Node->field_empty() &&
45+
// ...no methods that aren't pure virtual.
46+
llvm::none_of(Node->methods(), [](const CXXMethodDecl *M) {
47+
return M->isUserProvided() && !M->isPureVirtual() && !M->isStatic();
48+
});
49+
50+
InterfaceMap.try_emplace(Node, CurrentClassIsInterface);
7951
return CurrentClassIsInterface;
8052
}
8153

8254
void MultipleInheritanceCheck::registerMatchers(MatchFinder *Finder) {
83-
// Match declarations which have bases.
8455
Finder->addMatcher(cxxRecordDecl(hasBases(), isDefinition()).bind("decl"),
8556
this);
8657
}
8758

8859
void MultipleInheritanceCheck::check(const MatchFinder::MatchResult &Result) {
89-
if (const auto *D = Result.Nodes.getNodeAs<CXXRecordDecl>("decl")) {
90-
// Check against map to see if the class inherits from multiple
91-
// concrete classes
92-
unsigned NumConcrete = 0;
93-
for (const auto &I : D->bases()) {
94-
if (I.isVirtual())
95-
continue;
96-
const auto *Base = I.getType()->getAsCXXRecordDecl();
97-
if (!Base)
98-
continue;
99-
assert(Base->isCompleteDefinition());
100-
if (!isInterface(Base))
101-
NumConcrete++;
102-
}
103-
104-
// Check virtual bases to see if there is more than one concrete
105-
// non-virtual base.
106-
for (const auto &V : D->vbases()) {
107-
const auto *Base = V.getType()->getAsCXXRecordDecl();
108-
if (!Base)
109-
continue;
110-
assert(Base->isCompleteDefinition());
111-
if (!isInterface(Base))
112-
NumConcrete++;
113-
}
114-
115-
if (NumConcrete > 1) {
116-
diag(D->getBeginLoc(), "inheriting multiple classes that aren't "
117-
"pure virtual is discouraged");
118-
}
119-
}
60+
const auto &D = *Result.Nodes.getNodeAs<CXXRecordDecl>("decl");
61+
// Check to see if the class inherits from multiple concrete classes.
62+
unsigned NumConcrete =
63+
llvm::count_if(D.bases(), [&](const CXXBaseSpecifier &I) {
64+
return !I.isVirtual() && !isInterface(I);
65+
});
66+
67+
// Check virtual bases to see if there is more than one concrete
68+
// non-virtual base.
69+
NumConcrete += llvm::count_if(
70+
D.vbases(), [&](const CXXBaseSpecifier &V) { return !isInterface(V); });
71+
72+
if (NumConcrete > 1)
73+
diag(D.getBeginLoc(), "inheriting multiple classes that aren't "
74+
"pure virtual is discouraged");
12075
}
12176

12277
} // namespace clang::tidy::fuchsia

clang-tools-extra/clang-tidy/fuchsia/MultipleInheritanceCheck.h

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -30,10 +30,7 @@ class MultipleInheritanceCheck : public ClangTidyCheck {
3030
void onEndOfTranslationUnit() override { InterfaceMap.clear(); }
3131

3232
private:
33-
void addNodeToInterfaceMap(const CXXRecordDecl *Node, bool IsInterface);
34-
bool getInterfaceStatus(const CXXRecordDecl *Node, bool &IsInterface) const;
35-
bool isCurrentClassInterface(const CXXRecordDecl *Node) const;
36-
bool isInterface(const CXXRecordDecl *Node);
33+
bool isInterface(const CXXBaseSpecifier &Base);
3734

3835
// Contains the identity of each named CXXRecord as an interface. This is
3936
// used to memoize lookup speeds and improve performance from O(N^2) to O(N),

clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 30 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -910,75 +910,75 @@ TARGET_BUILTIN(__builtin_amdgcn_cooperative_atomic_store_8x16B, "vV4i*V4iIicC*",
910910
// Image builtins
911911
//===----------------------------------------------------------------------===//
912912
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f32_i32, "V4fiiQtii", "nc", "image-insts")
913-
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4hiiQtii", "nc", "image-insts")
913+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1d_v4f16_i32, "V4xiiQtii", "nc", "image-insts")
914914
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
915-
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
915+
TARGET_BUILTIN(__builtin_amdgcn_image_load_1darray_v4f16_i32, "V4xiiiQtii", "nc", "image-insts")
916916
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_f32_i32, "fiiiQtii", "nc", "image-insts")
917917
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
918-
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
918+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2d_v4f16_i32, "V4xiiiQtii", "nc", "image-insts")
919919
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_f32_i32, "fiiiiQtii", "nc", "image-insts")
920920
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
921-
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
921+
TARGET_BUILTIN(__builtin_amdgcn_image_load_2darray_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
922922
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
923-
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
923+
TARGET_BUILTIN(__builtin_amdgcn_image_load_3d_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
924924
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
925-
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
925+
TARGET_BUILTIN(__builtin_amdgcn_image_load_cube_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
926926
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f32_i32, "V4fiiiQtii", "nc", "image-insts")
927-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4hiiiQtii", "nc", "image-insts")
927+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1d_v4f16_i32, "V4xiiiQtii", "nc", "image-insts")
928928
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
929-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
929+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_1darray_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
930930
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_f32_i32, "fiiiiQtii", "nc", "image-insts")
931931
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f32_i32, "V4fiiiiQtii", "nc", "image-insts")
932-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4hiiiiQtii", "nc", "image-insts")
932+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2d_v4f16_i32, "V4xiiiiQtii", "nc", "image-insts")
933933
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_f32_i32, "fiiiiiQtii", "nc", "image-insts")
934934
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
935-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
935+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_2darray_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts")
936936
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
937-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
937+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_3d_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts")
938938
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f32_i32, "V4fiiiiiQtii", "nc", "image-insts")
939-
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4hiiiiiQtii", "nc", "image-insts")
939+
TARGET_BUILTIN(__builtin_amdgcn_image_load_mip_cube_v4f16_i32, "V4xiiiiiQtii", "nc", "image-insts")
940940
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f32_i32, "vV4fiiQtii", "nc", "image-insts")
941-
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4hiiQtii", "nc", "image-insts")
941+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1d_v4f16_i32, "vV4xiiQtii", "nc", "image-insts")
942942
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
943-
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
943+
TARGET_BUILTIN(__builtin_amdgcn_image_store_1darray_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts")
944944
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_f32_i32, "vfiiiQtii", "nc", "image-insts")
945945
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
946-
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
946+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2d_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts")
947947
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_f32_i32, "vfiiiiQtii", "nc", "image-insts")
948948
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
949-
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
949+
TARGET_BUILTIN(__builtin_amdgcn_image_store_2darray_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
950950
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
951-
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
951+
TARGET_BUILTIN(__builtin_amdgcn_image_store_3d_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
952952
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
953-
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
953+
TARGET_BUILTIN(__builtin_amdgcn_image_store_cube_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
954954
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f32_i32, "vV4fiiiQtii", "nc", "image-insts")
955-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4hiiiQtii", "nc", "image-insts")
955+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1d_v4f16_i32, "vV4xiiiQtii", "nc", "image-insts")
956956
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
957-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
957+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_1darray_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
958958
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_f32_i32, "vfiiiiQtii", "nc", "image-insts")
959959
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f32_i32, "vV4fiiiiQtii", "nc", "image-insts")
960-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4hiiiiQtii", "nc", "image-insts")
960+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2d_v4f16_i32, "vV4xiiiiQtii", "nc", "image-insts")
961961
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_f32_i32, "vfiiiiiQtii", "nc", "image-insts")
962962
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
963-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
963+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_2darray_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts")
964964
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
965-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
965+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_3d_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts")
966966
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f32_i32, "vV4fiiiiiQtii", "nc", "image-insts")
967-
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4hiiiiiQtii", "nc", "image-insts")
967+
TARGET_BUILTIN(__builtin_amdgcn_image_store_mip_cube_v4f16_i32, "vV4xiiiiiQtii", "nc", "image-insts")
968968
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "image-insts")
969-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4hifQtV4ibii", "nc", "image-insts")
969+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1d_v4f16_f32, "V4xifQtV4ibii", "nc", "image-insts")
970970
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts")
971-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts")
971+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_1darray_v4f16_f32, "V4xiffQtV4ibii", "nc", "image-insts")
972972
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_f32_f32, "fiffQtV4ibii", "nc", "image-insts")
973973
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f32_f32, "V4fiffQtV4ibii", "nc", "image-insts")
974-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4hiffQtV4ibii", "nc", "image-insts")
974+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2d_v4f16_f32, "V4xiffQtV4ibii", "nc", "image-insts")
975975
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_f32_f32, "fifffQtV4ibii", "nc", "image-insts")
976976
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
977-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
977+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_2darray_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts")
978978
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
979-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
979+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_3d_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts")
980980
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f32_f32, "V4fifffQtV4ibii", "nc", "image-insts")
981-
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4hifffQtV4ibii", "nc", "image-insts")
981+
TARGET_BUILTIN(__builtin_amdgcn_image_sample_cube_v4f16_f32, "V4xifffQtV4ibii", "nc", "image-insts")
982982
TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f32_f32, "V4fifQtV4ibii", "nc", "extended-image-insts")
983983
TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1d_v4f16_f32, "V4eifQtV4ibii", "nc", "extended-image-insts")
984984
TARGET_BUILTIN(__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32, "V4fiffQtV4ibii", "nc", "extended-image-insts")

clang/lib/AST/ASTContext.cpp

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10532,6 +10532,21 @@ bool ASTContext::areCompatibleVectorTypes(QualType FirstVec,
1053210532
Second->getVectorKind() != VectorKind::RVVFixedLengthMask_4)
1053310533
return true;
1053410534

10535+
// In OpenCL, treat half and _Float16 vector types as compatible.
10536+
if (getLangOpts().OpenCL &&
10537+
First->getNumElements() == Second->getNumElements()) {
10538+
QualType FirstElt = First->getElementType();
10539+
QualType SecondElt = Second->getElementType();
10540+
10541+
if ((FirstElt->isFloat16Type() && SecondElt->isHalfType()) ||
10542+
(FirstElt->isHalfType() && SecondElt->isFloat16Type())) {
10543+
if (First->getVectorKind() != VectorKind::AltiVecPixel &&
10544+
First->getVectorKind() != VectorKind::AltiVecBool &&
10545+
Second->getVectorKind() != VectorKind::AltiVecPixel &&
10546+
Second->getVectorKind() != VectorKind::AltiVecBool)
10547+
return true;
10548+
}
10549+
}
1053510550
return false;
1053610551
}
1053710552

clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -362,6 +362,20 @@ class OpenACCRoutineClauseEmitter final
362362
for (const DeviceTypeArgument &arg : clause.getArchitectures())
363363
lastDeviceTypeValues.push_back(decodeDeviceType(arg.getIdentifierInfo()));
364364
}
365+
366+
void VisitBindClause(const OpenACCBindClause &clause) {
367+
if (clause.isStringArgument()) {
368+
mlir::StringAttr value =
369+
builder.getStringAttr(clause.getStringArgument()->getString());
370+
371+
routineOp.addBindStrName(builder.getContext(), lastDeviceTypeValues,
372+
value);
373+
} else {
374+
assert(clause.isIdentifierArgument());
375+
cgm.errorNYI(clause.getSourceRange(),
376+
"Bind with an identifier argument is not yet supported");
377+
}
378+
}
365379
};
366380
} // namespace
367381

0 commit comments

Comments
 (0)