clang/AMDGPU: Respect __launch_bounds__ attribute (#215615)
Currently the HIP headers manually implement this with a
macro setting amdgpu attributes, and the proper clang attribute
is silently ignored. Directly map the proper attribute into
the target IR attributes. The first argument sets
"amdgpu-flat-work-group-size" and the second (reinterpreted by HIP
as minimum waves per EU) sets "amdgpu-waves-per-eu". An explicit
amdgpu_flat_work_group_size / amdgpu_waves_per_eu attribute takes
precedence. This matches the launch_bounds macro in the HIP headers,
which can now be dropped.
The 3rd maxclusterrank argument is only handled for NVPTX, so restrict
the sm_90 arch check to NVPTX targets and ignore the third argument on
other targets.
Fixes #91468
Co-authored-by: Claude (Claude-Opus-4.8) <noreply@anthropic.com>