[SYCL] Lift restriction that prevents us from using free function kernel queries in device compilation mode (#17398)

Free function kernel queries can, at the moment, only be used in host
compilation mode requiring us to add `#ifndef __SYCL_DEVICE_ONLY__ `
directives to get code that uses them to compile. This PR removes this
restriction and the instances of the workaround used in our tests.

My understanding is the following:

The current implementation employs SFINAE using the predicate
`is_kernel_v<Func>` to delete at compile time queries that are
instantiated with a `Func` that is not a free function kernel. This is
certainly a reasonable interpretation of the spec which says :
`Constraints: Available only if is_kernel_v<Func> is true.` for each
query.

However, this causes a problem. During the device front-end compiler
phase of our toolchain, information is gathered regarding what the free
function kernels are as defined in the code provided by the user. This
means during this phase itself, `Func` will be treated as an invalid
free function kernel name since this information becomes known only
after the front-end has finished executing. And because SFINAE interacts
with the front-end mechanisms, we get a hard error in this phase of
compilation about a missing function.

One simple fix, proposed in this PR, is to simply define the trait
`is_kernel_v<Func>` to be constantly true in device compilation mode
regardless of `Func`.
This will ensure device front-end succeeds since a declaration of the
query exists and it still interprets the spec reasonably because if
`Func` does not represent a kernel, then during host compilation the
declaration will not be available because the `is_kernel` trait fails
which will print a nice message to the user saying what went wrong.

Author: lbushi25

sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp

@@ -29,12 +29,21 @@ inline constexpr bool is_single_task_kernel_v =
     is_single_task_kernel<Func>::value;
 
 template <auto *Func> struct is_kernel {
+// During device compilation mode the compiler does not yet know
+// what the kernels are named because that is exactly what its trying to
+// figure out during this phase. Therefore, we set the is_kernel trait to true
+// by default during device compilation in order to not get missing functions
+// errors.
+#ifdef __SYCL_DEVICE_ONLY__
+  static constexpr bool value = true;
+#else
   static constexpr bool value = false;
+#endif
 };
 
 template <auto *Func>
 inline constexpr bool is_kernel_v = is_kernel<Func>::value;
 
 } // namespace ext::oneapi::experimental
 } // namespace _V1
-} // namespace sycl
+} // namespace sycl

sycl/include/sycl/kernel_bundle.hpp

@@ -491,7 +491,6 @@ class kernel_bundle : public detail::kernel_bundle_plain,
     return detail::kernel_bundle_plain::ext_oneapi_has_kernel(name);
   }
 
-  // For free functions.
   template <auto *Func>
   std::enable_if_t<ext::oneapi::experimental::is_kernel_v<Func>, bool>
   ext_oneapi_has_kernel() {
@@ -838,7 +837,6 @@ bool has_kernel_bundle(const context &Ctx, const std::vector<device> &Devs) {
   return has_kernel_bundle<State>(Ctx, Devs, {get_kernel_id<KernelName>()});
 }
 
-// For free functions.
 namespace ext::oneapi::experimental {
 template <auto *Func, bundle_state State>
 std::enable_if_t<is_kernel_v<Func>, bool>
@@ -866,7 +864,6 @@ template <typename KernelName> bool is_compatible(const device &Dev) {
   return is_compatible({get_kernel_id<KernelName>()}, Dev);
 }
 
-// For free functions.
 namespace ext::oneapi::experimental {
 template <auto *Func>
 std::enable_if_t<is_kernel_v<Func>, bool> is_compatible(const device &Dev) {

sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp

@@ -6,6 +6,7 @@
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/backend/level_zero.hpp>
 #include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 #include <vector>
 
@@ -22,10 +23,6 @@ int main() {
   sycl::queue q;
   sycl::context ctxt = q.get_context();
   sycl::device d = ctxt.get_devices()[0];
-  // The following ifndef is required due to a number of limitations of free
-  // function kernels. See CMPLRLLVM-61498.
-  // TODO: Remove it once these limitations are no longer there.
-#ifndef __SYCL_DEVICE_ONLY__
   // First, run the kernel using the SYCL API.
   auto bundle = sycl::get_kernel_bundle<sycl::bundle_state::executable>(ctxt);
   sycl::kernel_id iota_id = syclexp::get_kernel_id<iota>();
@@ -90,5 +87,4 @@ int main() {
   assert(*ptr_twin == *ptr);
   sycl::free(ptr, q);
   sycl::free(ptr_twin, q);
-#endif
 }

sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp

@@ -8,6 +8,7 @@
 #include <sycl/detail/cl.h>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 #include <vector>
 
@@ -24,10 +25,6 @@ int main() {
   sycl::queue q;
   sycl::context ctxt = q.get_context();
   sycl::device d = ctxt.get_devices()[0];
-  // The following ifndef is required due to a number of limitations of free
-  // function kernels. See CMPLRLLVM-61498.
-  // TODO: Remove it once these limitations are no longer there.
-#ifndef __SYCL_DEVICE_ONLY__
   // First, run the kernel using the SYCL API.
 
   auto bundle = sycl::get_kernel_bundle<sycl::bundle_state::executable>(ctxt);
@@ -77,5 +74,4 @@ int main() {
   assert(*ptr_twin == *ptr);
   sycl::free(ptr, q);
   sycl::free(ptr_twin, q);
-#endif
 }

sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/free_function_kernels.cpp

@@ -14,6 +14,7 @@
 #include <iostream>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 using namespace sycl;
@@ -79,8 +80,6 @@ bool test_0(queue Queue) {
   std::cout << "Test 0a: " << (PassA ? "PASS" : "FAIL") << std::endl;
 
   bool PassB = false;
-  // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle =
       get_kernel_bundle<bundle_state::executable>(Queue.get_context());
   kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<ff_0>();
@@ -98,7 +97,6 @@ bool test_0(queue Queue) {
   std::cout << "Test 0b: " << (PassB ? "PASS" : "FAIL") << std::endl;
 
   free(usmPtr, Queue);
-#endif
   return PassA && PassB;
 }
 
@@ -131,8 +129,6 @@ bool test_1(queue Queue) {
   std::cout << "Test 1a: " << (PassA ? "PASS" : "FAIL") << std::endl;
 
   bool PassB = false;
-  // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle =
       get_kernel_bundle<bundle_state::executable>(Queue.get_context());
   kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<(
@@ -151,7 +147,6 @@ bool test_1(queue Queue) {
   std::cout << "Test 1b: " << (PassB ? "PASS" : "FAIL") << std::endl;
 
   free(usmPtr, Queue);
-#endif
   return PassA && PassB;
 }
 
@@ -189,8 +184,6 @@ bool test_2(queue Queue) {
   std::cout << "Test 2a: " << (PassA ? "PASS" : "FAIL") << std::endl;
 
   bool PassB = false;
-  // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle =
       get_kernel_bundle<bundle_state::executable>(Queue.get_context());
   kernel_id Kernel_id =
@@ -208,7 +201,6 @@ bool test_2(queue Queue) {
   std::cout << "Test 2b: " << (PassB ? "PASS" : "FAIL") << std::endl;
 
   free(usmPtr, Queue);
-#endif
   return PassA && PassB;
 }
 
@@ -250,8 +242,6 @@ bool test_3(queue Queue) {
   std::cout << "Test 3a: " << (PassA ? "PASS" : "FAIL") << std::endl;
 
   bool PassB = false;
-  // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle =
       get_kernel_bundle<bundle_state::executable>(Queue.get_context());
   kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<(
@@ -269,7 +259,6 @@ bool test_3(queue Queue) {
   std::cout << "Test 3b: " << (PassB ? "PASS" : "FAIL") << std::endl;
 
   free(usmPtr, Queue);
-#endif
   return PassA && PassB;
 }
 

sycl/test-e2e/DeviceImageDependencies/free_function_kernels.cpp

@@ -13,6 +13,7 @@
 #include <iostream>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 using namespace sycl;
@@ -78,8 +79,6 @@ bool test_0(queue Queue) {
   std::cout << "Test 0a: " << (PassA ? "PASS" : "FAIL") << std::endl;
 
   bool PassB = false;
-  // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle =
       get_kernel_bundle<bundle_state::executable>(Queue.get_context());
   kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<ff_0>();
@@ -97,7 +96,6 @@ bool test_0(queue Queue) {
   std::cout << "Test 0b: " << (PassB ? "PASS" : "FAIL") << std::endl;
 
   free(usmPtr, Queue);
-#endif
   return PassA && PassB;
 }
 
@@ -130,8 +128,6 @@ bool test_1(queue Queue) {
   std::cout << "Test 1a: " << (PassA ? "PASS" : "FAIL") << std::endl;
 
   bool PassB = false;
-  // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle =
       get_kernel_bundle<bundle_state::executable>(Queue.get_context());
   kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<(
@@ -150,7 +146,6 @@ bool test_1(queue Queue) {
   std::cout << "Test 1b: " << (PassB ? "PASS" : "FAIL") << std::endl;
 
   free(usmPtr, Queue);
-#endif
   return PassA && PassB;
 }
 
@@ -188,8 +183,6 @@ bool test_2(queue Queue) {
   std::cout << "Test 2a: " << (PassA ? "PASS" : "FAIL") << std::endl;
 
   bool PassB = false;
-  // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle =
       get_kernel_bundle<bundle_state::executable>(Queue.get_context());
   kernel_id Kernel_id =
@@ -207,7 +200,6 @@ bool test_2(queue Queue) {
   std::cout << "Test 2b: " << (PassB ? "PASS" : "FAIL") << std::endl;
 
   free(usmPtr, Queue);
-#endif
   return PassA && PassB;
 }
 
@@ -249,8 +241,6 @@ bool test_3(queue Queue) {
   std::cout << "Test 3a: " << (PassA ? "PASS" : "FAIL") << std::endl;
 
   bool PassB = false;
-  // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle =
       get_kernel_bundle<bundle_state::executable>(Queue.get_context());
   kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<(
@@ -268,7 +258,6 @@ bool test_3(queue Queue) {
   std::cout << "Test 3b: " << (PassB ? "PASS" : "FAIL") << std::endl;
 
   free(usmPtr, Queue);
-#endif
   return PassA && PassB;
 }
 

sycl/test-e2e/Graph/Inputs/free_function_kernels.cpp

@@ -21,7 +21,6 @@ int main() {
 
   Queue.memset(PtrA, 0, Size * sizeof(int)).wait();
 
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle = get_kernel_bundle<bundle_state::executable>(Ctxt);
   kernel_id Kernel_id = exp_ext::get_kernel_id<ff_0>();
   kernel Kernel = Bundle.get_kernel(Kernel_id);
@@ -38,7 +37,6 @@ int main() {
   for (size_t i = 0; i < Size; i++) {
     assert(HostDataA[i] == i);
   }
-#endif
   sycl::free(PtrA, Queue);
 
   return 0;

sycl/test-e2e/Graph/Inputs/work_group_memory_free_function.cpp

@@ -27,7 +27,6 @@ int main() {
 
   const size_t LocalSize = 128;
 
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle =
       get_kernel_bundle<bundle_state::executable>(Queue.get_context());
   kernel_id Kernel_id = exp_ext::get_kernel_id<ff_local_mem>();
@@ -57,7 +56,6 @@ int main() {
     int Ref = 10 + i + (Iterations * (i * 2));
     assert(check_value(i, Ref, HostData[i], "Ptr"));
   }
-#endif
 
   free(Ptr, Queue);
   return 0;

sycl/test-e2e/Graph/Update/FreeFunctionKernels/update_before_finalize.cpp

@@ -30,7 +30,6 @@ int main() {
 
   exp_ext::dynamic_parameter InputParam(Graph, PtrA);
 
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle = get_kernel_bundle<bundle_state::executable>(Ctxt);
   kernel_id Kernel_id = exp_ext::get_kernel_id<ff_0>();
   kernel Kernel = Bundle.get_kernel(Kernel_id);
@@ -52,7 +51,6 @@ int main() {
     assert(HostDataA[i] == 0);
     assert(HostDataB[i] == i);
   }
-#endif
   sycl::free(PtrA, Queue);
   sycl::free(PtrB, Queue);
 

sycl/test-e2e/Graph/Update/FreeFunctionKernels/update_with_indices_multiple_exec_graphs.cpp

@@ -31,7 +31,6 @@ int main() {
 
   exp_ext::dynamic_parameter InputParam(Graph, PtrA);
 
-#ifndef __SYCL_DEVICE_ONLY__
   kernel_bundle Bundle = get_kernel_bundle<bundle_state::executable>(Ctxt);
   kernel_id Kernel_id = exp_ext::get_kernel_id<ff_1>();
   kernel Kernel = Bundle.get_kernel(Kernel_id);
@@ -69,7 +68,6 @@ int main() {
     assert(HostDataA[i] == i * 3);
     assert(HostDataB[i] == i);
   }
-#endif
   sycl::free(PtrA, Queue);
   sycl::free(PtrB, Queue);