[clang] [llvm] [mlir] [AMDGPU] Verify AMDGPU required workgroup size matches flat workgroup size (PR #200989)

[Shilei Tian via cfe-commits]

================
@@ -51,6 +51,9 @@ __attribute__((amdgpu_waves_per_eu(0, 4))) kernel void kernel_waves_per_eu_0_4()
 __attribute__((amdgpu_flat_work_group_size(64, 32))) kernel void kernel_flat_work_group_size_64_32() {} // expected-error {{'amdgpu_flat_work_group_size' attribute argument is invalid: min must not be greater than max}}
 __attribute__((amdgpu_waves_per_eu(4, 2))) kernel void kernel_waves_per_eu_4_2() {} // expected-error {{'amdgpu_waves_per_eu' attribute argument is invalid: min must not be greater than max}}
 
+__attribute__((reqd_work_group_size(32, 2, 1), amdgpu_flat_work_group_size(16, 128))) kernel void kernel_reqd_work_group_size_32_2_1_flat_work_group_size_16_128() {} // expected-error {{'amdgpu_flat_work_group_size' attribute must match 'reqd_work_group_size' product}}
+__attribute__((amdgpu_flat_work_group_size(16, 128), reqd_work_group_size(32, 2, 1))) kernel void kernel_flat_work_group_size_16_128_reqd_work_group_size_32_2_1() {} // expected-error {{'amdgpu_flat_work_group_size' attribute must match 'reqd_work_group_size' product}}
+
----------------
shiltian wrote:

> As a side note, SYCL 2020 has the attribute as well, though under `sycl::`, which implementations may just use as an alias.

I'm not going to comment on what SYCL is relative to OpenCL. :-)

> Clang-based implementations of that may also prefer that `reqd_work_group_size` is properly maintained in the C++ path.

The definition of `reqd_work_group_size` clearly states that it is an OpenCL-related attribute.

```
def ReqdWorkGroupSize : InheritableAttr {
  // Does not have a [[]] spelling because it is an OpenCL-related attribute.
...
```

> I'd argue that we should add it to HIP and copy over the OpenCL semantics

> The attribute should be respected. We should not arbitrarily disable attributes based on the language.

> It has obvious semantics, and the good QoI would be to respect it. And merely warn if you specify Y/Z dimensions

I don't see it has obvious semantics, because it doesn't have any semantics defined anywhere else. If we really want to support that beyond OpenCL, it should be done properly and is also beyond the scope of this PR.

https://github.com/llvm/llvm-project/pull/200989