[SYCL] Add default argument support for work_group_size_hint attribute

[Michoumichmich]

This commit adds support for default arguments in [[sycl::work_group_size_hint]] according to the specification (5.8.1 Kernel attributes, table 180).

Signed-off by: Michel Migdal mm@mgdl.fr

[AaronBallman]

This is missing all of its test coverage.

[AaronBallman]

This still needs test coverage for the changes. In addition to the functional change to making the arguments optional, we should have some tests for declaration merging. e.g.,

[[sycl::work_group_size_hint(1, 2, 3)]] void func();
[[sycl::work_group_size_hint(1)]] void func() {} // Should still have 1, 2, 3, right?

[Michoumichmich]

This still needs test coverage for the changes. In addition to the functional change to making the arguments optional, we should have some tests for declaration merging. e.g.,

I wasn't sure whether I was "allowed" to add tests for myself (as I was surprised that there were no prior tests), but I will.

[[sycl::work_group_size_hint(1, 2, 3)]] void func();
[[sycl::work_group_size_hint(1)]] void func() {} // Should still have 1, 2, 3, right?

Well, according to this comment:

llvm/clang/test/SemaSYCL/reqd-work-group-size-device.cpp

Line 29 in bc3cf09

// *different* attributes. This means that you should not be able to redeclare

this should be an error. No? (Furthermore it suggests that the dimensionality should match the kernel. But currently this is not diagnosed, but it's better like that). Maybe I will wait for #4224 to be resolved (dimX and dimZ are std::swapped in the implementation, but it's not done on this attribute)

[AaronBallman]

This still needs test coverage for the changes. In addition to the functional change to making the arguments optional, we should have some tests for declaration merging. e.g.,

I wasn't sure whether I was "allowed" to add tests for myself (as I was surprised that there were no prior tests), but I will.

Oh, we definitely love new test coverage! :-)

[[sycl::work_group_size_hint(1, 2, 3)]] void func();
[[sycl::work_group_size_hint(1)]] void func() {} // Should still have 1, 2, 3, right?

Well, according to this comment:

llvm/clang/test/SemaSYCL/reqd-work-group-size-device.cpp

Line 29 in bc3cf09

// *different* attributes. This means that you should not be able to redeclare

this should be an error. No? (Furthermore it suggests that the dimensionality should match the kernel. But currently this is not diagnosed, but it's better like that). Maybe I will wait for #4224 to be resolved (dimX and dimZ are std::swapped in the implementation, but it's not done on this attribute)

Oh, great catch, I had completely forgotten about that really weird behavior. That should be an error situation. I'd recommend adding it as a test case with the same set of FIXME comments associated with it.

[AaronBallman]

I'm adding some more reviewers for visibility.

[smanna12]

@Michoumichmich, could you please update your PR description? Thanks

[Michoumichmich]

@Michoumichmich, could you please update your PR description? Thanks

Yes, sorry, it's up-to-date now!

[Michoumichmich]

I added the workaround. However the same one should be applied reqd_work_group_size, but it never came up because tests are missing for a case like: __attribute__((reqd_work_group_size(1, 1, 1, 1))) in SemaOpenCL.

[Michoumichmich]

I also updated the documentation for work_group_size_hint and removed the part about handling reqd_work_group_size with num_simd_work_items (that was misplaced I guess and may already be found in the doc for intel::num_simd_work_items).

[premanandrao]

The test file work-group-size-hint-host.cpp is misnamed. I would

[Fznamznon]

LGTM

[smanna12]

LGTM

[Michoumichmich]

Hum, that's surprising, check-clang-semasycl passed successfully on my side

[bader]

Hum, that's surprising, check-clang-semasycl passed successfully on my side

It looks like your changes imply that PropagateAndDiagnoseDeviceAttr should be updated as well. Have you uploaded all the changes?

[Michoumichmich]

Hum, that's surprising, check-clang-semasycl passed successfully on my side

It looks like your changes imply that PropagateAndDiagnoseDeviceAttr should be updated as well. Have you uploaded all the changes?

@bader @smanna12 thanks for the help, I fixed it, I believe that llvm_unreachable did not trigger the error because I built in release

[bader]

Okay.
NOTE: you can use Release for CMAKE_BUILD_TYPE and have assertions enabled LLVM_ENABLE_ASSERTIONS at the same time.

[bader]

Pre-commit failures are not related to this patch and being investigated. @AaronBallman, please, check if you are okay with the changes.

[AaronBallman]

LGTM aside from some typos!

[AaronBallman]

LGTM!