Skip to content

Commit ef8e401

Browse files
AaronBallmanAaronBallman
authored
[SYCL] Allow [[sycl::work_group_size_hint]] to accept constant expr args (#3785)
This changes the arguments from taking simple integer values to instead accept arbitrary integer constant expression arguments. Additionally, it implements the logic to correctly handle merging attributes on redeclarations.
1 parent f68f787 commit ef8e401

File tree

8 files changed

+192
-33
lines changed

8 files changed

+192
-33
lines changed

‎clang/include/clang/Basic/Attr.td

Lines changed: 20 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2975,10 +2975,27 @@ def ReqdWorkGroupSize : InheritableAttr {
29752975
def WorkGroupSizeHint : InheritableAttr {
29762976
let Spellings = [GNU<"work_group_size_hint">,
29772977
CXX11<"sycl", "work_group_size_hint">];
2978-
let Args = [UnsignedArgument<"XDim">,
2979-
UnsignedArgument<"YDim">,
2980-
UnsignedArgument<"ZDim">];
2978+
let Args = [ExprArgument<"XDim">,
2979+
ExprArgument<"YDim">,
2980+
ExprArgument<"ZDim">];
29812981
let Subjects = SubjectList<[Function], ErrorDiag>;
2982+
let AdditionalMembers = [{
2983+
Optional<llvm::APSInt> getXDimVal() const {
2984+
if (const auto *CE = dyn_cast<ConstantExpr>(getXDim()))
2985+
return CE->getResultAsAPSInt();
2986+
return None;
2987+
}
2988+
Optional<llvm::APSInt> getYDimVal() const {
2989+
if (const auto *CE = dyn_cast<ConstantExpr>(getYDim()))
2990+
return CE->getResultAsAPSInt();
2991+
return None;
2992+
}
2993+
Optional<llvm::APSInt> getZDimVal() const {
2994+
if (const auto *CE = dyn_cast<ConstantExpr>(getZDim()))
2995+
return CE->getResultAsAPSInt();
2996+
return None;
2997+
}
2998+
}];
29822999
let Documentation = [WorkGroupSizeHintAttrDocs];
29833000
}
29843001

‎clang/include/clang/Sema/Sema.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10306,6 +10306,10 @@ class Sema final {
1030610306
template <typename AttrType>
1030710307
void addIntelTripleArgAttr(Decl *D, const AttributeCommonInfo &CI,
1030810308
Expr *XDimExpr, Expr *YDimExpr, Expr *ZDimExpr);
10309+
void AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
10310+
Expr *XDim, Expr *YDim, Expr *ZDim);
10311+
WorkGroupSizeHintAttr *
10312+
MergeWorkGroupSizeHintAttr(Decl *D, const WorkGroupSizeHintAttr &A);
1030910313
void AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
1031010314
Expr *E);
1031110315
IntelReqdSubGroupSizeAttr *

‎clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -622,9 +622,9 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
622622

623623
if (const WorkGroupSizeHintAttr *A = FD->getAttr<WorkGroupSizeHintAttr>()) {
624624
llvm::Metadata *AttrMDArgs[] = {
625-
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())),
626-
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())),
627-
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))};
625+
llvm::ConstantAsMetadata::get(Builder.getInt(*A->getXDimVal())),
626+
llvm::ConstantAsMetadata::get(Builder.getInt(*A->getYDimVal())),
627+
llvm::ConstantAsMetadata::get(Builder.getInt(*A->getZDimVal()))};
628628
Fn->setMetadata("work_group_size_hint", llvm::MDNode::get(Context, AttrMDArgs));
629629
}
630630

‎clang/lib/Sema/SemaDecl.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2641,6 +2641,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
26412641
NewAttr = S.MergeIntelFPGAForcePow2DepthAttr(D, *A);
26422642
else if (const auto *A = dyn_cast<SYCLIntelFPGAInitiationIntervalAttr>(Attr))
26432643
NewAttr = S.MergeSYCLIntelFPGAInitiationIntervalAttr(D, *A);
2644+
else if (const auto *A = dyn_cast<WorkGroupSizeHintAttr>(Attr))
2645+
NewAttr = S.MergeWorkGroupSizeHintAttr(D, *A);
26442646
else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr))
26452647
NewAttr = cast<InheritableAttr>(Attr->clone(S.Context));
26462648

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 100 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -3160,31 +3160,113 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
31603160
S.addIntelTripleArgAttr<WorkGroupAttr>(D, AL, XDimExpr, YDimExpr, ZDimExpr);
31613161
}
31623162

3163-
// Handles work_group_size_hint.
3164-
static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
3165-
S.CheckDeprecatedSYCLAttributeSpelling(AL);
3163+
// Returns a DupArgResult value; Same means the args have the same value,
3164+
// Different means the args do not have the same value, and Unknown means that
3165+
// the args cannot (yet) be compared.
3166+
enum class DupArgResult { Unknown, Same, Different };
3167+
static DupArgResult AreArgValuesIdentical(const Expr *LHS, const Expr *RHS) {
3168+
// If either operand is still value dependent, we can't test anything.
3169+
const auto *LHSCE = dyn_cast<ConstantExpr>(LHS);
3170+
const auto *RHSCE = dyn_cast<ConstantExpr>(RHS);
3171+
if (!LHSCE || !RHSCE)
3172+
return DupArgResult::Unknown;
3173+
3174+
// Otherwise, test that the values.
3175+
return LHSCE->getResultAsAPSInt() == RHSCE->getResultAsAPSInt()
3176+
? DupArgResult::Same
3177+
: DupArgResult::Different;
3178+
}
3179+
3180+
void Sema::AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
3181+
Expr *XDim, Expr *YDim, Expr *ZDim) {
3182+
// Returns nullptr if diagnosing, otherwise returns the original expression
3183+
// or the original expression converted to a constant expression.
3184+
auto CheckAndConvertArg = [&](Expr *E) -> Expr * {
3185+
// We can only check if the expression is not value dependent.
3186+
if (!E->isValueDependent()) {
3187+
llvm::APSInt ArgVal;
3188+
ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal);
3189+
if (Res.isInvalid())
3190+
return nullptr;
3191+
E = Res.get();
31663192

3167-
uint32_t WGSize[3];
3168-
for (unsigned i = 0; i < AL.getNumArgs(); ++i) {
3169-
if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i,
3170-
/*StrictlyUnsigned=*/true))
3171-
return;
3193+
// This attribute requires a strictly positive value.
3194+
if (ArgVal <= 0) {
3195+
Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
3196+
<< CI << /*positive*/ 0;
3197+
return nullptr;
3198+
}
3199+
}
3200+
3201+
return E;
3202+
};
31723203

3173-
if (WGSize[i] == 0) {
3174-
S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero)
3175-
<< AL << AL.getArgAsExpr(i)->getSourceRange();
3204+
// Check all three argument values, and if any are bad, bail out. This will
3205+
// convert the given expressions into constant expressions when possible.
3206+
XDim = CheckAndConvertArg(XDim);
3207+
YDim = CheckAndConvertArg(YDim);
3208+
ZDim = CheckAndConvertArg(ZDim);
3209+
if (!XDim || !YDim || !ZDim)
3210+
return;
3211+
3212+
// If the attribute was already applied with different arguments, then
3213+
// diagnose the second attribute as a duplicate and don't add it.
3214+
if (const auto *Existing = D->getAttr<WorkGroupSizeHintAttr>()) {
3215+
DupArgResult Results[] = {AreArgValuesIdentical(XDim, Existing->getXDim()),
3216+
AreArgValuesIdentical(YDim, Existing->getYDim()),
3217+
AreArgValuesIdentical(ZDim, Existing->getZDim())};
3218+
// If any of the results are known to be different, we can diagnose at this
3219+
// point and drop the attribute.
3220+
if (llvm::is_contained(Results, DupArgResult::Different)) {
3221+
Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
3222+
Diag(Existing->getLoc(), diag::note_previous_attribute);
31763223
return;
31773224
}
3225+
// If all of the results are known to be the same, we can silently drop the
3226+
// attribute. Otherwise, we have to add the attribute and resolve its
3227+
// differences later.
3228+
if (llvm::all_of(Results,
3229+
[](DupArgResult V) { return V == DupArgResult::Same; }))
3230+
return;
31783231
}
31793232

3180-
WorkGroupSizeHintAttr *Existing = D->getAttr<WorkGroupSizeHintAttr>();
3181-
if (Existing &&
3182-
!(Existing->getXDim() == WGSize[0] && Existing->getYDim() == WGSize[1] &&
3183-
Existing->getZDim() == WGSize[2]))
3184-
S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL;
3233+
D->addAttr(::new (Context)
3234+
WorkGroupSizeHintAttr(Context, CI, XDim, YDim, ZDim));
3235+
}
3236+
3237+
WorkGroupSizeHintAttr *
3238+
Sema::MergeWorkGroupSizeHintAttr(Decl *D, const WorkGroupSizeHintAttr &A) {
3239+
// Check to see if there's a duplicate attribute already applied.
3240+
if (const auto *DeclAttr = D->getAttr<WorkGroupSizeHintAttr>()) {
3241+
DupArgResult Results[] = {
3242+
AreArgValuesIdentical(DeclAttr->getXDim(), A.getXDim()),
3243+
AreArgValuesIdentical(DeclAttr->getYDim(), A.getYDim()),
3244+
AreArgValuesIdentical(DeclAttr->getZDim(), A.getZDim())};
3245+
3246+
// If any of the results are known to be different, we can diagnose at this
3247+
// point and drop the attribute.
3248+
if (llvm::is_contained(Results, DupArgResult::Different)) {
3249+
Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
3250+
Diag(A.getLoc(), diag::note_previous_attribute);
3251+
return nullptr;
3252+
}
3253+
// If all of the results are known to be the same, we can silently drop the
3254+
// attribute. Otherwise, we have to add the attribute and resolve its
3255+
// differences later.
3256+
if (llvm::all_of(Results,
3257+
[](DupArgResult V) { return V == DupArgResult::Same; }))
3258+
return nullptr;
3259+
}
3260+
return ::new (Context)
3261+
WorkGroupSizeHintAttr(Context, A, A.getXDim(), A.getYDim(), A.getZDim());
3262+
}
3263+
3264+
// Handles work_group_size_hint.
3265+
static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
3266+
S.CheckDeprecatedSYCLAttributeSpelling(AL);
31853267

3186-
D->addAttr(::new (S.Context) WorkGroupSizeHintAttr(S.Context, AL, WGSize[0],
3187-
WGSize[1], WGSize[2]));
3268+
S.AddWorkGroupSizeHintAttr(D, AL, AL.getArgAsExpr(0), AL.getArgAsExpr(1),
3269+
AL.getArgAsExpr(2));
31883270
}
31893271

31903272
void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -743,6 +743,25 @@ static void instantiateSYCLIntelESimdVectorizeAttr(
743743
S.AddSYCLIntelESimdVectorizeAttr(New, *A, Result.getAs<Expr>());
744744
}
745745

746+
static void instantiateWorkGroupSizeHintAttr(
747+
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
748+
const WorkGroupSizeHintAttr *A, Decl *New) {
749+
EnterExpressionEvaluationContext Unevaluated(
750+
S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
751+
ExprResult XResult = S.SubstExpr(A->getXDim(), TemplateArgs);
752+
if (XResult.isInvalid())
753+
return;
754+
ExprResult YResult = S.SubstExpr(A->getYDim(), TemplateArgs);
755+
if (YResult.isInvalid())
756+
return;
757+
ExprResult ZResult = S.SubstExpr(A->getZDim(), TemplateArgs);
758+
if (ZResult.isInvalid())
759+
return;
760+
761+
S.AddWorkGroupSizeHintAttr(New, *A, XResult.get(), YResult.get(),
762+
ZResult.get());
763+
}
764+
746765
/// Determine whether the attribute A might be relevent to the declaration D.
747766
/// If not, we can skip instantiating it. The attribute may or may not have
748767
/// been instantiated yet.
@@ -986,6 +1005,10 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
9861005
SYCLIntelESimdVectorize, New);
9871006
continue;
9881007
}
1008+
if (const auto *A = dyn_cast<WorkGroupSizeHintAttr>(TmplAttr)) {
1009+
instantiateWorkGroupSizeHintAttr(*this, TemplateArgs, A, New);
1010+
continue;
1011+
}
9891012
// Existing DLL attribute on the instantiation takes precedence.
9901013
if (TmplAttr->getKind() == attr::DLLExport ||
9911014
TmplAttr->getKind() == attr::DLLImport) {

‎clang/test/SemaOpenCL/invalid-kernel-attrs.cl

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,7 @@ kernel __attribute__((vec_type_hint(int))) __attribute__((vec_type_hint(float)))
1212

1313
kernel __attribute__((work_group_size_hint(8,16,32,4))) void kernel6() {} //expected-error{{'work_group_size_hint' attribute requires exactly 3 arguments}}
1414

15-
kernel __attribute__((work_group_size_hint(1,2,3))) __attribute__((work_group_size_hint(3,2,1))) void kernel7() {} //expected-warning{{attribute 'work_group_size_hint' is already applied with different arguments}}
15+
kernel __attribute__((work_group_size_hint(1,2,3))) __attribute__((work_group_size_hint(3,2,1))) void kernel7() {} //expected-warning{{attribute 'work_group_size_hint' is already applied with different arguments}} expected-note {{previous attribute is here}}
1616

1717
__attribute__((reqd_work_group_size(8,16,32))) void kernel8(){} // expected-error {{attribute 'reqd_work_group_size' can only be applied to an OpenCL kernel}}
1818

@@ -42,7 +42,7 @@ kernel __attribute__((intel_reqd_sub_group_size(-1))) void kernel16() {} // expe
4242
kernel __attribute__((intel_reqd_sub_group_size(8))) __attribute__((intel_reqd_sub_group_size(16))) void kernel17() {} //expected-warning{{attribute 'intel_reqd_sub_group_size' is already applied with different arguments}} \
4343
// expected-note {{previous attribute is here}}
4444

45-
__kernel __attribute__((work_group_size_hint(8,-16,32))) void neg1() {} //expected-error{{'work_group_size_hint' attribute requires a non-negative integral compile time constant expression}}
45+
__kernel __attribute__((work_group_size_hint(8,-16,32))) void neg1() {} //expected-error{{'work_group_size_hint' attribute requires a positive integral compile time constant expression}}
4646
__kernel __attribute__((reqd_work_group_size(8, 16, -32))) void neg2() {} //expected-warning{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
4747

4848
// 4294967294 is a negative integer if treated as signed.

‎clang/test/SemaSYCL/work-group-size-hint.cpp

Lines changed: 38 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -3,13 +3,12 @@
33
// Check the basics.
44
[[sycl::work_group_size_hint]] void f0(); // expected-error {{'work_group_size_hint' attribute requires exactly 3 arguments}}
55
[[sycl::work_group_size_hint(12, 12, 12, 12)]] void f1(); // expected-error {{'work_group_size_hint' attribute requires exactly 3 arguments}}
6-
[[sycl::work_group_size_hint("derp", 1, 2)]] void f2(); // expected-error {{'work_group_size_hint' attribute requires parameter 0 to be an integer constant}}
6+
[[sycl::work_group_size_hint("derp", 1, 2)]] void f2(); // expected-error {{integral constant expression must have integral or unscoped enumeration type, not 'const char [5]'}}
77
[[sycl::work_group_size_hint(1, 1, 1)]] int i; // expected-error {{'work_group_size_hint' attribute only applies to functions}}
88

9-
// FIXME: this should produce a conflicting attribute warning but doesn't. It
10-
// is missing a merge method (and is also missing template instantiation logic).
11-
[[sycl::work_group_size_hint(4, 1, 1)]] void f3();
12-
[[sycl::work_group_size_hint(32, 1, 1)]] void f3() {}
9+
// Produce a conflicting attribute warning when the args are different.
10+
[[sycl::work_group_size_hint(4, 1, 1)]] void f3(); // expected-note {{previous attribute is here}}
11+
[[sycl::work_group_size_hint(32, 1, 1)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
1312

1413
// FIXME: the attribute is like reqd_work_group_size in that it has a one, two,
1514
// and three arg form that needs to be supported.
@@ -24,11 +23,43 @@ __attribute__((work_group_size_hint(4, 1, 1))) void f6(); // expected-warning {{
2423

2524
// Catch the easy case where the attributes are all specified at once with
2625
// different arguments.
27-
[[sycl::work_group_size_hint(4, 1, 1), sycl::work_group_size_hint(32, 1, 1)]] void f7(); // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
26+
[[sycl::work_group_size_hint(4, 1, 1), sycl::work_group_size_hint(32, 1, 1)]] void f7(); // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} expected-note {{previous attribute is here}}
2827

2928
// Show that the attribute works on member functions.
3029
class Functor {
3130
public:
3231
[[sycl::work_group_size_hint(16, 1, 1)]] [[sycl::work_group_size_hint(16, 1, 1)]] void operator()() const;
33-
[[sycl::work_group_size_hint(16, 1, 1)]] [[sycl::work_group_size_hint(32, 1, 1)]] void operator()(int) const; // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
32+
[[sycl::work_group_size_hint(16, 1, 1)]] [[sycl::work_group_size_hint(32, 1, 1)]] void operator()(int) const; // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} expected-note {{previous attribute is here}}
3433
};
34+
35+
// Ensure that template arguments behave appropriately based on instantiations.
36+
template <int N>
37+
[[sycl::work_group_size_hint(N, 1, 1)]] void f8(); // #f8
38+
39+
// Test that template redeclarations also get diagnosed properly.
40+
template <int X, int Y, int Z>
41+
[[sycl::work_group_size_hint(1, 1, 1)]] void f9(); // #f9prev
42+
43+
template <int X, int Y, int Z>
44+
[[sycl::work_group_size_hint(X, Y, Z)]] void f9() {} // #f9
45+
46+
// Test that a template redeclaration where the difference is known up front is
47+
// diagnosed immediately, even without instantiation.
48+
template <int X, int Y, int Z>
49+
[[sycl::work_group_size_hint(X, 1, Z)]] void f10(); // expected-note {{previous attribute is here}}
50+
template <int X, int Y, int Z>
51+
[[sycl::work_group_size_hint(X, 2, Z)]] void f10(); // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
52+
53+
void instantiate() {
54+
f8<1>(); // OK
55+
// expected-error@#f8 {{'work_group_size_hint' attribute requires a positive integral compile time constant expression}}
56+
f8<-1>(); // expected-note {{in instantiation}}
57+
// expected-error@#f8 {{'work_group_size_hint' attribute requires a positive integral compile time constant expression}}
58+
f8<0>(); // expected-note {{in instantiation}}
59+
60+
f9<1, 1, 1>(); // OK, args are the same on the redecl.
61+
62+
// expected-warning@#f9 {{attribute 'work_group_size_hint' is already applied with different arguments}}
63+
// expected-note@#f9prev {{previous attribute is here}}
64+
f9<1, 2, 3>(); // expected-note {{in instantiation}}
65+
}

0 commit comments

Comments
 (0)