[SYCL][DOC] Allow captures in "if_device_has", etc (#8840)

Lift the restrictions in `if_device_has` and `if_architecture_is`
disallowing captures in the lamdba expressions. I originally added these
restrictions because I thought we could not implement these APIs in the
JIT case if they had captures. However, I have an outline for the JIT
design now, which will support captures.

Since we allow captures now, there isn't any need for the lambdas to
accept user-supplied parameters because the user can just capture
whatever variables they need. Therefore, this PR also removes support
for the parameters.

Author: gmlueck

sycl/doc/design/DeviceIf.md

@@ -207,7 +207,7 @@ listed architectures.  The following code snippet illustrates the technique:
 
 ```
 namespace sycl {
-namespace ext::oneapi::exprimental {
+namespace ext::oneapi::experimental {
 
 enum class architecture {
   x86_64,
@@ -217,7 +217,7 @@ enum class architecture {
   // ...
 };
 
-} // namespace ext::oneapi::exprimental
+} // namespace ext::oneapi::experimental
 
 namespace detail {
 
@@ -281,44 +281,43 @@ constexpr static bool device_architecture_is() {
 template<bool MakeCall>
 class if_architecture_helper {
  public:
-  template<ext::oneapi::exprimental::architecture ...Archs, typename T,
-           typename ...Args>
-  constexpr auto else_if_architecture_is(T fnTrue, Args ...args) {
+  template<ext::oneapi::experimental::architecture ...Archs, typename T>
+  constexpr auto else_if_architecture_is(T fnTrue) {
     if constexpr (MakeCall && device_architecture_is<Archs...>()) {
-      fnTrue(args...);
+      fnTrue();
       return if_architecture_helper<false>{};
     } else {
       return if_architecture_helper<MakeCall>{};
     }
   }
 
-  template<typename T, typename ...Args>
-  constexpr void otherwise(T fn, Args ...args) {
+  template<typename T>
+  constexpr void otherwise(T fn) {
     if constexpr (MakeCall) {
-      fn(args...);
+      fn();
     }
   }
 };
 
 } // namespace detail
 
-namespace ext::oneapi::exprimental {
+namespace ext::oneapi::experimental {
 
-template<architecture ...Archs, typename T, typename ...Args>
-constexpr static auto if_architecture_is(T fnTrue, Args ...args) {
+template<architecture ...Archs, typename T>
+constexpr static auto if_architecture_is(T fnTrue) {
   static_assert(detail::allowable_aot_mode<Archs...>(),
     "The if_architecture_is function may only be used when AOT "
     "compiling with '-fsycl-targets=spir64_x86_64' or "
     "'-fsycl-targets=intel_gpu_*'");
   if constexpr (detail::device_architecture_is<Archs...>()) {
-    fnTrue(args...);
+    fnTrue();
     return detail::if_architecture_helper<false>{};
   } else {
     return detail::if_architecture_helper<true>{};
   }
 }
 
-} // namespace ext::oneapi::exprimental
+} // namespace ext::oneapi::experimental
 } // namespace sycl
 ```
 

sycl/doc/extensions/experimental/sycl_ext_oneapi_device_architecture.asciidoc

@@ -485,8 +485,8 @@ code.  This function is not available in host code.
 ```
 namespace sycl::ext::oneapi::experimental {
 
-template<architecture ...Archs, typename ...Args, typename T>
-/* unspecified */ if_architecture_is(T fn, Args ...args);
+template<architecture ...Archs, typename T>
+/* unspecified */ if_architecture_is(T fn);
 
 } // namespace sycl::ext::oneapi::experimental
 ```
@@ -504,11 +504,11 @@ type, which provides the following member functions:
 ```
 class /* unspecified */ {
  public:
-  template<architecture ...Archs, typename ...Args, typename T>
-  /* unspecified */ else_if_architecture_is(T fn, Args ...args);
+  template<architecture ...Archs, typename T>
+  /* unspecified */ else_if_architecture_is(T fn);
 
-  template<typename T, typename ...Args>
-  void otherwise(T fn, Args ...args);
+  template<typename T>
+  void otherwise(T fn);
 };
 ```
 

sycl/doc/extensions/proposed/sycl_ext_oneapi_device_if.asciidoc

@@ -20,7 +20,7 @@
 == Notice
 
 [%hardbreaks]
-Copyright (C) 2021-2022 Intel Corporation.  All rights reserved.
+Copyright (C) 2021-2023 Intel Corporation.  All rights reserved.
 
 Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
 of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
@@ -135,20 +135,16 @@ code.  This function is not available in host code.
 ```
 namespace sycl::ext::oneapi::experimental {
 
-template<aspect ...Aspects, typename ...Args, typename T>
-/* unspecified */ if_device_has(T fn, Args ...args);
+template<aspect ...Aspects, typename T>
+/* unspecified */ if_device_has(T fn);
 
 } // namespace sycl::ext::oneapi::experimental
 ```
 
-The parameter `fn` must be a C++ `Callable` object which is invocable with the
-parameters specified in `args`.  Normal SYCL restrictions apply, so `fn` must
-not be a function pointer or a pointer to a member function.  (The expectation
-is that most applications will pass a lambda expression for this parameter.)
-
-If `fn` is a lambda expression, it must not have any captures.  Applications
-can work around this limitation by passing variables referenced by the lambda
-as function arguments.
+The parameter `fn` must be a C++ `Callable` object which is invocable with an
+empty parameter list.  Normal SYCL restrictions apply, so `fn` must not be a
+function pointer or a pointer to a member function.  (The expectation is that
+most applications will pass a lambda expression for this parameter.)
 
 The `Aspects` parameter pack identifies the condition that gates execution of
 the callable object `fn`.  This condition is `true` only if the device which
@@ -164,11 +160,11 @@ which provides the following member functions:
 ```
 class /* unspecified */ {
  public:
-  template<aspect ...Aspects, typename ...Args, typename T>
-  /* unspecified */ else_if_device_has(T fn, Args ...args);
+  template<aspect ...Aspects, typename T>
+  /* unspecified */ else_if_device_has(T fn);
 
-  template<typename T, typename ...Args>
-  void otherwise(T fn, Args ...args);
+  template<typename T>
+  void otherwise(T fn);
 };
 ```
 
@@ -214,10 +210,10 @@ void frob() {
 As specified above, the function `fn` may be discarded if the condition
 associated with the call to `if_device_has`, `else_if_device_has`, or
 `otherwise` is `false`.  More formally, this means that `fn` is potentially
-discarded (if `fn` is a function) or `+operator(Args...)+` of `fn` is
-potentially discarded (if `fn` is a callable object).  In addition, any other
-functions they call (and functions called by those functions etc.) are
-potentially discarded.
+discarded (if `fn` is a function) or `operator()()` of `fn` is potentially
+discarded (if `fn` is a callable object).  In addition, any other functions
+they call (and functions called by those functions etc.) are potentially
+discarded.
 
 These functions are discarded if all calls to them are reachable only from
 `if_device_has`, `else_if_device_has`, or `otherwise` whose associated