Skip to content

Commit ba51bb6

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Refactor SYCL path for work_group_size_hint (#7961)
This commit refactors the SYCL path for work_group_size_hint, similar to the refactoring for reqd_work_group_size in #7450. It also fixes an issue where the dimensions of the work-group hint were not correctly reversed. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 13adc40 commit ba51bb6

15 files changed

+363
-80
lines changed

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

Lines changed: 24 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -3497,35 +3497,55 @@ def SYCLReqdWorkGroupSize : InheritableAttr, LanguageOptionsSpecificAttr {
34973497
let HasCustomParsing = 1;
34983498
}
34993499

3500-
def WorkGroupSizeHint : InheritableAttr {
3500+
def WorkGroupSizeHint : InheritableAttr, LanguageOptionsSpecificAttr {
3501+
let Spellings = [GNU<"work_group_size_hint">,
3502+
CXX11<"sycl", "work_group_size_hint">];
3503+
let Args = [UnsignedArgument<"XDim">,
3504+
UnsignedArgument<"YDim">,
3505+
UnsignedArgument<"ZDim">];
3506+
let Subjects = SubjectList<[Function], ErrorDiag>;
3507+
let Documentation = [WorkGroupSizeHintAttrDocs];
3508+
let LangOpts = [NotSYCL];
3509+
let SupportsNonconformingLambdaSyntax = 1;
3510+
let ParseKind = "WorkGroupSizeHint";
3511+
let HasCustomParsing = 1;
3512+
}
3513+
3514+
def SYCLWorkGroupSizeHint : InheritableAttr, LanguageOptionsSpecificAttr {
35013515
let Spellings = [GNU<"work_group_size_hint">,
35023516
CXX11<"sycl", "work_group_size_hint">];
35033517
let Args = [ExprArgument<"XDim">,
35043518
ExprArgument<"YDim", /*optional*/1>,
35053519
ExprArgument<"ZDim", /*optional*/1>];
35063520
let Subjects = SubjectList<[Function], ErrorDiag>;
35073521
// In OpenCL C this attribute takes no default values whereas in SYCL it does.
3508-
// To avoid confusing diagnostics, the checks are deferred to "handleWorkGroupSizeHint".
3522+
// To avoid confusing diagnostics, the checks are deferred to
3523+
// "handleWorkGroupSizeHint".
35093524
let HasCustomParsing = 1;
35103525
let AdditionalMembers = [{
35113526
Optional<llvm::APSInt> getXDimVal() const {
3527+
// X-dimension is not optional.
35123528
if (const auto *CE = dyn_cast<ConstantExpr>(getXDim()))
35133529
return CE->getResultAsAPSInt();
35143530
return std::nullopt;
35153531
}
35163532
Optional<llvm::APSInt> getYDimVal() const {
3517-
if (const auto *CE = dyn_cast<ConstantExpr>(getYDim()))
3533+
// Y-dimension is optional so a nullptr value is allowed.
3534+
if (const auto *CE = dyn_cast_or_null<ConstantExpr>(getYDim()))
35183535
return CE->getResultAsAPSInt();
35193536
return std::nullopt;
35203537
}
35213538
Optional<llvm::APSInt> getZDimVal() const {
3522-
if (const auto *CE = dyn_cast<ConstantExpr>(getZDim()))
3539+
// Z-dimension is optional so a nullptr value is allowed.
3540+
if (const auto *CE = dyn_cast_or_null<ConstantExpr>(getZDim()))
35233541
return CE->getResultAsAPSInt();
35243542
return std::nullopt;
35253543
}
35263544
}];
35273545
let Documentation = [WorkGroupSizeHintAttrDocs];
3546+
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
35283547
let SupportsNonconformingLambdaSyntax = 1;
3548+
let ParseKind = "WorkGroupSizeHint";
35293549
}
35303550

35313551
def InitPriority : InheritableAttr, TargetSpecificAttr<TargetSupportsInitPriority> {

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -10997,10 +10997,10 @@ class Sema final {
1099710997
bool AllWorkGroupSizesSame(const Expr *LHSXDim, const Expr *LHSYDim,
1099810998
const Expr *LHSZDim, const Expr *RHSXDim,
1099910999
const Expr *RHSYDim, const Expr *RHSZDim);
11000-
void AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
11001-
Expr *XDim, Expr *YDim, Expr *ZDim);
11002-
WorkGroupSizeHintAttr *
11003-
MergeWorkGroupSizeHintAttr(Decl *D, const WorkGroupSizeHintAttr &A);
11000+
void AddSYCLWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
11001+
Expr *XDim, Expr *YDim, Expr *ZDim);
11002+
SYCLWorkGroupSizeHintAttr *
11003+
MergeSYCLWorkGroupSizeHintAttr(Decl *D, const SYCLWorkGroupSizeHintAttr &A);
1100411004
void AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
1100511005
Expr *E);
1100611006
IntelReqdSubGroupSizeAttr *

‎clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 24 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -613,12 +613,33 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
613613

614614
if (const WorkGroupSizeHintAttr *A = FD->getAttr<WorkGroupSizeHintAttr>()) {
615615
llvm::Metadata *AttrMDArgs[] = {
616-
llvm::ConstantAsMetadata::get(Builder.getInt(*A->getXDimVal())),
617-
llvm::ConstantAsMetadata::get(Builder.getInt(*A->getYDimVal())),
618-
llvm::ConstantAsMetadata::get(Builder.getInt(*A->getZDimVal()))};
616+
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())),
617+
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())),
618+
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))};
619619
Fn->setMetadata("work_group_size_hint", llvm::MDNode::get(Context, AttrMDArgs));
620620
}
621621

622+
if (const SYCLWorkGroupSizeHintAttr *A =
623+
FD->getAttr<SYCLWorkGroupSizeHintAttr>()) {
624+
llvm::Optional<llvm::APSInt> XDimVal = A->getXDimVal();
625+
llvm::Optional<llvm::APSInt> YDimVal = A->getYDimVal();
626+
llvm::Optional<llvm::APSInt> ZDimVal = A->getZDimVal();
627+
llvm::SmallVector<llvm::Metadata *, 3> AttrMDArgs;
628+
629+
// On SYCL target the dimensions are reversed if present.
630+
if (ZDimVal)
631+
AttrMDArgs.push_back(
632+
llvm::ConstantAsMetadata::get(Builder.getInt(*ZDimVal)));
633+
if (YDimVal)
634+
AttrMDArgs.push_back(
635+
llvm::ConstantAsMetadata::get(Builder.getInt(*YDimVal)));
636+
AttrMDArgs.push_back(
637+
llvm::ConstantAsMetadata::get(Builder.getInt(*XDimVal)));
638+
639+
Fn->setMetadata("work_group_size_hint",
640+
llvm::MDNode::get(Context, AttrMDArgs));
641+
}
642+
622643
if (const ReqdWorkGroupSizeAttr *A = FD->getAttr<ReqdWorkGroupSizeAttr>()) {
623644
llvm::Metadata *AttrMDArgs[] = {
624645
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())),

‎clang/lib/Sema/SemaDecl.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2971,8 +2971,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
29712971
NewAttr = S.MergeSYCLIntelForcePow2DepthAttr(D, *A);
29722972
else if (const auto *A = dyn_cast<SYCLIntelInitiationIntervalAttr>(Attr))
29732973
NewAttr = S.MergeSYCLIntelInitiationIntervalAttr(D, *A);
2974-
else if (const auto *A = dyn_cast<WorkGroupSizeHintAttr>(Attr))
2975-
NewAttr = S.MergeWorkGroupSizeHintAttr(D, *A);
2974+
else if (const auto *A = dyn_cast<SYCLWorkGroupSizeHintAttr>(Attr))
2975+
NewAttr = S.MergeSYCLWorkGroupSizeHintAttr(D, *A);
29762976
else if (const auto *A = dyn_cast<SYCLIntelMaxGlobalWorkDimAttr>(Attr))
29772977
NewAttr = S.MergeSYCLIntelMaxGlobalWorkDimAttr(D, *A);
29782978
else if (const auto *BTFA = dyn_cast<BTFDeclTagAttr>(Attr))

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 44 additions & 38 deletions
Original file line numberDiff line numberDiff line change
@@ -3275,6 +3275,9 @@ static void handleWeakImportAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
32753275
// Handles reqd_work_group_size and work_group_size_hint.
32763276
template <typename WorkGroupAttr>
32773277
static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
3278+
if (!AL.checkExactlyNumArgs(S, 3))
3279+
return;
3280+
32783281
uint32_t WGSize[3];
32793282
for (unsigned i = 0; i < 3; ++i) {
32803283
const Expr *E = AL.getArgAsExpr(i);
@@ -3348,24 +3351,24 @@ bool Sema::AllWorkGroupSizesSame(const Expr *LHSXDim, const Expr *LHSYDim,
33483351
[](DupArgResult V) { return V == DupArgResult::Same; });
33493352
}
33503353

3351-
void Sema::AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
3352-
Expr *XDim, Expr *YDim, Expr *ZDim) {
3354+
void Sema::AddSYCLWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
3355+
Expr *XDim, Expr *YDim, Expr *ZDim) {
33533356
// Returns nullptr if diagnosing, otherwise returns the original expression
33543357
// or the original expression converted to a constant expression.
3355-
auto CheckAndConvertArg = [&](Expr *E) -> Expr * {
3358+
auto CheckAndConvertArg = [&](Expr *E) -> Optional<Expr *> {
33563359
// We can only check if the expression is not value dependent.
3357-
if (!E->isValueDependent()) {
3360+
if (E && !E->isValueDependent()) {
33583361
llvm::APSInt ArgVal;
33593362
ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal);
33603363
if (Res.isInvalid())
3361-
return nullptr;
3364+
return std::nullopt;
33623365
E = Res.get();
33633366

33643367
// This attribute requires a strictly positive value.
33653368
if (ArgVal <= 0) {
33663369
Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
33673370
<< CI << /*positive*/ 0;
3368-
return nullptr;
3371+
return std::nullopt;
33693372
}
33703373
}
33713374

@@ -3374,15 +3377,18 @@ void Sema::AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
33743377

33753378
// Check all three argument values, and if any are bad, bail out. This will
33763379
// convert the given expressions into constant expressions when possible.
3377-
XDim = CheckAndConvertArg(XDim);
3378-
YDim = CheckAndConvertArg(YDim);
3379-
ZDim = CheckAndConvertArg(ZDim);
3380-
if (!XDim || !YDim || !ZDim)
3380+
Optional<Expr *> XDimConvert = CheckAndConvertArg(XDim);
3381+
Optional<Expr *> YDimConvert = CheckAndConvertArg(YDim);
3382+
Optional<Expr *> ZDimConvert = CheckAndConvertArg(ZDim);
3383+
if (!XDimConvert || !YDimConvert || !ZDimConvert)
33813384
return;
3385+
XDim = XDimConvert.value();
3386+
YDim = YDimConvert.value();
3387+
ZDim = ZDimConvert.value();
33823388

33833389
// If the attribute was already applied with different arguments, then
33843390
// diagnose the second attribute as a duplicate and don't add it.
3385-
if (const auto *Existing = D->getAttr<WorkGroupSizeHintAttr>()) {
3391+
if (const auto *Existing = D->getAttr<SYCLWorkGroupSizeHintAttr>()) {
33863392
// If any of the results are known to be different, we can diagnose at this
33873393
// point and drop the attribute.
33883394
if (AnyWorkGroupSizesDiffer(XDim, YDim, ZDim, Existing->getXDim(),
@@ -3400,13 +3406,14 @@ void Sema::AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
34003406
}
34013407

34023408
D->addAttr(::new (Context)
3403-
WorkGroupSizeHintAttr(Context, CI, XDim, YDim, ZDim));
3409+
SYCLWorkGroupSizeHintAttr(Context, CI, XDim, YDim, ZDim));
34043410
}
34053411

3406-
WorkGroupSizeHintAttr *
3407-
Sema::MergeWorkGroupSizeHintAttr(Decl *D, const WorkGroupSizeHintAttr &A) {
3412+
SYCLWorkGroupSizeHintAttr *
3413+
Sema::MergeSYCLWorkGroupSizeHintAttr(Decl *D,
3414+
const SYCLWorkGroupSizeHintAttr &A) {
34083415
// Check to see if there's a duplicate attribute already applied.
3409-
if (const auto *DeclAttr = D->getAttr<WorkGroupSizeHintAttr>()) {
3416+
if (const auto *DeclAttr = D->getAttr<SYCLWorkGroupSizeHintAttr>()) {
34103417
// If any of the results are known to be different, we can diagnose at this
34113418
// point and drop the attribute.
34123419
if (AnyWorkGroupSizesDiffer(DeclAttr->getXDim(), DeclAttr->getYDim(),
@@ -3424,40 +3431,36 @@ Sema::MergeWorkGroupSizeHintAttr(Decl *D, const WorkGroupSizeHintAttr &A) {
34243431
A.getZDim()))
34253432
return nullptr;
34263433
}
3427-
return ::new (Context)
3428-
WorkGroupSizeHintAttr(Context, A, A.getXDim(), A.getYDim(), A.getZDim());
3434+
return ::new (Context) SYCLWorkGroupSizeHintAttr(Context, A, A.getXDim(),
3435+
A.getYDim(), A.getZDim());
34293436
}
34303437

3431-
// Handles work_group_size_hint.
3432-
static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
3438+
// Handles SYCL work_group_size_hint.
3439+
static void handleSYCLWorkGroupSizeHint(Sema &S, Decl *D,
3440+
const ParsedAttr &AL) {
34333441
S.CheckDeprecatedSYCLAttributeSpelling(AL);
34343442

34353443
// __attribute__((work_group_size_hint) requires exactly three arguments.
34363444
if (AL.getSyntax() == ParsedAttr::AS_GNU || !AL.hasScope() ||
34373445
(AL.hasScope() && !AL.getScopeName()->isStr("sycl"))) {
34383446
if (!AL.checkExactlyNumArgs(S, 3))
34393447
return;
3440-
}
3441-
3442-
// FIXME: NumArgs checking is disabled in Attr.td to keep consistent
3443-
// disgnostics with OpenCL C that does not have optional values here.
3444-
if (!AL.checkAtLeastNumArgs(S, 1) || !AL.checkAtMostNumArgs(S, 3))
3448+
} else if (!AL.checkAtLeastNumArgs(S, 1) || !AL.checkAtMostNumArgs(S, 3))
34453449
return;
34463450

3447-
// Handles default arguments in [[sycl::work_group_size_hint]] attribute.
3448-
auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL) {
3449-
assert(AL.getKind() == ParsedAttr::AT_WorkGroupSizeHint && AL.hasScope() &&
3450-
AL.getScopeName()->isStr("sycl"));
3451-
return IntegerLiteral::Create(S.Context, llvm::APInt(32, 1),
3452-
S.Context.IntTy, AL.getLoc());
3453-
};
3451+
size_t NumArgs = AL.getNumArgs();
3452+
Expr *XDimExpr = NumArgs > 0 ? AL.getArgAsExpr(0) : nullptr;
3453+
Expr *YDimExpr = NumArgs > 1 ? AL.getArgAsExpr(1) : nullptr;
3454+
Expr *ZDimExpr = NumArgs > 2 ? AL.getArgAsExpr(2) : nullptr;
3455+
S.AddSYCLWorkGroupSizeHintAttr(D, AL, XDimExpr, YDimExpr, ZDimExpr);
3456+
}
34543457

3455-
Expr *XDimExpr = AL.getArgAsExpr(0);
3456-
Expr *YDimExpr =
3457-
AL.isArgExpr(1) ? AL.getArgAsExpr(1) : SetDefaultValue(S, AL);
3458-
Expr *ZDimExpr =
3459-
AL.isArgExpr(2) ? AL.getArgAsExpr(2) : SetDefaultValue(S, AL);
3460-
S.AddWorkGroupSizeHintAttr(D, AL, XDimExpr, YDimExpr, ZDimExpr);
3458+
static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
3459+
// Handle the attribute based on whether we are targeting SYCL or not.
3460+
if (S.getLangOpts().SYCLIsDevice || S.getLangOpts().SYCLIsHost)
3461+
handleSYCLWorkGroupSizeHint(S, D, AL);
3462+
else
3463+
handleWorkGroupSize<WorkGroupSizeHintAttr>(S, D, AL);
34613464
}
34623465

34633466
// Checks correctness of mutual usage of different work_group_size attributes:
@@ -7912,7 +7915,7 @@ void Sema::CheckSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) {
79127915
for (const auto *Attr : std::vector<AttributeCommonInfo *>{
79137916
D->getAttr<SYCLReqdWorkGroupSizeAttr>(),
79147917
D->getAttr<IntelReqdSubGroupSizeAttr>(),
7915-
D->getAttr<WorkGroupSizeHintAttr>(),
7918+
D->getAttr<SYCLWorkGroupSizeHintAttr>(),
79167919
D->getAttr<SYCLDeviceHasAttr>()})
79177920
if (Attr)
79187921
Diag(Attr->getLoc(), diag::warn_sycl_old_and_new_kernel_attributes)
@@ -12057,6 +12060,9 @@ void Sema::ProcessDeclAttributeList(
1205712060
} else if (const auto *A = D->getAttr<SYCLReqdWorkGroupSizeAttr>()) {
1205812061
Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
1205912062
D->setInvalidDecl();
12063+
} else if (const auto *A = D->getAttr<SYCLWorkGroupSizeHintAttr>()) {
12064+
Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
12065+
D->setInvalidDecl();
1206012066
} else if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
1206112067
Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
1206212068
D->setInvalidDecl();

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -531,7 +531,7 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD,
531531
llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
532532
// FIXME: Make this list self-adapt as new SYCL attributes are added.
533533
return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
534-
SYCLReqdWorkGroupSizeAttr, WorkGroupSizeHintAttr,
534+
SYCLReqdWorkGroupSizeAttr, SYCLWorkGroupSizeHintAttr,
535535
SYCLIntelKernelArgsRestrictAttr, SYCLIntelNumSimdWorkItemsAttr,
536536
SYCLIntelSchedulerTargetFmaxMhzAttr,
537537
SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
@@ -4383,9 +4383,9 @@ static void PropagateAndDiagnoseDeviceAttr(
43834383
}
43844384
break;
43854385
}
4386-
case attr::Kind::WorkGroupSizeHint: {
4387-
auto *WGSH = cast<WorkGroupSizeHintAttr>(A);
4388-
if (auto *Existing = SYCLKernel->getAttr<WorkGroupSizeHintAttr>()) {
4386+
case attr::Kind::SYCLWorkGroupSizeHint: {
4387+
auto *WGSH = cast<SYCLWorkGroupSizeHintAttr>(A);
4388+
if (auto *Existing = SYCLKernel->getAttr<SYCLWorkGroupSizeHintAttr>()) {
43894389
if (S.AnyWorkGroupSizesDiffer(Existing->getXDim(), Existing->getYDim(),
43904390
Existing->getZDim(), WGSH->getXDim(),
43914391
WGSH->getYDim(), WGSH->getZDim())) {

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -842,9 +842,9 @@ static void instantiateSYCLAddIRAnnotationsMemberAttr(
842842
S.AddSYCLAddIRAnnotationsMemberAttr(New, *A, Args);
843843
}
844844

845-
static void instantiateWorkGroupSizeHintAttr(
845+
static void instantiateSYCLWorkGroupSizeHintAttr(
846846
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
847-
const WorkGroupSizeHintAttr *A, Decl *New) {
847+
const SYCLWorkGroupSizeHintAttr *A, Decl *New) {
848848
EnterExpressionEvaluationContext Unevaluated(
849849
S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
850850
ExprResult XResult = S.SubstExpr(A->getXDim(), TemplateArgs);
@@ -857,8 +857,8 @@ static void instantiateWorkGroupSizeHintAttr(
857857
if (ZResult.isInvalid())
858858
return;
859859

860-
S.AddWorkGroupSizeHintAttr(New, *A, XResult.get(), YResult.get(),
861-
ZResult.get());
860+
S.AddSYCLWorkGroupSizeHintAttr(New, *A, XResult.get(), YResult.get(),
861+
ZResult.get());
862862
}
863863

864864
static void instantiateSYCLIntelMaxWorkGroupSizeAttr(
@@ -1200,8 +1200,8 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
12001200
*this, TemplateArgs, SYCLAddIRAnnotationsMember, New);
12011201
continue;
12021202
}
1203-
if (const auto *A = dyn_cast<WorkGroupSizeHintAttr>(TmplAttr)) {
1204-
instantiateWorkGroupSizeHintAttr(*this, TemplateArgs, A, New);
1203+
if (const auto *A = dyn_cast<SYCLWorkGroupSizeHintAttr>(TmplAttr)) {
1204+
instantiateSYCLWorkGroupSizeHintAttr(*this, TemplateArgs, A, New);
12051205
continue;
12061206
}
12071207
if (const auto *A = dyn_cast<SYCLDeviceHasAttr>(TmplAttr)) {

‎clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -367,4 +367,4 @@ int main() {
367367
// CHECK: ![[NUM32]] = !{i32 16, i32 16, i32 32}
368368
// CHECK: ![[NUM88]] = !{i32 8, i32 8, i32 8}
369369
// CHECK: ![[NUM22]] = !{i32 2, i32 2, i32 2}
370-
// CHECK: ![[NUM123]] = !{i32 1, i32 2, i32 3}
370+
// CHECK: ![[NUM123]] = !{i32 3, i32 2, i32 1}

‎clang/test/CodeGenSYCL/check-work-group-attributes-match.cpp

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2020 -opaque-pointers -emit-llvm -o - %s | FileCheck %s
2+
// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2020 -no-opaque-pointers -emit-llvm -o - %s | FileCheck %s
3+
4+
// Tests that work_group_size_hint and reqd_work_group_size generate the same
5+
// metadata nodes for the same arguments.
6+
7+
#include "sycl.hpp"
8+
9+
using namespace sycl;
10+
11+
int main() {
12+
queue q;
13+
14+
q.submit([&](handler &h) {
15+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_1d() #0 {{.*}} !work_group_size_hint ![[WG1D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WG1D]]
16+
h.single_task<class kernel_1d>([]() [[sycl::work_group_size_hint(8)]] [[sycl::reqd_work_group_size(8)]] {});
17+
});
18+
19+
q.submit([&](handler &h) {
20+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_2d() #0 {{.*}} !work_group_size_hint ![[WG2D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WG2D]]
21+
h.single_task<class kernel_2d>([]() [[sycl::work_group_size_hint(8, 16)]] [[sycl::reqd_work_group_size(8, 16)]] {});
22+
});
23+
24+
q.submit([&](handler &h) {
25+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_3d() #0 {{.*}} !work_group_size_hint ![[WG3D:[0-9]+]]{{.*}} !reqd_work_group_size ![[WG3D]]
26+
h.single_task<class kernel_3d>([]() [[sycl::work_group_size_hint(8, 16, 32)]] [[sycl::reqd_work_group_size(8, 16, 32)]] {});
27+
});
28+
}
29+
30+
// CHECK: ![[WG1D]] = !{i32 8}
31+
// CHECK: ![[WG2D]] = !{i32 16, i32 8}
32+
// CHECK: ![[WG3D]] = !{i32 32, i32 16, i32 8}

‎clang/test/CodeGenSYCL/no_opaque_check-direct-attribute-propagation.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -367,4 +367,4 @@ int main() {
367367
// CHECK: ![[NUM32]] = !{i32 16, i32 16, i32 32}
368368
// CHECK: ![[NUM88]] = !{i32 8, i32 8, i32 8}
369369
// CHECK: ![[NUM22]] = !{i32 2, i32 2, i32 2}
370-
// CHECK: ![[NUM123]] = !{i32 1, i32 2, i32 3}
370+
// CHECK: ![[NUM123]] = !{i32 3, i32 2, i32 1}

0 commit comments

Comments
 (0)