Skip to content

[SYCL] Change sycl::reqd_work_group_size with optional dimensions #7450

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 5 commits into from
Dec 15, 2022
Merged

[SYCL] Change sycl::reqd_work_group_size with optional dimensions #7450

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
3b673fb
[SYCL] Change sycl::reqd_work_group_size with optional dimensions
steffenlarsen Nov 18, 2022
dde66cf
Respond to comments
steffenlarsen Nov 21, 2022
1942f22
Split ReqdWorkGroupSizeAttr by language options
steffenlarsen Nov 21, 2022
98414e7
Fix clang-tidy and minor refactoring
steffenlarsen Nov 22, 2022
f5e44b5
Remove wrong not in comments
steffenlarsen Dec 1, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,8 +57,8 @@ void SingleWorkItemBarrierCheck::check(const MatchFinder::MatchResult &Result) {
bool IsNDRange = false;
if (MatchedDecl->hasAttr<ReqdWorkGroupSizeAttr>()) {
const auto *Attribute = MatchedDecl->getAttr<ReqdWorkGroupSizeAttr>();
if (*Attribute->getXDimVal() > 1 || *Attribute->getYDimVal() > 1 ||
*Attribute->getZDimVal() > 1)
if (Attribute->getXDim() > 1 || Attribute->getYDim() > 1 ||
Attribute->getZDim() > 1)
IsNDRange = true;
}
if (IsNDRange) // No warning if kernel is treated as an NDRange.
Expand Down
48 changes: 45 additions & 3 deletions clang/include/clang/Basic/Attr.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -370,6 +370,7 @@ def SYCL : LangOpt<"SYCL">;
def SYCLIsDevice : LangOpt<"SYCLIsDevice">;
def SYCLIsHost : LangOpt<"SYCLIsHost">;
def SilentlyIgnoreSYCLIsHost : LangOpt<"SYCLIsHost", "", 1>;
def NotSYCL : LangOpt<"", "!LangOpts.SYCLIsDevice && !LangOpts.SYCLIsHost">;
def COnly : LangOpt<"", "!LangOpts.CPlusPlus">;
def CPlusPlus : LangOpt<"CPlusPlus">;
def OpenCL : LangOpt<"OpenCL">;
Expand Down Expand Up @@ -665,6 +666,27 @@ class TargetSpecificAttr<TargetSpec target> {
string ParseKind;
}

/// A language-option-specific attribute. This class is meant to be used as a
/// mixin with InheritableAttr or Attr depending on the attribute's needs.
class LanguageOptionsSpecificAttr {
// Attributes are generally required to have unique spellings for their names
// so that the parser can determine what kind of attribute it has parsed.
// However, language-option-specific attributes are special as they have
// different semantics based on the language options specified. To support
// this, a Kind can be explicitly specified for a language-option-specific
// attribute. This corresponds to the ParsedAttr::AT_* enum that is generated
// and it should contain a shared value between the attributes.
// The language options these attributes are unique for are specified in the
// LangOpts member of Attr.
//
// Language-option-specific attributes which use this feature should ensure
// that the spellings match exactly between the attributes, and if the
// arguments or subjects differ, should specify HasCustomParsing = 1 and
// implement their own parsing and semantic handling requirements as-needed.
// Additionally, they should ensure that the language options do not overlap.
string ParseKind;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should this class hold the language option used to distinguish the attribute variants? It seems a bit strange to me that we can use this mixin but never specify a language option (which is different from target-specific attributes where you have to specify the target).

That said, it might be confusing for this class and Attr to both accept a language option, so we might want tablegen to yell at you if you do something wrong. e.g., if we leave the design as-is, maybe tablegen should bark if you don't specify the language option in the attribute definition, and if we switch the design up, maybe tablegen should bark if you specify the language option in two places (once here and once on the attribute).

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The first design I tried had a new LangOpts-like member as part of this class, but I decided for the current design because it did feel redundant when we already had LangOpts in Attr, which the attributes would be deriving from anyway. As such, I prefer to keep it as is and have checks for LangOpts being set in the attributes deriving from it. See #7947.

}

/// An inheritable parameter attribute is inherited by later
/// redeclarations, even when it's written on a parameter.
class InheritableParamAttr : InheritableAttr;
Expand Down Expand Up @@ -3419,33 +3441,53 @@ def NoDeref : TypeAttr {

// Default arguments can only be used with the sycl::reqd_work_group_size
// spelling.
def ReqdWorkGroupSize : InheritableAttr {
def ReqdWorkGroupSize : InheritableAttr, LanguageOptionsSpecificAttr {
let Spellings = [GNU<"reqd_work_group_size">,
CXX11<"cl", "reqd_work_group_size">,
CXX11<"sycl", "reqd_work_group_size">];
let Args = [UnsignedArgument<"XDim">, UnsignedArgument<"YDim">,
UnsignedArgument<"ZDim">];
let Subjects = SubjectList<[Function], ErrorDiag>;
let LangOpts = [NotSYCL];
let Documentation = [ReqdWorkGroupSizeAttrDocs];
let SupportsNonconformingLambdaSyntax = 1;
let ParseKind = "ReqdWorkGroupSize";
let HasCustomParsing = 1;
}

def SYCLReqdWorkGroupSize : InheritableAttr, LanguageOptionsSpecificAttr {
let Spellings = [GNU<"reqd_work_group_size">,
CXX11<"cl", "reqd_work_group_size">,
CXX11<"sycl", "reqd_work_group_size">];
let Args = [ExprArgument<"XDim">,
ExprArgument<"YDim", /*optional*/1>,
ExprArgument<"ZDim", /*optional*/1>];
let Subjects = SubjectList<[Function], ErrorDiag>;
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let AdditionalMembers = [{
Optional<llvm::APSInt> getXDimVal() const {
// X-dimension is not optional.
if (const auto *CE = dyn_cast<ConstantExpr>(getXDim()))
return CE->getResultAsAPSInt();
return None;
}
Optional<llvm::APSInt> getYDimVal() const {
if (const auto *CE = dyn_cast<ConstantExpr>(getYDim()))
// Y-dimension is optional so a nullptr value is allowed.
if (const auto *CE = dyn_cast_or_null<ConstantExpr>(getYDim()))
return CE->getResultAsAPSInt();
return None;
}
Optional<llvm::APSInt> getZDimVal() const {
if (const auto *CE = dyn_cast<ConstantExpr>(getZDim()))
// Z-dimension is optional so a nullptr value is allowed.
if (const auto *CE = dyn_cast_or_null<ConstantExpr>(getZDim()))
return CE->getResultAsAPSInt();
return None;
}
}];
let Documentation = [ReqdWorkGroupSizeAttrDocs];
let SupportsNonconformingLambdaSyntax = 1;
let ParseKind = "ReqdWorkGroupSize";
let HasCustomParsing = 1;
}

def WorkGroupSizeHint : InheritableAttr {
Expand Down
17 changes: 13 additions & 4 deletions clang/include/clang/Sema/Sema.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10975,6 +10975,12 @@ class Sema final {

void AddIntelFPGABankBitsAttr(Decl *D, const AttributeCommonInfo &CI,
Expr **Exprs, unsigned Size);
bool AnyWorkGroupSizesDiffer(const Expr *LHSXDim, const Expr *LHSYDim,
const Expr *LHSZDim, const Expr *RHSXDim,
const Expr *RHSYDim, const Expr *RHSZDim);
bool AllWorkGroupSizesSame(const Expr *LHSXDim, const Expr *LHSYDim,
const Expr *LHSZDim, const Expr *RHSXDim,
const Expr *RHSYDim, const Expr *RHSZDim);
void AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
Expr *XDim, Expr *YDim, Expr *ZDim);
WorkGroupSizeHintAttr *
Expand Down Expand Up @@ -11049,6 +11055,9 @@ class Sema final {
const SYCLUsesAspectsAttr &A);
void AddSYCLUsesAspectsAttr(Decl *D, const AttributeCommonInfo &CI,
Expr **Exprs, unsigned Size);
bool CheckMaxAllowedWorkGroupSize(const Expr *RWGSXDim, const Expr *RWGSYDim,
const Expr *RWGSZDim, const Expr *MWGSXDim,
const Expr *MWGSYDim, const Expr *MWGSZDim);
void AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI,
Expr *XDim, Expr *YDim, Expr *ZDim);
SYCLIntelMaxWorkGroupSizeAttr *
Expand Down Expand Up @@ -11077,10 +11086,10 @@ class Sema final {
const SYCLAddIRAnnotationsMemberAttr &A);
void AddSYCLAddIRAnnotationsMemberAttr(Decl *D, const AttributeCommonInfo &CI,
MutableArrayRef<Expr *> Args);
void AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI,
Expr *XDim, Expr *YDim, Expr *ZDim);
ReqdWorkGroupSizeAttr *
MergeReqdWorkGroupSizeAttr(Decl *D, const ReqdWorkGroupSizeAttr &A);
void AddSYCLReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI,
Expr *XDim, Expr *YDim, Expr *ZDim);
SYCLReqdWorkGroupSizeAttr *
MergeSYCLReqdWorkGroupSizeAttr(Decl *D, const SYCLReqdWorkGroupSizeAttr &A);

SYCLTypeAttr *MergeSYCLTypeAttr(Decl *D, const AttributeCommonInfo &CI,
SYCLTypeAttr::SYCLType TypeName);
Expand Down
30 changes: 24 additions & 6 deletions clang/lib/CodeGen/CodeGenFunction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -619,13 +619,31 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
}

if (const ReqdWorkGroupSizeAttr *A = FD->getAttr<ReqdWorkGroupSizeAttr>()) {
// Attributes arguments (first and third) are reversed on SYCLDevice.
llvm::Metadata *AttrMDArgs[] = {
llvm::ConstantAsMetadata::get(Builder.getInt(
getLangOpts().SYCLIsDevice ? *A->getZDimVal() : *A->getXDimVal())),
llvm::ConstantAsMetadata::get(Builder.getInt(*A->getYDimVal())),
llvm::ConstantAsMetadata::get(Builder.getInt(
getLangOpts().SYCLIsDevice ? *A->getXDimVal() : *A->getZDimVal()))};
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())),
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())),
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))};
Fn->setMetadata("reqd_work_group_size",
llvm::MDNode::get(Context, AttrMDArgs));
}

if (const SYCLReqdWorkGroupSizeAttr *A =
FD->getAttr<SYCLReqdWorkGroupSizeAttr>()) {
llvm::Optional<llvm::APSInt> XDimVal = A->getXDimVal();
llvm::Optional<llvm::APSInt> YDimVal = A->getYDimVal();
llvm::Optional<llvm::APSInt> ZDimVal = A->getZDimVal();
llvm::SmallVector<llvm::Metadata *, 3> AttrMDArgs;

// On SYCL target the dimensions are reversed if present.
if (ZDimVal)
AttrMDArgs.push_back(
llvm::ConstantAsMetadata::get(Builder.getInt(*ZDimVal)));
if (YDimVal)
AttrMDArgs.push_back(
llvm::ConstantAsMetadata::get(Builder.getInt(*YDimVal)));
AttrMDArgs.push_back(
llvm::ConstantAsMetadata::get(Builder.getInt(*XDimVal)));

Fn->setMetadata("reqd_work_group_size",
llvm::MDNode::get(Context, AttrMDArgs));
}
Expand Down
37 changes: 16 additions & 21 deletions clang/lib/CodeGen/TargetInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8496,16 +8496,16 @@ void TCETargetCodeGenInfo::setTargetAttributes(

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

Operands.push_back(llvm::ConstantAsMetadata::get(
llvm::Constant::getIntegerValue(M.Int32Ty, llvm::APInt(32, XDim))));
Operands.push_back(llvm::ConstantAsMetadata::get(
llvm::Constant::getIntegerValue(M.Int32Ty, llvm::APInt(32, YDim))));
Operands.push_back(llvm::ConstantAsMetadata::get(
llvm::Constant::getIntegerValue(M.Int32Ty, llvm::APInt(32, ZDim))));
Operands.push_back(
llvm::ConstantAsMetadata::get(llvm::Constant::getIntegerValue(
M.Int32Ty, llvm::APInt(32, Attr->getXDim()))));
Operands.push_back(
llvm::ConstantAsMetadata::get(llvm::Constant::getIntegerValue(
M.Int32Ty, llvm::APInt(32, Attr->getYDim()))));
Operands.push_back(
llvm::ConstantAsMetadata::get(llvm::Constant::getIntegerValue(
M.Int32Ty, llvm::APInt(32, Attr->getZDim()))));

// Add a boolean constant operand for "required" (true) or "hint"
// (false) for implementing the work_group_size_hint attr later.
Expand Down Expand Up @@ -9380,21 +9380,16 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
if (ReqdWGS || FlatWGS) {
unsigned Min = 0;
unsigned Max = 0;
unsigned XDim = 0;
unsigned YDim = 0;
unsigned ZDim = 0;
ASTContext &Ctx = M.getContext();
if (FlatWGS) {
Min = FlatWGS->getMin()->EvaluateKnownConstInt(Ctx).getExtValue();
Max = FlatWGS->getMax()->EvaluateKnownConstInt(Ctx).getExtValue();
}
if (ReqdWGS) {
XDim = ReqdWGS->getXDimVal()->getZExtValue();
YDim = ReqdWGS->getYDimVal()->getZExtValue();
ZDim = ReqdWGS->getZDimVal()->getZExtValue();
Min = FlatWGS->getMin()
->EvaluateKnownConstInt(M.getContext())
.getExtValue();
Max = FlatWGS->getMax()
->EvaluateKnownConstInt(M.getContext())
.getExtValue();
}
if (ReqdWGS && Min == 0 && Max == 0)
Min = Max = XDim * YDim * ZDim;
Min = Max = ReqdWGS->getXDim() * ReqdWGS->getYDim() * ReqdWGS->getZDim();

if (Min != 0) {
assert(Min <= Max && "Min must be less than or equal Max");
Expand Down
4 changes: 2 additions & 2 deletions clang/lib/Sema/SemaDecl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3000,8 +3000,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
NewAttr = S.MergeSYCLAddIRAttributesGlobalVariableAttr(D, *A);
else if (const auto *A = dyn_cast<SYCLAddIRAnnotationsMemberAttr>(Attr))
NewAttr = S.MergeSYCLAddIRAnnotationsMemberAttr(D, *A);
else if (const auto *A = dyn_cast<ReqdWorkGroupSizeAttr>(Attr))
NewAttr = S.MergeReqdWorkGroupSizeAttr(D, *A);
else if (const auto *A = dyn_cast<SYCLReqdWorkGroupSizeAttr>(Attr))
NewAttr = S.MergeSYCLReqdWorkGroupSizeAttr(D, *A);
else if (const auto *NT = dyn_cast<HLSLNumThreadsAttr>(Attr))
NewAttr =
S.mergeHLSLNumThreadsAttr(D, *NT, NT->getX(), NT->getY(), NT->getZ());
Expand Down
Loading