[SYCL] Refactor reqd_work_group_size attribute implementation

[smanna12]

This patch

1. refactors SYCL kernel function attribute reqd_work_group_size to better fit for community standards.

2. refactors the way we handled duplicate attributes logic with when present on a given declaration.

3. refactors the way we handled mutually exclusive attributes logic with when present on a given declaration with
   reqd_work_group_size, max_work_group_size, max_global_work_dim, and num_simd_work_items attributes.

4. handles redeclarations or template instantiations properly.

5. stores ConstantExpr into the semantic attribute rather than calling getIntegerConstantExpr() for attribute dimension values.

6. fixes crash with optional arguments on ReqdWorkGroupAttr and adds missing diagnostics for merging cases when present
   on a given declaration with reqd_work_group_size, max_work_group_size, max_global_work_dim, and num_simd_work_items
   attributes.

7. updates clang-tidy build after reqd_work_group_size attribute changes.

8. adds tests.

This patch fixes
#5736, #5735, #5734, #5732, #5731, #3223, #3361, #3175, #3742, and #5733

Signed-off-by: Soumi Manna [email protected]

[AaronBallman]

Looks like it's heading in the right direction at a high level. One thing that would be nice (but perhaps not easy to do) would be to see how much attribute checking code can be refactored to share between the Add and Merge functions here. They're doing the same things notionally (as best I saw), so it'd be nice to only write that logic once if possible. However, that might require more refactoring that's better left until later.

[smanna12]

One thing that would be nice (but perhaps not easy to do) would be to see how much attribute checking code can be refactored to share between the Add and Merge functions here. They're doing the same things notionally (as best I saw), so it'd be nice to only write that logic once if possible. However, that might require more refactoring that's better left until later.

Agreed. Existing implementation of "reqd_work_group_size" is very messy condition. This refactorization will fix several missing diagnostics and compilation crashes. I would like to do this (refactor the login once) in a separate PR. Not easy way to do this. @AaronBallman, would it be OK ?

[smanna12]

Sorry for the force push. I was having issues to push my changes due to recent pulldown.

[elizabethandrews]

LGTM. Thanks!

Just wanted to point out that #3175 might still need more work. I see tests with FIXME for default value still. So we probably should not close this bug till we're sure of optional argument behavior.

// Same for the default values.
// FIXME: This turns out to be wrong as there aren't really default values
// (that is an implementation detail we use but shouldn't expose to the user).
// Instead, the dimensionality of the attribute needs to match that of the
// kernel, so the one, two, and three arg forms of the attribute are actually
// *different* attributes. This means that you should not be able to redeclare
// the function with a different dimensionality.
[[sycl::reqd_work_group_size(4)]] void four_again();
[[sycl::reqd_work_group_size(4)]] void four_again(); // OK
[[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK
[[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK
[[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK
[[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK

[smanna12]

LGTM. Thanks!

Just wanted to point out that #3175 might still need more work. I see tests with FIXME for default value still. So we probably should not close this bug till we're sure of optional argument behavior.

// Same for the default values.
// FIXME: This turns out to be wrong as there aren't really default values
// (that is an implementation detail we use but shouldn't expose to the user).
// Instead, the dimensionality of the attribute needs to match that of the
// kernel, so the one, two, and three arg forms of the attribute are actually
// *different* attributes. This means that you should not be able to redeclare
// the function with a different dimensionality.
[[sycl::reqd_work_group_size(4)]] void four_again();
[[sycl::reqd_work_group_size(4)]] void four_again(); // OK
[[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK
[[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK
[[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK
[[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK

Thanks for the reviews!
This issues is covered by #3743

Just wanted to point out that #3175 might still need more work

This is related to NULL value check issue and assertion, so this is covered by current patch.

[smanna12]

ping @premanandrao and @AaronBallman

[premanandrao]

All minor nits. Changes are getting way better @smanna12!

[smanna12]

@AaronBallman, are you OK with the patch? Do you have any more concerns? Thank you

[AaronBallman]

@AaronBallman, are you OK with the patch? Do you have any more concerns? Thank you

Nope, this LGTM (I didn't spot anything with my quick once over from today).

[smanna12]

Thank you everyone for the reviews!

[keryell]

We should not push new symmetrical examples/test like this.
Any mixing X, Y and Z would just pass the test.