[SYCL][CUDA] No attribute form of __launch_bounds__ left: intel::max_work_group_size removed in #21864, and min_work_groups_per_cu / max_work_groups_per_mp reject kernel-lambda syntax #23322

Describe the bug

#8080 documents how to express CUDA's __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor, maxBlocksPerCluster) in SYCL with three attributes:

This causes two problems.

1. Removing intel::max_work_group_size silently drops the launch bound. It was an FPGA attribute, but on NVPTX it was also the documented way to emit maxntid. Existing code still compiles, with only an unknown-attribute warning. The bound is lost, and a kernel launched with 1024 work-items fails at run time:

   SYCL error: Exceeded the number of registers available on the hardware.
   	The number registers per work-group cannot exceed 65536 for this kernel on this device.
   	The kernel uses 254 registers per work-item for a total of 1024 work-items per work-group.
   

   The replacement is the max_work_group_size<...> kernel property from sycl_ext_oneapi_kernel_properties, returned from get(properties_tag). With it, the same kernel launches and maxntidx is emitted. However, neither the warning nor [SYCL][CUDA][HIP] Guidance for __launch_bounds__ in SYCL  #8080 points to the property.

2. intel::min_work_groups_per_cu and intel::max_work_groups_per_mp can't be written on a kernel lambda. AttrDocs.td says they apply to "a device function/lambda function". Unlike [[sycl::reqd_sub_group_size]] (and the removed intel::max_work_group_size), they don't set SupportsNonconformingLambdaSyntax, so the usual SYCL lambda placement is an error. There is also no kernel property for them. So a lambda kernel has no way to express the second and third __launch_bounds__ arguments, and the first one is only available as a property.

To reproduce

Compiler: intel/llvm sycl at fbf7d1fd2cbf40b346e8109a81b3a13cf87d7ebd (clang 23.0.0git).

Lambda form (lb.cpp):

#include <sycl/sycl.hpp>
int main() {
  sycl::queue q;
  q.submit([&](sycl::handler& h) {
    h.parallel_for(sycl::nd_range<1>(1024, 1024), [=](sycl::nd_item<1>)
        [[intel::max_work_group_size(1, 1, 1024), intel::min_work_groups_per_cu(1),
          intel::max_work_groups_per_mp(2)]] {});
  });
}

$ clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fsycl-device-only -S -emit-llvm lb.cpp -o lb.ll
lb.cpp:6:11: warning: unknown attribute 'intel::max_work_group_size' ignored [-Wunknown-attributes]
lb.cpp:6:51: error: 'intel::min_work_groups_per_cu' attribute cannot be applied to types
lb.cpp:7:11: error: 'intel::max_work_groups_per_mp' attribute cannot be applied to types

Functor form (lb2.cpp) works, but only when combined with the property:

#include <sycl/sycl.hpp>
namespace syclex = sycl::ext::oneapi::experimental;
struct K {
  [[intel::min_work_groups_per_cu(2), intel::max_work_groups_per_mp(4)]]
  void operator()(sycl::nd_item<1>) const {}
  auto get(syclex::properties_tag) const {
    return syclex::properties{syclex::max_work_group_size<256>};
  }
};
int main() {
  sycl::queue q;
  q.submit([&](sycl::handler& h) { h.parallel_for(sycl::nd_range<1>(256, 256), K{}); });
}

$ clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda -Xsycl-target-backend --cuda-gpu-arch=sm_90 \
    -fsycl-device-only -S -emit-llvm lb2.cpp -o lb2.ll
$ grep -oE '!"(maxntidx|minctasm|maxclusterrank)", i32 [0-9]+' lb2.ll | sort -u
!"maxclusterrank", i32 4
!"maxntidx", i32 256
!"minctasm", i32 2

Expected behavior

• Code that used [[intel::max_work_group_size]] as a launch bound should get a diagnostic that names the max_work_group_size kernel property as the replacement, rather than a generic unknown-attribute warning. Alternatively, the attribute could be kept for non-FPGA targets.
• intel::min_work_groups_per_cu and intel::max_work_groups_per_mp should accept the SYCL kernel-lambda syntax, as their documentation says. Alternatively, they could be exposed as kernel properties next to max_work_group_size, so that all of __launch_bounds__ can be expressed in one place.
• The guidance in [SYCL][CUDA][HIP] Guidance for __launch_bounds__ in SYCL  #8080 should be updated for the post-[SYCL] PR 8 - Remove FPGA Attributes from SYCL FE #21864 state.

Environment

• OS: Linux x86_64 (RHEL 8)
• Target: NVIDIA A100 (sm_80) at run time; sm_90 for the maxclusterrank check above
• DPC++: intel/llvm sycl branch, fbf7d1fd2cbf40b346e8109a81b3a13cf87d7ebd, built from source

Additional context

Found while porting a CUDA kernel that uses __launch_bounds__(1024) to SYCL. Switching to get(properties_tag) returning max_work_group_size<1024> fixed the launch on the A100 and did not change AMD (gfx942) performance.