[SYCL] Address Issue about Unable to Specify Kernel Properties when Reductions are Present

[HPS-1]

Basically to fix this issue. Note: seems that other than nd_launch, some variants (overloads) of parallel_for are also having similar problem and thus blocking the revision of the test case(s). Investigating on this.

[HPS-1]

@slawekptak Just to let you know that my work in this PR might address part of the problem you mentioned in https://github.com/intel/llvm/blame/40a9999b5536ac739e5f850c05d6d9dc1aa6e58b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp#L274, and thus pinging you to avoid repetitive work. Also feel free to correct me if I got anything wrong here. TY!

[slawekptak]

@slawekptak Just to let you know that my work in this PR might address part of the problem you mentioned in https://github.com/intel/llvm/blame/40a9999b5536ac739e5f850c05d6d9dc1aa6e58b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp#L274, and thus pinging you to avoid repetitive work. Also feel free to correct me if I got anything wrong here. TY!

@HPS-1 Thanks for letting me know. I'm adding @uditagarwal97 who works on properties support for the handler-less path. This patch might not be directly related, as this addresses the handler path, but it is good to be aware of this change.

[HPS-1]

@slawekptak Just to let you know that my work in this PR might address part of the problem you mentioned in https://github.com/intel/llvm/blame/40a9999b5536ac739e5f850c05d6d9dc1aa6e58b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp#L274, and thus pinging you to avoid repetitive work. Also feel free to correct me if I got anything wrong here. TY!

@HPS-1 Thanks for letting me know. I'm adding @uditagarwal97 who works on properties support for the handler-less path. This patch might not be directly related, as this addresses the handler path, but it is good to be aware of this change.

Thanks! And just a mention that I believe this is actually related to the handler-less path since the handler-less path implicitly calls the handler path sometimes: as you can see in the nd_launch implementation for the handler-less path:

template <typename KernelName = sycl::detail::auto_name, int Dimensions,
          typename KernelType, typename... ReductionsT>
void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
               ReductionsT &&...Reductions) {
  // TODO The handler-less path does not support reductions, kernel
  // function properties and kernel functions with the kernel_handler
  // type argument yet.
  if constexpr (sizeof...(ReductionsT) == 0 &&
                !(ext::oneapi::experimental::detail::
                      HasKernelPropertiesGetMethod<
                          const KernelType &>::value) &&
                !(detail::KernelLambdaHasKernelHandlerArgT<
                    KernelType, sycl::nd_item<Dimensions>>::value)) {
    detail::submit_kernel_direct_parallel_for<KernelName>(
        std::move(Q), empty_properties_t{}, Range, KernelObj);
  } else {
    submit(std::move(Q), [&](handler &CGH) {
      nd_launch<KernelName>(CGH, Range, KernelObj,
                            std::forward<ReductionsT>(Reductions)...);
    });
  }

Here in the if statement, it actually calls the nd_launch variant for the handler path when the kernel functor uses a get(properties_tag) to specifies its properties (i.e. when ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<const KernelType &>::value evaluates to true).