[SYCL] Support intel::reqd_work_group_size #1328

fadeeval · 2020-03-16T10:03:38Z

Implementing intel::reqd_work_group_size.

Signed-off-by: Aleksander Fadeev aleksander.fadeev@intel.com

erichkeane · 2020-03-16T13:26:35Z

clang/lib/Sema/SemaDeclAttr.cpp

@@ -2923,14 +2923,20 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
    return;

  uint32_t WGSize[3];
+  if (AL.getNormalizedFullName() == "intel::reqd_work_group_size") {


getAttributeSpellingListIndex exists for this purpose.

But in this situation this statement sifts not just ReqdWorkGroupSize spellings, but also different WorkGroupAttr.

I think, that getNormalizedFullName() more applicable here.

We have 2 attributes with the same spelling!?! That is completely unacceptable.

No, but getAttributeSpellingListIndex() returns unsigned and an attribute spelling index is a just unsigned in fact. And in result we get the same outcome (true) with diff spellings indices of diff attributes.

Ah... I think I see the problem in TableGen. I'm going to push a fix to LLOrg in the next hour or so to make it so that this enum still gets emitted, since it is useful here.

Suggested a review, but my change is here: https://reviews.llvm.org/D76289

There, got it in! rG661c950630fb. Once that comes down, this line should be AL.getKind == ?? && AL.getSpelling == ReqdWorkGroupSize::CXX11_intel_reqd_work_group_size.

Should I wait until the merge?

erichkeane · 2020-03-16T13:36:02Z

clang/test/SemaSYCL/intel-reqd-work-group-size.cpp

+
+class Functor16 {
+public:
+  [[intel::reqd_work_group_size(16)]] void operator()() {}


Any examples with 2? What happens when I pass a negative to these? How about when things conflict with '1'?

@fadeeval please apply this.

erichkeane · 2020-03-16T13:36:46Z

clang/test/SemaSYCL/intel-reqd-work-group-size.cpp

+// expected-note@-1 {{conflicting attribute is here}}
+[[intel::reqd_work_group_size(32, 1, 1)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}}
+
+[[intel::reqd_work_group_size(16, 1, 1)]] void f16x1x1() {}   // expected-note {{conflicting attribute is here}}


Since the defaults are 1, why set them so here (and around here)?

Since the advantage to this form of attribute is the 'default' values, perhaps we should have more tests evaluating that.

clang/include/clang/Basic/Attr.td

AlexeySachkov · 2020-03-16T15:23:18Z

clang/include/clang/Basic/AttrDocs.td

@@ -1983,6 +1983,17 @@ device kernel, the attribute is ignored and it is not propagated to a kernel.
  }];
 }

+def ReqdWorkGroupSizeAttrDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "cl::reqd_work_group_size, intel::reqd_work_group_size";


What about GNU spelling?

I don't sure, I will ask about it.

The answer of John Pennycook is following:
"I hadn't thought about this before, but I'm going to say "No" – at least for now. If we're not using the C++11 attribute specifier syntax we can't have a namespace, and I don't want to set a precedent for converting between intel:: and intel_ that might accidentally introduce name conflicts."

So that a GNU spelling doesn't need here.

@Pennycook, please note that this is documentation for the compiler itself, not for the particular language. It is used to generated the following documentation: https://clang.llvm.org/docs/AttributeReference.html

So, this attribute is accepted by the compiler in GNU spelling and therefore, must be documented here as well if you adding such documentation. In SYCL-specific documents we can omit GNU spelling, but not here.

BTW, looking at the web-page I referenced above, it would be okay for me to either mention all declared spellings or just leave reqd_work_group_size generic string

@fadeeval, this comment is the only one I have. I would like header to be aligned with the actual documentation, which describes all three spellings

clang/include/clang/Basic/AttrDocs.td

clang/lib/Sema/SemaDeclAttr.cpp

erichkeane

As a side note, please give meaningful names to your commits/updates. I have no idea what 10.2 means.

erichkeane · 2020-03-16T15:31:03Z

clang/lib/Sema/SemaDeclAttr.cpp

@@ -2923,14 +2923,20 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
    return;

  uint32_t WGSize[3];
+  if (AL.getAttributeSpellingListIndex() == 2) {


There is an enum class that gets created in the attribute type itself that has these lists. Something like ReqdWorkGroupSizeAttr::CXX11_intel_reqd_work_group_size

I cannot see the exact spelling because I don't have your patch installed, so i sort of guessed based on the spelling. Look in your build directory at:
whatever/build/tools/clang/include/clang/Sema/Attrs.inc
You'll see something like this:

class SectionAttr : public InheritableAttr { unsigned nameLength; char *name; public: enum Spelling { GNU_section = 0, CXX11_gnu_section = 1, Declspec_allocate = 2, SpellingNotCalculated = 15 }; ...

See the Spelling enum which should have the entry for your new spelling.

I watched in this file and I didn't find enum Spelling in class ReqdWorkGroupSize. It is very strange.

erichkeane · 2020-03-17T14:44:52Z

clang/lib/Sema/SemaDeclAttr.cpp

@@ -2923,14 +2923,20 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
    return;

  uint32_t WGSize[3];


Note: Every comparison with this gives me a -Wsign-compare warning. XDim, YDim, and ZDim are all 'int' in type. This should be so as well likely.

The problem is TableGen has no unsigned type.

Okey, this problem I solved, it seems.

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

erichkeane

I'm OK with this, but @AlexeySachkov has a number of comments that I think are unaddressed. Once he's OK, I am as well.

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

fadeeval · 2020-03-18T22:22:04Z

clang-format-check doesn't pass, because I don't like the option, it suggests to me.

clang/include/clang/Basic/Attr.td

AlexeySachkov · 2020-03-19T08:15:35Z

clang/include/clang/Basic/AttrDocs.td

@@ -1983,6 +1983,17 @@ device kernel, the attribute is ignored and it is not propagated to a kernel.
  }];
 }

+def ReqdWorkGroupSizeAttrDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "cl::reqd_work_group_size, intel::reqd_work_group_size";


@Pennycook, please note that this is documentation for the compiler itself, not for the particular language. It is used to generated the following documentation: https://clang.llvm.org/docs/AttributeReference.html

So, this attribute is accepted by the compiler in GNU spelling and therefore, must be documented here as well if you adding such documentation. In SYCL-specific documents we can omit GNU spelling, but not here.

BTW, looking at the web-page I referenced above, it would be okay for me to either mention all declared spellings or just leave reqd_work_group_size generic string

clang/include/clang/Basic/AttrDocs.td

clang/test/SemaSYCL/intel-reqd-work-group-size.cpp

AlexeySachkov · 2020-03-19T08:39:20Z

clang-format-check doesn't pass, because I don't like the option, it suggests to me.

There is nothing wrong with suggested formatting, I suggest to apply it. Tagging @bader to comment

bader · 2020-03-19T09:01:37Z

clang-format-check doesn't pass, because I don't like the option, it suggests to me.

There is nothing wrong with suggested formatting, I suggest to apply it. Tagging @bader to comment

Absolutely agree with @AlexeySachkov.

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

AlexeySachkov · 2020-03-19T16:44:52Z

clang/include/clang/Basic/AttrDocs.td

+  let Heading = "GNU<reqd_work_group_size>, CXX11<cl::reqd_work_group_size, 
+                CXX11<intel::reqd_work_group_size>";


Suggested change

let Heading = "GNU<reqd_work_group_size>, CXX11<cl::reqd_work_group_size,

CXX11<intel::reqd_work_group_size>";

let Heading = "__attribute__((reqd_work_group_size)), [[cl::reqd_work_group_size]],

[[intel::reqd_work_group_size]]";

buildbot/Build_PR_Win didn't pass because Heading, and writes:
"D:/buildbot/Product_worker_intel/Build_PR_With_Lit_Win/llvm.src/clang/include\clang/Basic/AttrDocs.td:2008:17: error: Unknown token when parsing a value
let Heading = "attribute((reqd_work_group_size)), [[cl::reqd_work_group_size]],"

Maybe just "reqd_work_group_size" will be better?

Just "reqd_work_group_size" works for me

unknown token when parsing a value
let Heading = "attribute((reqd_work_group_size)), [[cl::reqd_work_group_size]],"

Probably this is some kind of keyword which is being parsed regardless of the fact that it is within a string literal

Fznamznon · 2020-03-19T16:52:18Z

clang/include/clang/Basic/AttrDocs.td

+and allows to specify exact *local_work_size* which must be used as
+argument to **clEnqueueNDRangeKernel** (in OpenCL) or to
+**parallel_for** in SYCL. This allows the compiler to optimize the
+generated code appropriately for this kernel.


I'm just a dumb purple bee, so

This allows the compiler to optimize the
generated code appropriately for this kernel.

Which kernel?

The kernel to which attribute is applied. Maybe "the kernel" will be more correct then "this kernel", I think.

Or better to write: "for the kernel to which attribute is applied"?

Both variants looks good to me. I just copied this from OpenCL spec, but apparently it has more context so "this" is applicable there

Fznamznon · 2020-03-19T16:53:59Z

clang/include/clang/Basic/AttrDocs.td

+spelling is a bit different: 
+
+SYCL 1.2.1 describes ``[[cl::reqd_work_group_size(X, Y, Z)]]`` spelling: this
+attribute is legal on device functions and their specification is propagated


Do you mean specification of the attribute is propagated somewhere?

No, it says about function specification, that the attribute setup.

Fznamznon · 2020-03-19T17:01:28Z

clang/include/clang/Basic/AttrDocs.td

+down to any caller of those device functions, such that the kernel attributes
+are the sum of all the kernel attributes of all device functions called.
+See section 6.7 Attributes for more details.


I see that this phrase is actually copy-pasted from the spec, but it looks ugly. Maybe we could ask someone with good enough English to rewrite it...

clang/test/SemaSYCL/intel-reqd-work-group-size.cpp

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Fznamznon · 2020-03-23T13:14:58Z

clang/include/clang/Basic/AttrDocs.td

+As Intel extension, ``[[intel::reqd_work_group_size(X, Y, Z)]]`` spelling is allowed
+which features optional arguments `Y` and `Z`, which simplifies its usage if
+only 1- or 2-dimensional ND-range is assumed by a programmer. These arguments
+defaults to ``1``.


This is from the spec too? which after which feels ugly too. :)

Fznamznon · 2020-03-23T13:15:29Z

clang/include/clang/Basic/AttrDocs.td

+(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see, for example, section
+6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification


Suggested change

(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see, for example, section

6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification

(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see section

6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification for details.

clang/test/SemaSYCL/intel-reqd-work-group-size.cpp

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

…hinx * upstream/sycl: (357 commits) [Support] Implement a simple tabular data management library (intel#1358) [Support] Implement a property set I/O library (intel#1357) [SYCL] Fix buffer constructor using iterators (intel#1386) [SYCL][FPGA] Enable a set of loop attributes (intel#1312) [Driver][SYCL][FPGA] Proper dependency output location when given /Fo<dir> (intel#1346) [SPIR-V] Enabling SPIR-V builtin lookup in device SYCL mode (intel#1384) [SYCL][NFC] Unify setting kernel arguments (intel#1379) [SYCL][Doc] First revision of standard layout relaxation extension (intel#1344) [SYCL] Fixed sub-buffer alloca search (intel#1385) [SYCL][FPGA] Emit multiple IR variants for the IVDep attribute (intel#1383) [SYCL] Add experimental flag to enable front-end optimizations (intel#1376) [SYCL] Remove unexpected double in complex SPIR-V for float support (intel#1381) [SYCL] Default work-group sizes based on max (intel#952) [SYCL][CUDA] Fix usage of multiple backends in the same program (intel#1252) [SPIR-V] Add SPIR-V builtin definitions to the builtin lookup. [SPIR-V] Add macro definition when -fdeclare-spirv-builtins is activated [SYCL] Fix sycl_generic printing [SYCL] Support intel::reqd_work_group_size (intel#1328) [SYCL][NFC] Make the RT::PiPlugin object private (intel#1375) [SPIRV] Add convergent attribute to SPIR-V built-ins (intel#1373) ...

fadeeval requested review from asavonic, Fznamznon and erichkeane March 16, 2020 10:37

erichkeane reviewed Mar 16, 2020

View reviewed changes

AlexeySachkov reviewed Mar 16, 2020

View reviewed changes

erichkeane reviewed Mar 16, 2020

View reviewed changes

erichkeane reviewed Mar 17, 2020

View reviewed changes

fadeeval added 16 commits March 18, 2020 15:23

[SYCL] Support intel::reqd_work_group_size (fix intel#6)

4555f44

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#7)

c704d1a

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#7.2)

d6e3575

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#7.3)

18b0e94

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#7.4)

6a313bf

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#8)

995352e

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#8.2)

421217a

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#8.3)

f9c5822

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#9)

eb5c797

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#9.2)

9c1c130

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#9.3)

a8b1121

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#10)

93a022d

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

[SYCL] Support intel::reqd_work_group_size (fix intel#10.2)

fb7e14e

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

The test commit, didn' have practical benefit

8a3f1b8

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Implementing DefaultUnsignedArgument class

dc0eaac

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Impelmention getAttributeSpellingListIndex()

04a4c83

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

erichkeane reviewed Mar 18, 2020

View reviewed changes

fadeeval added 2 commits March 18, 2020 16:10

Resolving cinflict ClangAttrEmitter.cpp

fc10805

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

The AttrDocs.td changing and other little fixes

971eaa4

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

fadeeval force-pushed the private/fadeeval/Intel-reqd_work_group_size branch from ca62510 to 971eaa4 Compare March 18, 2020 22:07

Formatting

65fc0ba

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

AlexeySachkov reviewed Mar 19, 2020

View reviewed changes

fadeeval added 2 commits March 19, 2020 17:29

AttrDocs.td and formatting

2776fe5

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

adding GNU in AttrDocs.td

545cb9f

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

AlexeySachkov reviewed Mar 19, 2020

View reviewed changes

romanovvlad assigned AlexeySachkov Mar 19, 2020

Fznamznon reviewed Mar 19, 2020

View reviewed changes

fadeeval added 4 commits March 20, 2020 10:51

Docs.td modifying

60541fd

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

fix AttrDocs.td heading

50e397d

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Minifix AttrDocs

cf200c4

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

Formatting

23e0bbb

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

fadeeval requested a review from Fznamznon March 23, 2020 12:18

Fznamznon requested a review from erichkeane March 23, 2020 13:14

Fznamznon reviewed Mar 23, 2020

View reviewed changes

AttrDocs.td fix

6d2bce7

Signed-off-by: Aleksander Fadeev <aleksander.fadeev@intel.com>

AlexeySachkov approved these changes Mar 23, 2020

View reviewed changes

romanovvlad merged commit 8eb588d into intel:sycl Mar 24, 2020

fadeeval deleted the private/fadeeval/Intel-reqd_work_group_size branch April 23, 2020 07:29

		@@ -2923,14 +2923,20 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
		return;

		uint32_t WGSize[3];

		let Heading = "GNU<reqd_work_group_size>, CXX11<cl::reqd_work_group_size,
		CXX11<intel::reqd_work_group_size>";

		(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see, for example, section
		6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification

[SYCL] Support intel::reqd_work_group_size #1328

[SYCL] Support intel::reqd_work_group_size #1328

Uh oh!

Conversation

fadeeval commented Mar 16, 2020

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!

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!

Fznamznon Mar 23, 2020 • 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!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

fadeeval Mar 18, 2020 • 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!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

erichkeane left a comment

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!

This comment was marked as resolved.

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!

erichkeane left a comment

Choose a reason for hiding this comment

Uh oh!

fadeeval commented Mar 18, 2020

Uh oh!

Uh oh!

Fznamznon Mar 23, 2020 •

edited

Loading

fadeeval Mar 18, 2020 •

edited

Loading

AlexeySachkov commented Mar 19, 2020 •

edited

Loading

Fznamznon Mar 19, 2020 •

edited

Loading