[SYCL] Introduce min_work_groups_per_cu and max_work_groups_per_mp #11192

jchlanda · 2023-09-15T12:27:52Z

The attributes match to CUDA's launch bounds minBlocksPerMultiprocessor
and maxBlocksPerCluster respectively. See: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#cluster-dimension-directives-maxclusterrank and
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#launch-bounds for details.

jchlanda · 2023-09-27T07:05:02Z

This is now ready for review, as llvm/llvm-project#66496 went into upstream llvm. Not sure what is the practice with changes depending on the upstream (e0a9fcfeec8606db7f70a5ea19ba44b2b0d18fd9) we could cherry-pick or make this PR wait till the pull-down.

premanandrao · 2023-09-28T16:11:30Z

This is now ready for review, as llvm/llvm-project#66496 went into upstream llvm. Not sure what is the practice with changes depending on the upstream (e0a9fcf) we could cherry-pick or make this PR wait till the pull-down.

Unless there is urgency, I would prefer the PR wait until after the pulldown.

jchlanda · 2023-10-03T06:36:45Z

This is now ready for review, as llvm/llvm-project#66496 went into upstream llvm. Not sure what is the practice with changes depending on the upstream (e0a9fcf) we could cherry-pick or make this PR wait till the pull-down.

Unless there is urgency, I would prefer the PR wait until after the pulldown.

Sure, that's fine with me.

The attributes match to CUDA's launch bounds minBlocksPerMultiprocessor and maxBlocksPerCluster respectively.

jchlanda · 2023-10-30T15:40:06Z

@elizabethandrews @smanna12 is there anything else that needs to be address for this PR?

clang/include/clang/Basic/DiagnosticSemaKinds.td

elizabethandrews · 2023-10-30T22:30:28Z

clang/lib/Sema/SemaDeclAttr.cpp

@@ -22,6 +22,7 @@
 #include "clang/AST/Mangle.h"
 #include "clang/AST/RecursiveASTVisitor.h"
 #include "clang/AST/Type.h"
+#include "clang/Basic/AttributeCommonInfo.h"


Do you actually need to add this? Unless I'm missing something obvious, I thought we've been using AttributeCommonInfo in this file prior to this PR ? I just assumed it was included by one of the other headers or something. What necessitated the addition in this PR?

Apologies, that seems to be IDE residue.

elizabethandrews · 2023-10-30T22:33:57Z

clang/lib/Sema/SemaDeclAttr.cpp

  return AL.getLocation();
 }
+template <typename T,
+          std::enable_if_t<std::is_same_v<AttributeCommonInfo, T>, bool> = true>
+static SourceLocation getAttrLoc(const T &AL) {


Where is this being called?

In checkUInt32Argument. And it is being called from https://github.com/jchlanda/llvm/blob/jakub/launch_bounds/clang/lib/Sema/SemaDeclAttr.cpp#L4466 with the const AttributeCommonInfo &CI as AttrInfotemplate parameter.

elizabethandrews · 2023-10-30T22:47:22Z

clang/lib/Sema/SemaDeclAttr.cpp

+                             true /* StrictlyUnsigned */))
+      return;
+
+    // Validate that we have an integer constant expression and then store the


@smanna12 please confirm but I do not think you need to do both checkUInt32Argument and VerifyIntegerConstantExpression. I could be wrong but I feel we're repeating code by calling both these functions.

Perhaps as a side note, VerifyIntegerConstantExpression has this nice property (I'm slightly out of my depths with clang, so forgive if I misname the things) that it rewrites the AST, so for each use of the attributes, it always follows the same pattern:

ConstantExpr{{.*}}'int' value: Int N IntegerLiteral {{.*}} 'int' N

Can you just get the Value from VerifyIntegerConstantExpression and replicate this diagnostic from checkUInt32Argument -

if (!I->isIntN(32)) { S.Diag(Expr->getExprLoc(), diag::err_ice_too_large) << toString(*I, 10, false) << 32 << /* Unsigned */ 1; return false; } if (StrictlyUnsigned && I->isSigned() && I->isNegative()) { S.Diag(getAttrLoc(AI), diag::err_attribute_requires_positive_integer) << &AI << /*non-negative*/ 1; return false; }

I'm guessing there are other attributes which do similar things. @smanna will know better since she works with these a lot.

Hm, that's exactly how it used to be, I had it in a helper function used by both attributes: 9431687#diff-2a5bdb2d9f07f8d77de51d5403d349c22978141b6de6bd87fc5e449f5ed95becL4450-L4469

Yea sorry for some confusion over this. I thought you would be able to just call checkUInt32Argument (without calling VerifyIntegerConstantExpression, but I am not so sure about this anymore. I've asked a colleague about this and will get back to you once he replies). In any case though, the issue with the code you first had and the current one now is with calling getIntegerConstantExpr or it's equivalent twice. VerifyIntegerConstantExpression should give you the value you need right?

// This attribute requires a strictly positive value. if (ArgVal <= 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << CI << /*positive*/ 0; return std::nullopt; } }

I'm guessing there are other attributes which do similar things. @smanna will know better since she works with these a lot.

SYCLIntelBankBitsAttr is doing similar thing in void Sema::AddSYCLIntelBankBitsAttr()

@jchlanda I spoke to Aaron who has implemented/worked with a lot of attribute code in community clang to get a better understanding of when to call checkUInt32Argument vs VerifyIntegerConstantExpression. Apparently checkUInt32Argument was implemented when template instantiations of arguments were not really a thing. Since we're now doing more with template instantiations, it makes sense to catch the constant expression (checkUInt32Argument discards it after evaluating it) and keep it in AST. So VerifyIntegerConstantExpression + required diagnostics as suggested above is the right way to handle this. I apologize for the confusion here. This area needs some cleanup in community clang.

Sure, not a problem at all. I got rid of the extra overload to make checkUInt32Argument work, and encapsulated those check in a static helper function that is called instead of checkUInt32Argument (while keeping VerifyIntegerConstantExpression), it required some test changes, all commited here: acba95f

// This attribute requires a strictly positive value. if (ArgVal <= 0) { Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer) << CI << /*positive*/ 0; return std::nullopt; } }

I'm guessing there are other attributes which do similar things. @smanna will know better since she works with these a lot.

SYCLIntelBankBitsAttr is doing similar thing in void Sema::AddSYCLIntelBankBitsAttr()

Please stop using @smanna.

clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp

clang/test/CodeGenSYCL/lb_sm_90.cpp

clang/test/SemaSYCL/lb_sm_70.cpp

clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp

clang/test/SemaSYCL/lb_sm_70.cpp

elizabethandrews · 2023-10-31T17:32:24Z

clang/lib/Sema/SemaDeclAttr.cpp

+                             true /* StrictlyUnsigned */))
+      return;
+
+    // Validate that we have an integer constant expression and then store the


Can you just get the Value from VerifyIntegerConstantExpression and replicate this diagnostic from checkUInt32Argument -

if (!I->isIntN(32)) { S.Diag(Expr->getExprLoc(), diag::err_ice_too_large) << toString(*I, 10, false) << 32 << /* Unsigned */ 1; return false; } if (StrictlyUnsigned && I->isSigned() && I->isNegative()) { S.Diag(getAttrLoc(AI), diag::err_attribute_requires_positive_integer) << &AI << /*non-negative*/ 1; return false; }

I'm guessing there are other attributes which do similar things. @smanna will know better since she works with these a lot.

clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp

elizabethandrews · 2023-10-31T21:26:57Z

clang/test/SemaSYCL/lb_sm_70.cpp

@@ -0,0 +1,60 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_70 -fsycl-is-device -Wno-c++23-extensions -S -emit-llvm %s -o -fsyntax-only -verify %s


You can remove the -S -emit-llvm here

clang/lib/CodeGen/CodeGenFunction.cpp

clang/lib/CodeGen/Targets/NVPTX.cpp

smanna · 2023-11-03T12:31:38Z

Please do not tag me ***@***.***). You should be tagging @smanna12. I‘m getting spammed.On 1 Nov 2023, at 15:18, smanna12 ***@***.***> wrote: @smanna12 commented on this pull request. In clang/lib/CodeGen/Targets/NVPTX.cpp:

@@ -245,6 +245,32 @@ void NVPTXTargetCodeGenInfo::setTargetAttributes(

// And kernel functions are not subject to inlining F->addFnAttr(llvm::Attribute::NoInline); } + if (const auto *MWGS = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) { + auto MaxThreads = (*MWGS->getZDimVal()).getExtValue() * + (*MWGS->getYDimVal()).getExtValue() * + (*MWGS->getXDimVal()).getExtValue(); + if (MaxThreads > 0) + addNVVMMetadata(F, "maxntidx", MaxThreads); + + auto attrValue = [&](Expr *E) { + const auto *CE = cast<ConstantExpr>(E); + std::optional<llvm::APInt> Val = CE->getResultAsAPSInt(); + assert(Val.has_value() && "Failed to get attribute value."); assert does not need here. cast<> will give an assert for null attribute value. —Reply to this email directly, view it on GitHub, or unsubscribe.You are receiving this because you were mentioned.Message ID: ***@***.***>

elizabethandrews · 2023-11-03T16:16:44Z

clang/lib/Sema/SemaDeclAttr.cpp

-template <typename AttrInfo>
-static std::enable_if_t<std::is_base_of_v<Attr, AttrInfo>, SourceLocation>
-getAttrLoc(const AttrInfo &AL) {
+/// Helper functions to provide Attribute Location for the Attr types,


These changes can be removed now right?

My bad, I removed the overload but did not remember about that change, fixed now.

elizabethandrews · 2023-11-03T16:23:58Z

clang/test/SemaSYCL/lb_sm_70.cpp

+
+    constexpr float A = 2.0;
+    // expected-warning@+5{{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute}}
+    // expected-error@+3 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}


Ugh. Not a huge fan of the new diagnostics but I guess we can't do anything since I assume these are thrown from VerifyIntegerConstantExpression?

Yes, that's generated in VerifyIntegerConstantExpression

elizabethandrews

LGTM! Thank you for making all the changes!

jchlanda · 2023-11-07T11:58:00Z

LGTM! Thank you for making all the changes!

Not a problem, than you for a thorough review :)

jchlanda requested review from a team as code owners September 15, 2023 12:27

jchlanda marked this pull request as draft September 15, 2023 12:28

jchlanda force-pushed the jakub/launch_bounds branch 2 times, most recently from 3f48159 to ecf8358 Compare September 15, 2023 12:35

This was referenced Sep 15, 2023

[SYCL][CUDA][HIP] Guidance for __launch_bounds__ in SYCL #8080

Closed

[SYCL][CUDA] support launch bounds #9772

Closed

jchlanda force-pushed the jakub/launch_bounds branch 2 times, most recently from 2742daa to ff95dde Compare September 15, 2023 13:23

jchlanda force-pushed the jakub/launch_bounds branch from ff95dde to e0a9fcf Compare September 27, 2023 07:03

jchlanda marked this pull request as ready for review September 27, 2023 07:03

jchlanda requested a review from a team as a code owner September 27, 2023 07:03

jchlanda requested a review from bso-intel September 27, 2023 07:03

jchlanda had a problem deploying to WindowsCILock September 27, 2023 07:04 — with GitHub Actions Failure

jchlanda temporarily deployed to WindowsCILock September 27, 2023 08:20 — with GitHub Actions Inactive

jchlanda had a problem deploying to WindowsCILock October 2, 2023 16:52 — with GitHub Actions Failure

jchlanda temporarily deployed to WindowsCILock October 2, 2023 18:11 — with GitHub Actions Inactive

jchlanda force-pushed the jakub/launch_bounds branch from e0a9fcf to fd13ad4 Compare October 5, 2023 12:23

jchlanda had a problem deploying to WindowsCILock October 5, 2023 12:24 — with GitHub Actions Failure

jchlanda force-pushed the jakub/launch_bounds branch from fd13ad4 to ed14966 Compare October 10, 2023 11:21

jchlanda had a problem deploying to WindowsCILock October 10, 2023 12:08 — with GitHub Actions Failure

jchlanda force-pushed the jakub/launch_bounds branch from ed14966 to 1264d97 Compare October 10, 2023 12:37

jchlanda had a problem deploying to WindowsCILock October 10, 2023 12:38 — with GitHub Actions Failure

jchlanda temporarily deployed to WindowsCILock October 10, 2023 13:30 — with GitHub Actions Inactive

[SYCL] Introduce min_work_groups_per_cu and max_work_groups_per_mp

bbe994d

The attributes match to CUDA's launch bounds minBlocksPerMultiprocessor and maxBlocksPerCluster respectively.

jchlanda force-pushed the jakub/launch_bounds branch from 1264d97 to bbe994d Compare October 11, 2023 06:01

jchlanda temporarily deployed to WindowsCILock October 11, 2023 06:02 — with GitHub Actions Inactive

jchlanda temporarily deployed to WindowsCILock October 25, 2023 10:06 — with GitHub Actions Inactive

jchlanda temporarily deployed to WindowsCILock October 25, 2023 11:03 — with GitHub Actions Inactive

jchlanda added 2 commits October 26, 2023 23:43

Remove SupportsNonconformingLambdaSyntax and adjust tests

7744f69

Merge remote-tracking branch 'upstream/sycl' into jakub/launch_bounds

80cc279

jchlanda temporarily deployed to WindowsCILock October 27, 2023 07:03 — with GitHub Actions Inactive

jchlanda temporarily deployed to WindowsCILock October 27, 2023 08:03 — with GitHub Actions Inactive

elizabethandrews reviewed Oct 31, 2023

View reviewed changes

PR feedback 2.

0d6b772

jchlanda requested a review from elizabethandrews October 31, 2023 11:51

jchlanda temporarily deployed to WindowsCILock October 31, 2023 11:52 — with GitHub Actions Inactive

jchlanda temporarily deployed to WindowsCILock October 31, 2023 12:24 — with GitHub Actions Inactive

elizabethandrews reviewed Oct 31, 2023

View reviewed changes

smanna12 reviewed Nov 1, 2023

View reviewed changes

clang/lib/CodeGen/CodeGenFunction.cpp Outdated Show resolved Hide resolved

smanna12 reviewed Nov 1, 2023

View reviewed changes

clang/lib/CodeGen/Targets/NVPTX.cpp Outdated Show resolved Hide resolved

PR: attribute check

acba95f

jchlanda temporarily deployed to WindowsCILock November 3, 2023 10:52 — with GitHub Actions Inactive

PR feedback 3

7691bba

jchlanda requested review from smanna12 and elizabethandrews November 3, 2023 11:07

jchlanda temporarily deployed to WindowsCILock November 3, 2023 11:27 — with GitHub Actions Inactive

jchlanda temporarily deployed to WindowsCILock November 3, 2023 12:29 — with GitHub Actions Inactive

elizabethandrews reviewed Nov 3, 2023

View reviewed changes

Remove changes to getAttrLoc

1342224

jchlanda temporarily deployed to WindowsCILock November 6, 2023 09:11 — with GitHub Actions Inactive

jchlanda temporarily deployed to WindowsCILock November 6, 2023 09:39 — with GitHub Actions Inactive

elizabethandrews approved these changes Nov 6, 2023

View reviewed changes

againull merged commit 52556d8 into intel:sycl Nov 6, 2023

		@@ -0,0 +1,60 @@
		// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_70 -fsycl-is-device -Wno-c++23-extensions -S -emit-llvm %s -o -fsyntax-only -verify %s

[SYCL] Introduce min_work_groups_per_cu and max_work_groups_per_mp #11192

[SYCL] Introduce min_work_groups_per_cu and max_work_groups_per_mp #11192

Uh oh!

Conversation

jchlanda commented Sep 15, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

jchlanda commented Sep 27, 2023

Uh oh!

premanandrao commented Sep 28, 2023

Uh oh!

jchlanda commented Oct 3, 2023

Uh oh!

jchlanda commented Oct 30, 2023

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jchlanda Nov 1, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

elizabethandrews Nov 1, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

smanna12 Nov 1, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

smanna commented Nov 3, 2023 via email

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

elizabethandrews left a comment

Choose a reason for hiding this comment

Uh oh!

jchlanda commented Nov 7, 2023

Uh oh!

Uh oh!

jchlanda commented Sep 15, 2023 •

edited

Loading

jchlanda Nov 1, 2023 •

edited

Loading

elizabethandrews Nov 1, 2023 •

edited

Loading

smanna12 Nov 1, 2023 •

edited

Loading