Skip to content

[SYCL] Remove arbitrary range Part 1 #3133

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

Closed
wants to merge 1 commit into from
Closed

[SYCL] Remove arbitrary range Part 1 #3133

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
16 changes: 0 additions & 16 deletions clang/include/clang/Basic/Attr.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1269,14 +1269,6 @@ def SYCLIntelSchedulerTargetFmaxMhz : InheritableAttr {
let LangOpts = [SYCLIsDevice, SYCLIsHost];
let Subjects = SubjectList<[Function], ErrorDiag>;
let Documentation = [SYCLIntelSchedulerTargetFmaxMhzAttrDocs];
let AdditionalMembers = [{
static unsigned getMinValue() {
return 0;
}
static unsigned getMaxValue() {
return 1024*1024;
}
}];
}

def SYCLIntelMaxWorkGroupSize : InheritableAttr {
Expand Down Expand Up @@ -1331,14 +1323,6 @@ def SYCLIntelLoopFuse : InheritableAttr {
let Accessors = [Accessor<"isIndependent",
[CXX11<"intel", "loop_fuse_independent">]>];
let Documentation = [SYCLIntelLoopFuseDocs];
let AdditionalMembers = [{
static unsigned getMinValue() {
return 0;
}
static unsigned getMaxValue() {
return 1024*1024;
}
}];
}

def C11NoReturn : InheritableAttr {
Expand Down
7 changes: 7 additions & 0 deletions clang/include/clang/Sema/Sema.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12970,6 +12970,13 @@ void Sema::addIntelSYCLSingleArgFunctionAttr(Decl *D,
return;
}
}
if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz) {
if (ArgInt < 0) {
Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
<< CI.getAttrName() << /*non-negative*/ 1;
return;
}
}
Comment on lines +12973 to +12979
Copy link
Contributor

@smanna12 smanna12 Feb 2, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

You can combine this inside addIntelSYCLSingleArgFunctionAttr () function :

if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim ||
CI.getParsedKind() == ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz ) {
if (ArgInt < 0) {
Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
<< CI.getAttrName() << /non-negative/ 1;
return;
}

Copy link
Contributor Author

Choose a reason for hiding this comment

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

AT_SYCLIntelMaxGlobalWorkDim also has an upper-bound check, which is why I kept AT_SYCLIntelSchedulerTargetFmaxMhz in a separate check. There are actually also more FPGA attributes I am adding irefactoring in another patch, which will change this function even more. I will refactor this function then if that sounds ok to you.

Copy link
Contributor

@smanna12 smanna12 Feb 2, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

We definitely need to keep upper bound check for max_global_work_dim kernel attribute (which has a natural limitation of '3'). Several diagnostics and work_gorup_size attributes are dependent on this.

I agree with @MrSidims comments here #2911 (comment)

We should only remove the arbitrary upper bound here:
#2911 (comment)

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Yes we need to keep the upper bound of 3 for max_global_work_dim kernel attribute. I am removing only arbitrary upper bounds of 1024*1024 for the attributes mentioned here - #2911.

Copy link
Contributor

@smanna12 smanna12 Feb 6, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

addIntelSYCLSingleArgFunctionAttr () - I added this function to handle template parameter support for SYCL FPGA function attributes. So we can add other new function attributes in future here.

The following attributes below are FPGA memory attributes, i think it would be better if we add something like addIntelSingleArgMemoryAttr () instead of adding them inside addIntelSYCLSingleArgFunctionAttr ().

IntelFPGAPrivateCopies
IntelFPGABankBits

1 - 1024*1024
IntelFPGABankWidth
IntelFPGANumBanks
IntelFPGAMaxReplicates

@AaronBallman any thought?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Thanks for taking a look @smanna12 . I do plan to refactor the functions when I work on the IntelFPGA attributes. You can see #3134. I've uploaded a draft PR (with just one attribute. Others will follow) to get initial feedback from @AaronBallman.

}

D->addAttr(::new (Context) AttrType(Context, CI, E));
Expand Down
12 changes: 8 additions & 4 deletions clang/lib/Sema/SemaDeclAttr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3210,7 +3210,8 @@ static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
S.Diag(AL.getLoc(), diag::note_spelling_suggestion)
<< "'intel::scheduler_target_fmax_mhz'";

S.AddOneConstantValueAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>(D, AL, E);
S.addIntelSYCLSingleArgFunctionAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>(
D, AL, E);
}

// Handles max_global_work_dim.
Expand Down Expand Up @@ -3287,10 +3288,13 @@ static bool checkSYCLIntelLoopFuseArgument(Sema &S,
return true;
}

SYCLIntelLoopFuseAttr TmpAttr(S.Context, CI, E);
ExprResult ICE;
if (!ArgVal->isNonNegative()) {
S.Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
<< CI << /*non-negative*/ 1;
return true;
}

return S.checkRangedIntegralArgument<SYCLIntelLoopFuseAttr>(E, &TmpAttr, ICE);
return false;
}

void Sema::addSYCLIntelLoopFuseAttr(Decl *D, const AttributeCommonInfo &CI,
Expand Down
7 changes: 4 additions & 3 deletions clang/test/SemaSYCL/loop_fusion.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,8 +4,8 @@

[[intel::loop_fuse("foo")]] void func() {} // expected-error{{'loop_fuse' attribute requires an integer constant}}

[[intel::loop_fuse(1048577)]] void func1() {} // expected-error{{'loop_fuse' attribute requires integer constant between 0 and 1048576 inclusive}}
[[intel::loop_fuse_independent(-1)]] void func2() {} // expected-error{{'loop_fuse_independent' attribute requires integer constant between 0 and 1048576 inclusive}}
[[intel::loop_fuse(1048577)]] void func1() {} // OK
[[intel::loop_fuse_independent(-1)]] void func2() {} // expected-error{{'loop_fuse_independent' attribute requires a non-negative integral compile time constant expression}}

[[intel::loop_fuse(0, 1)]] void func3() {} // expected-error{{'loop_fuse' attribute takes no more than 1 argument}}
[[intel::loop_fuse_independent(2, 3)]] void func4() {} // expected-error{{'loop_fuse_independent' attribute takes no more than 1 argument}}
Expand Down Expand Up @@ -48,13 +48,14 @@
[[intel::loop_fuse]] void func16();

template <int N>
[[intel::loop_fuse(N)]] void func17(); // expected-error{{'loop_fuse' attribute requires integer constant between 0 and 1048576 inclusive}}
[[intel::loop_fuse(N)]] void func17(); // expected-error{{'loop_fuse' attribute requires a non-negative integral compile time constant expression}}

template <typename Ty>
[[intel::loop_fuse(Ty{})]] void func18() {} // expected-error{{'loop_fuse' attribute requires an integer constant}}

void checkTemplates() {
func17<-1>(); // expected-note{{in instantiation of}}
func17<0>(); // OK
func18<float>(); // expected-note{{in instantiation of}}
}

Expand Down
12 changes: 4 additions & 8 deletions clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,20 +10,16 @@ template <int N>
[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {}

int main() {
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 'void ()'
// CHECK: FunctionDecl {{.*}}test_kernel1 'void ()'
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
Copy link
Contributor Author

Choose a reason for hiding this comment

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

I think ConstantExpr is no longer generated here because of a difference in how attribute expression is obtained in addIntelSYCLSingleArgFunctionAttr vs AddOneConstantValueAttr. The latter calls VerifyIntegerConstantExpression which generates this AST.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

@AaronBallman @premanandrao Should the ConstantExpr be generated. It isn't consistently generated for IntelSYCL attributes, which is why I removed it here.

Copy link
Contributor

Choose a reason for hiding this comment

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

I think it should be -- see my reasoning in #3134 (comment)

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Ok. Thanks for the review! I'll make the change!

// CHECK-NEXT: value: Int 5
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5
// expected-warning@+3 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}}
// expected-note@+2 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}}
cl::sycl::kernel_single_task<class test_kernel1>(
[]() [[intelfpga::scheduler_target_fmax_mhz(5)]]{});

// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 'void ()'
// CHECK: FunctionDecl {{.*}}test_kernel2 'void ()'
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
// CHECK-NEXT: value: Int 2
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2
cl::sycl::kernel_single_task<class test_kernel2>(
[]() { func(); });
Expand All @@ -39,10 +35,10 @@ int main() {
[[intel::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}}

cl::sycl::kernel_single_task<class test_kernel4>(
[]() [[intel::scheduler_target_fmax_mhz(1048577)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}}
[]() [[intel::scheduler_target_fmax_mhz(1048577)]]{}); // OK

cl::sycl::kernel_single_task<class test_kernel5>(
[]() [[intel::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}}
[]() [[intel::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}

cl::sycl::kernel_single_task<class test_kernel6>(
[]() [[intel::scheduler_target_fmax_mhz(1), intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different parameters}}
Expand Down