Skip to content

Commit b211d0d

Browse files
[SYCL] Change sycl::reqd_work_group_size with optional dimensions (#7450)
In the current implementation the sycl::reqd_work_group_size attribute sets the Y and Z dimension arguments optional. However, when the internal representation of the attribute is created it will be padded with 1's in the additional dimensions. An effect of this padding is that dimensionality information is lost, which has three big drawbacks: 1. SYCL work-group sizes are reversed, but only the specified dimensions are reversed. For example `sycl::reqd_work_group_size(1, 2, 3)` is for the backends the same as a backend work-group of size `<3,2,1>`, but `sycl::reqd_work_group_size(3)` corresonds to a backend work-group of `<3,1,1>` rather than `<1,1,3>`. 2. The SYCL runtime is supposed to throw an exception when a kernel is launched with a number of dimensions that does not match the `sycl::reqd_work_group_size`'s dimensionality. sycl-post-link generates kernel meta-information for the runtime which could be used to diagnose these, but since the attribute is padded with 1's the `reqd_work_group_size` metadata node knows no difference between it and a user-specified attribute with trailing 1's. 3. Sema cannot know the difference between two attributes where one was padded with 1's by the user and one was not, so it currently thinks these are equivalent. To fix these, this patch changes Sema to not add the padding and instead consider cases where the Y and Z dimensions are unset. This only affects the SYCL spelling of the attribute. Additionally, when generating the `reqd_work_group_size` attribute CodeGen will only generate a metadata value for dimensions that have been set.
1 parent bbb9009 commit b211d0d

23 files changed

+637
-382
lines changed

clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -57,8 +57,8 @@ void SingleWorkItemBarrierCheck::check(const MatchFinder::MatchResult &Result) {
5757
bool IsNDRange = false;
5858
if (MatchedDecl->hasAttr<ReqdWorkGroupSizeAttr>()) {
5959
const auto *Attribute = MatchedDecl->getAttr<ReqdWorkGroupSizeAttr>();
60-
if (*Attribute->getXDimVal() > 1 || *Attribute->getYDimVal() > 1 ||
61-
*Attribute->getZDimVal() > 1)
60+
if (Attribute->getXDim() > 1 || Attribute->getYDim() > 1 ||
61+
Attribute->getZDim() > 1)
6262
IsNDRange = true;
6363
}
6464
if (IsNDRange) // No warning if kernel is treated as an NDRange.

clang/include/clang/Basic/Attr.td

Lines changed: 45 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -370,6 +370,7 @@ def SYCL : LangOpt<"SYCL">;
370370
def SYCLIsDevice : LangOpt<"SYCLIsDevice">;
371371
def SYCLIsHost : LangOpt<"SYCLIsHost">;
372372
def SilentlyIgnoreSYCLIsHost : LangOpt<"SYCLIsHost", "", 1>;
373+
def NotSYCL : LangOpt<"", "!LangOpts.SYCLIsDevice && !LangOpts.SYCLIsHost">;
373374
def COnly : LangOpt<"", "!LangOpts.CPlusPlus">;
374375
def CPlusPlus : LangOpt<"CPlusPlus">;
375376
def OpenCL : LangOpt<"OpenCL">;
@@ -665,6 +666,27 @@ class TargetSpecificAttr<TargetSpec target> {
665666
string ParseKind;
666667
}
667668

669+
/// A language-option-specific attribute. This class is meant to be used as a
670+
/// mixin with InheritableAttr or Attr depending on the attribute's needs.
671+
class LanguageOptionsSpecificAttr {
672+
// Attributes are generally required to have unique spellings for their names
673+
// so that the parser can determine what kind of attribute it has parsed.
674+
// However, language-option-specific attributes are special as they have
675+
// different semantics based on the language options specified. To support
676+
// this, a Kind can be explicitly specified for a language-option-specific
677+
// attribute. This corresponds to the ParsedAttr::AT_* enum that is generated
678+
// and it should contain a shared value between the attributes.
679+
// The language options these attributes are unique for are specified in the
680+
// LangOpts member of Attr.
681+
//
682+
// Language-option-specific attributes which use this feature should ensure
683+
// that the spellings match exactly between the attributes, and if the
684+
// arguments or subjects differ, should specify HasCustomParsing = 1 and
685+
// implement their own parsing and semantic handling requirements as-needed.
686+
// Additionally, they should ensure that the language options do not overlap.
687+
string ParseKind;
688+
}
689+
668690
/// An inheritable parameter attribute is inherited by later
669691
/// redeclarations, even when it's written on a parameter.
670692
class InheritableParamAttr : InheritableAttr;
@@ -3426,33 +3448,53 @@ def NoDeref : TypeAttr {
34263448

34273449
// Default arguments can only be used with the sycl::reqd_work_group_size
34283450
// spelling.
3429-
def ReqdWorkGroupSize : InheritableAttr {
3451+
def ReqdWorkGroupSize : InheritableAttr, LanguageOptionsSpecificAttr {
3452+
let Spellings = [GNU<"reqd_work_group_size">,
3453+
CXX11<"cl", "reqd_work_group_size">,
3454+
CXX11<"sycl", "reqd_work_group_size">];
3455+
let Args = [UnsignedArgument<"XDim">, UnsignedArgument<"YDim">,
3456+
UnsignedArgument<"ZDim">];
3457+
let Subjects = SubjectList<[Function], ErrorDiag>;
3458+
let LangOpts = [NotSYCL];
3459+
let Documentation = [ReqdWorkGroupSizeAttrDocs];
3460+
let SupportsNonconformingLambdaSyntax = 1;
3461+
let ParseKind = "ReqdWorkGroupSize";
3462+
let HasCustomParsing = 1;
3463+
}
3464+
3465+
def SYCLReqdWorkGroupSize : InheritableAttr, LanguageOptionsSpecificAttr {
34303466
let Spellings = [GNU<"reqd_work_group_size">,
34313467
CXX11<"cl", "reqd_work_group_size">,
34323468
CXX11<"sycl", "reqd_work_group_size">];
34333469
let Args = [ExprArgument<"XDim">,
34343470
ExprArgument<"YDim", /*optional*/1>,
34353471
ExprArgument<"ZDim", /*optional*/1>];
34363472
let Subjects = SubjectList<[Function], ErrorDiag>;
3473+
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
34373474
let AdditionalMembers = [{
34383475
Optional<llvm::APSInt> getXDimVal() const {
3476+
// X-dimension is not optional.
34393477
if (const auto *CE = dyn_cast<ConstantExpr>(getXDim()))
34403478
return CE->getResultAsAPSInt();
34413479
return None;
34423480
}
34433481
Optional<llvm::APSInt> getYDimVal() const {
3444-
if (const auto *CE = dyn_cast<ConstantExpr>(getYDim()))
3482+
// Y-dimension is optional so a nullptr value is allowed.
3483+
if (const auto *CE = dyn_cast_or_null<ConstantExpr>(getYDim()))
34453484
return CE->getResultAsAPSInt();
34463485
return None;
34473486
}
34483487
Optional<llvm::APSInt> getZDimVal() const {
3449-
if (const auto *CE = dyn_cast<ConstantExpr>(getZDim()))
3488+
// Z-dimension is optional so a nullptr value is allowed.
3489+
if (const auto *CE = dyn_cast_or_null<ConstantExpr>(getZDim()))
34503490
return CE->getResultAsAPSInt();
34513491
return None;
34523492
}
34533493
}];
34543494
let Documentation = [ReqdWorkGroupSizeAttrDocs];
34553495
let SupportsNonconformingLambdaSyntax = 1;
3496+
let ParseKind = "ReqdWorkGroupSize";
3497+
let HasCustomParsing = 1;
34563498
}
34573499

34583500
def WorkGroupSizeHint : InheritableAttr {

clang/include/clang/Sema/Sema.h

Lines changed: 13 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -10990,6 +10990,12 @@ class Sema final {
1099010990

1099110991
void AddSYCLIntelBankBitsAttr(Decl *D, const AttributeCommonInfo &CI,
1099210992
Expr **Exprs, unsigned Size);
10993+
bool AnyWorkGroupSizesDiffer(const Expr *LHSXDim, const Expr *LHSYDim,
10994+
const Expr *LHSZDim, const Expr *RHSXDim,
10995+
const Expr *RHSYDim, const Expr *RHSZDim);
10996+
bool AllWorkGroupSizesSame(const Expr *LHSXDim, const Expr *LHSYDim,
10997+
const Expr *LHSZDim, const Expr *RHSXDim,
10998+
const Expr *RHSYDim, const Expr *RHSZDim);
1099310999
void AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
1099411000
Expr *XDim, Expr *YDim, Expr *ZDim);
1099511001
WorkGroupSizeHintAttr *
@@ -11064,6 +11070,9 @@ class Sema final {
1106411070
const SYCLUsesAspectsAttr &A);
1106511071
void AddSYCLUsesAspectsAttr(Decl *D, const AttributeCommonInfo &CI,
1106611072
Expr **Exprs, unsigned Size);
11073+
bool CheckMaxAllowedWorkGroupSize(const Expr *RWGSXDim, const Expr *RWGSYDim,
11074+
const Expr *RWGSZDim, const Expr *MWGSXDim,
11075+
const Expr *MWGSYDim, const Expr *MWGSZDim);
1106711076
void AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI,
1106811077
Expr *XDim, Expr *YDim, Expr *ZDim);
1106911078
SYCLIntelMaxWorkGroupSizeAttr *
@@ -11092,10 +11101,10 @@ class Sema final {
1109211101
const SYCLAddIRAnnotationsMemberAttr &A);
1109311102
void AddSYCLAddIRAnnotationsMemberAttr(Decl *D, const AttributeCommonInfo &CI,
1109411103
MutableArrayRef<Expr *> Args);
11095-
void AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI,
11096-
Expr *XDim, Expr *YDim, Expr *ZDim);
11097-
ReqdWorkGroupSizeAttr *
11098-
MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A);
11104+
void AddSYCLReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI,
11105+
Expr *XDim, Expr *YDim, Expr *ZDim);
11106+
SYCLReqdWorkGroupSizeAttr *
11107+
MergeSYCLReqdWorkGroupSizeAttr(Decl *D, const SYCLReqdWorkGroupSizeAttr &A);
1109911108

1110011109
SYCLTypeAttr *MergeSYCLTypeAttr(Decl *D, const AttributeCommonInfo &CI,
1110111110
SYCLTypeAttr::SYCLType TypeName);

clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 24 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -620,13 +620,31 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
620620
}
621621

622622
if (const ReqdWorkGroupSizeAttr *A = FD->getAttr<ReqdWorkGroupSizeAttr>()) {
623-
// Attributes arguments (first and third) are reversed on SYCLDevice.
624623
llvm::Metadata *AttrMDArgs[] = {
625-
llvm::ConstantAsMetadata::get(Builder.getInt(
626-
getLangOpts().SYCLIsDevice ? *A->getZDimVal() : *A->getXDimVal())),
627-
llvm::ConstantAsMetadata::get(Builder.getInt(*A->getYDimVal())),
628-
llvm::ConstantAsMetadata::get(Builder.getInt(
629-
getLangOpts().SYCLIsDevice ? *A->getXDimVal() : *A->getZDimVal()))};
624+
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())),
625+
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())),
626+
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))};
627+
Fn->setMetadata("reqd_work_group_size",
628+
llvm::MDNode::get(Context, AttrMDArgs));
629+
}
630+
631+
if (const SYCLReqdWorkGroupSizeAttr *A =
632+
FD->getAttr<SYCLReqdWorkGroupSizeAttr>()) {
633+
llvm::Optional<llvm::APSInt> XDimVal = A->getXDimVal();
634+
llvm::Optional<llvm::APSInt> YDimVal = A->getYDimVal();
635+
llvm::Optional<llvm::APSInt> ZDimVal = A->getZDimVal();
636+
llvm::SmallVector<llvm::Metadata *, 3> AttrMDArgs;
637+
638+
// On SYCL target the dimensions are reversed if present.
639+
if (ZDimVal)
640+
AttrMDArgs.push_back(
641+
llvm::ConstantAsMetadata::get(Builder.getInt(*ZDimVal)));
642+
if (YDimVal)
643+
AttrMDArgs.push_back(
644+
llvm::ConstantAsMetadata::get(Builder.getInt(*YDimVal)));
645+
AttrMDArgs.push_back(
646+
llvm::ConstantAsMetadata::get(Builder.getInt(*XDimVal)));
647+
630648
Fn->setMetadata("reqd_work_group_size",
631649
llvm::MDNode::get(Context, AttrMDArgs));
632650
}

clang/lib/CodeGen/TargetInfo.cpp

Lines changed: 16 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -8498,16 +8498,16 @@ void TCETargetCodeGenInfo::setTargetAttributes(
84988498

84998499
SmallVector<llvm::Metadata *, 5> Operands;
85008500
Operands.push_back(llvm::ConstantAsMetadata::get(F));
8501-
unsigned XDim = Attr->getXDimVal()->getZExtValue();
8502-
unsigned YDim = Attr->getYDimVal()->getZExtValue();
8503-
unsigned ZDim = Attr->getZDimVal()->getZExtValue();
85048501

8505-
Operands.push_back(llvm::ConstantAsMetadata::get(
8506-
llvm::Constant::getIntegerValue(M.Int32Ty, llvm::APInt(32, XDim))));
8507-
Operands.push_back(llvm::ConstantAsMetadata::get(
8508-
llvm::Constant::getIntegerValue(M.Int32Ty, llvm::APInt(32, YDim))));
8509-
Operands.push_back(llvm::ConstantAsMetadata::get(
8510-
llvm::Constant::getIntegerValue(M.Int32Ty, llvm::APInt(32, ZDim))));
8502+
Operands.push_back(
8503+
llvm::ConstantAsMetadata::get(llvm::Constant::getIntegerValue(
8504+
M.Int32Ty, llvm::APInt(32, Attr->getXDim()))));
8505+
Operands.push_back(
8506+
llvm::ConstantAsMetadata::get(llvm::Constant::getIntegerValue(
8507+
M.Int32Ty, llvm::APInt(32, Attr->getYDim()))));
8508+
Operands.push_back(
8509+
llvm::ConstantAsMetadata::get(llvm::Constant::getIntegerValue(
8510+
M.Int32Ty, llvm::APInt(32, Attr->getZDim()))));
85118511

85128512
// Add a boolean constant operand for "required" (true) or "hint"
85138513
// (false) for implementing the work_group_size_hint attr later.
@@ -9382,21 +9382,16 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
93829382
if (ReqdWGS || FlatWGS) {
93839383
unsigned Min = 0;
93849384
unsigned Max = 0;
9385-
unsigned XDim = 0;
9386-
unsigned YDim = 0;
9387-
unsigned ZDim = 0;
9388-
ASTContext &Ctx = M.getContext();
93899385
if (FlatWGS) {
9390-
Min = FlatWGS->getMin()->EvaluateKnownConstInt(Ctx).getExtValue();
9391-
Max = FlatWGS->getMax()->EvaluateKnownConstInt(Ctx).getExtValue();
9392-
}
9393-
if (ReqdWGS) {
9394-
XDim = ReqdWGS->getXDimVal()->getZExtValue();
9395-
YDim = ReqdWGS->getYDimVal()->getZExtValue();
9396-
ZDim = ReqdWGS->getZDimVal()->getZExtValue();
9386+
Min = FlatWGS->getMin()
9387+
->EvaluateKnownConstInt(M.getContext())
9388+
.getExtValue();
9389+
Max = FlatWGS->getMax()
9390+
->EvaluateKnownConstInt(M.getContext())
9391+
.getExtValue();
93979392
}
93989393
if (ReqdWGS && Min == 0 && Max == 0)
9399-
Min = Max = XDim * YDim * ZDim;
9394+
Min = Max = ReqdWGS->getXDim() * ReqdWGS->getYDim() * ReqdWGS->getZDim();
94009395

94019396
if (Min != 0) {
94029397
assert(Min <= Max && "Min must be less than or equal Max");

clang/lib/Sema/SemaDecl.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3001,8 +3001,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
30013001
NewAttr = S.MergeSYCLAddIRAttributesGlobalVariableAttr(D, *A);
30023002
else if (const auto *A = dyn_cast<SYCLAddIRAnnotationsMemberAttr>(Attr))
30033003
NewAttr = S.MergeSYCLAddIRAnnotationsMemberAttr(D, *A);
3004-
else if (const auto *A = dyn_cast<ReqdWorkGroupSizeAttr>(Attr))
3005-
NewAttr = S.MergeReqdWorkGroupSizeAttr(D, *A);
3004+
else if (const auto *A = dyn_cast<SYCLReqdWorkGroupSizeAttr>(Attr))
3005+
NewAttr = S.MergeSYCLReqdWorkGroupSizeAttr(D, *A);
30063006
else if (const auto *NT = dyn_cast<HLSLNumThreadsAttr>(Attr))
30073007
NewAttr =
30083008
S.mergeHLSLNumThreadsAttr(D, *NT, NT->getX(), NT->getY(), NT->getZ());

0 commit comments

Comments
 (0)