[clang] [NVPTX] Add support for maxclusterrank in launch_bounds (PR #66496)

[Jakub Chlanda via cfe-commits]

================
@@ -11836,6 +11836,10 @@ def err_sycl_special_type_num_init_method : Error<
   "types with 'sycl_special_class' attribute must have one and only one '__init' "
   "method defined">;
 
+def warn_cuda_maxclusterrank_sm_90 : Warning<
+  "maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring "
+  "%1 attribute">, InGroup<IgnoredAttributes>;
----------------
jchlanda wrote:

The whole thing, this is analogous to how we currently handle:
```
__launch_bounds__(128, -2)
```

we issue a warning:
```
/home/dev/llvm/clang/test/SemaCUDA/launch_bounds_running_test.cu:5:24: warning: 'launch_bounds' attribute parameter 1 is negative and will be ignored [-Wcuda-compat]
    5 | __launch_bounds__(128, -2) void Test2Args(void);
      |                        ^~
/home/dev/llvm/clang/test/SemaCUDA/Inputs/cuda.h:14:61: note: expanded from macro '__launch_bounds__'
   14 | #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
      |                                                             ^~~~~~~~~~~
1 warning generated when compiling for host.
```
vs max cluster rank:
```
/home/dev/llvm/clang/test/SemaCUDA/launch_bounds_running_test.cu:5:27: warning: 'launch_bounds' attribute parameter 2 is negative and will be ignored [-Wcuda-compat]
    5 | __launch_bounds__(128, 2, -8) void Test2Args(void);
      |                           ^~
/home/dev/llvm/clang/test/SemaCUDA/Inputs/cuda.h:14:61: note: expanded from macro '__launch_bounds__'
   14 | #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
      |                                                             ^~~~~~~~~~~
1 warning generated when compiling for host.
```

and the resulting asm contains neither of the directives. 

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