From 1000e332405a63523111f4a3b7fc11a47faf5888 Mon Sep 17 00:00:00 2001
From: Alexander Flegontov <alexander.flegontov@intel.com>
Date: Mon, 27 Jul 2020 10:44:58 +0300
Subject: [PATCH 1/3] [SYCL] Fix not found kernel due to empty kernel name when
 using set_arg(s)

Signed-off-by: Alexander Flegontov <alexander.flegontov@intel.com>
---
 sycl/include/CL/sycl/handler.hpp | 26 +++++++++++++++++---------
 1 file changed, 17 insertions(+), 9 deletions(-)

diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index da60f8258a0bb..1f740b1b15f9f 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -726,6 +726,7 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(std::move(NumWorkItems));
     MCGType = detail::CG::KERNEL;
     extractArgsAndReqs();
+    MKernelName = getKernelName();
   }
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -1185,6 +1186,7 @@ class __SYCL_EXPORT handler {
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     MCGType = detail::CG::KERNEL;
     extractArgsAndReqs();
+    MKernelName = getKernelName();
   }
 
   void parallel_for(range<1> NumWorkItems, kernel Kernel) {
@@ -1218,6 +1220,7 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
     MCGType = detail::CG::KERNEL;
     extractArgsAndReqs();
+    MKernelName = getKernelName();
   }
 
   /// Defines and invokes a SYCL kernel function for the specified range and
@@ -1238,6 +1241,7 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(std::move(NDRange));
     MCGType = detail::CG::KERNEL;
     extractArgsAndReqs();
+    MKernelName = getKernelName();
   }
 
   /// Defines and invokes a SYCL kernel function.
@@ -1260,9 +1264,10 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(range<1>{1});
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     MCGType = detail::CG::KERNEL;
-    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>())
+    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
       extractArgsAndReqs();
-    else
+      MKernelName = getKernelName();
+    } else
       StoreLambda<NameT, KernelType, /*Dims*/ 0, void>(std::move(KernelFunc));
 #endif
   }
@@ -1300,9 +1305,10 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(std::move(NumWorkItems));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     MCGType = detail::CG::KERNEL;
-    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>())
+    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
       extractArgsAndReqs();
-    else
+      MKernelName = getKernelName();
+    } else
       StoreLambda<NameT, KernelType, Dims, LambdaArgType>(
           std::move(KernelFunc));
 #endif
@@ -1336,9 +1342,10 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     MCGType = detail::CG::KERNEL;
-    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>())
+    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
       extractArgsAndReqs();
-    else
+      MKernelName = getKernelName();
+    } else
       StoreLambda<NameT, KernelType, Dims, LambdaArgType>(
           std::move(KernelFunc));
 #endif
@@ -1373,9 +1380,10 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(std::move(NDRange));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     MCGType = detail::CG::KERNEL;
-    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>())
+    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
       extractArgsAndReqs();
-    else
+      MKernelName = getKernelName();
+    } else
       StoreLambda<NameT, KernelType, Dims, LambdaArgType>(
           std::move(KernelFunc));
 #endif
@@ -1791,7 +1799,7 @@ class __SYCL_EXPORT handler {
   unique_ptr_class<detail::HostKernelBase> MHostKernel;
   /// Storage for lambda/function when using HostTask
   unique_ptr_class<detail::HostTask> MHostTask;
-  detail::OSModuleHandle MOSModuleHandle;
+  detail::OSModuleHandle MOSModuleHandle = detail::OSUtil::ExeModuleHandle;
   // Storage for a lambda or function when using InteropTasks
   std::unique_ptr<detail::InteropTask> MInteropTask;
   /// The list of events that order this operation.

From 2ee73dbbd89328b80157e390e9a4e3cf70e319c8 Mon Sep 17 00:00:00 2001
From: Alexander Flegontov <alexander.flegontov@intel.com>
Date: Tue, 28 Jul 2020 23:25:39 +0300
Subject: [PATCH 2/3] [SYCL] Added test using set_args with handler API

Signed-off-by: Alexander Flegontov <alexander.flegontov@intel.com>
---
 .../basic_tests/handler/handler_set_args.cpp  | 255 ++++++++++++++++++
 1 file changed, 255 insertions(+)
 create mode 100644 sycl/test/basic_tests/handler/handler_set_args.cpp

diff --git a/sycl/test/basic_tests/handler/handler_set_args.cpp b/sycl/test/basic_tests/handler/handler_set_args.cpp
new file mode 100644
index 0000000000000..3589e36c0c27f
--- /dev/null
+++ b/sycl/test/basic_tests/handler/handler_set_args.cpp
@@ -0,0 +1,255 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//==--------------- handler_set_args.cpp -------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#include <CL/sycl.hpp>
+#include <cassert>
+
+struct use_offset {
+  static const int no = 0;
+  static const int yes = 1;
+};
+
+using accessor_t =
+    cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                       cl::sycl::access::target::global_buffer>;
+
+struct single_task_functor {
+  single_task_functor(accessor_t acc) : acc(acc) {}
+
+  void operator()() { acc[0] = 10; }
+
+  accessor_t acc;
+};
+
+struct single_task_new_functor {
+  single_task_new_functor(accessor_t acc) : acc(acc) {}
+
+  void operator()() { acc[0] = 10; }
+
+  accessor_t acc;
+};
+
+template <int useOffset> struct parallel_for_range_id_functor {
+  parallel_for_range_id_functor(accessor_t acc) : acc(acc) {}
+
+  void operator()(cl::sycl::id<1> id) { acc[0] = 10; }
+
+  accessor_t acc;
+};
+
+template <int useOffset> struct parallel_for_range_item_functor {
+  parallel_for_range_item_functor(accessor_t acc) : acc(acc) {}
+
+  void operator()(cl::sycl::item<1> item) { acc[0] = 10; }
+
+  accessor_t acc;
+};
+
+struct parallel_for_nd_range_functor {
+  parallel_for_nd_range_functor(accessor_t acc) : acc(acc) {}
+
+  void operator()(cl::sycl::nd_item<1> ndItem) { acc[0] = 10; }
+
+  accessor_t acc;
+};
+
+template <class kernel_name>
+cl::sycl::kernel get_prebuilt_kernel(cl::sycl::queue &queue) {
+  cl::sycl::program program(queue.get_context());
+  program.build_with_kernel_type<kernel_name>();
+  return program.get_kernel<kernel_name>();
+}
+
+const cl::sycl::range<1> range = 1;
+
+template <class kernel_wrapper>
+void check_api_call(cl::sycl::queue &queue, kernel_wrapper &&kernelWrapper) {
+  int result = 0;
+  {
+    auto buf = cl::sycl::buffer<int, 1>(&result, range);
+    queue.submit([&](cl::sycl::handler &cgh) {
+      auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+      kernelWrapper(cgh, acc);
+    });
+  }
+  assert(result == 10);
+}
+
+int main() {
+  cl::sycl::queue queue;
+  const cl::sycl::id<1> offset(0);
+
+  const cl::sycl::nd_range<1> ndRange(range, range);
+
+  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    cgh.single_task(single_task_functor(acc));
+  });
+
+  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    cgh.parallel_for(range, parallel_for_range_id_functor<use_offset::no>(acc));
+  });
+
+  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    cgh.parallel_for(range, offset,
+                     parallel_for_range_id_functor<use_offset::yes>(acc));
+  });
+
+  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    cgh.parallel_for(range,
+                     parallel_for_range_item_functor<use_offset::no>(acc));
+  });
+
+  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    cgh.parallel_for(range, offset,
+                     parallel_for_range_item_functor<use_offset::yes>(acc));
+  });
+
+  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    cgh.parallel_for(ndRange, parallel_for_nd_range_functor(acc));
+  });
+
+  {
+    auto preBuiltKernel = get_prebuilt_kernel<single_task_functor>(queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.single_task(preBuiltKernel);
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::no>>(
+            queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for(range, preBuiltKernel);
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::yes>>(
+            queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for(range, offset, preBuiltKernel);
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::no>>(
+            queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for(range, preBuiltKernel);
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::yes>>(
+            queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for(range, offset, preBuiltKernel);
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_nd_range_functor>(queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for(ndRange, preBuiltKernel);
+    });
+  }
+
+  {
+    auto preBuiltKernel = get_prebuilt_kernel<single_task_functor>(queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.single_task<class other_kernel_name1>(preBuiltKernel,
+                                                [=]() { acc[0] = 10; });
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::no>>(
+            queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for<class other_kernel_name2>(
+          preBuiltKernel, range, [=](cl::sycl::id<1> id) { acc[0] = 10; });
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::yes>>(
+            queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for<class other_kernel_name3>(
+          preBuiltKernel, range, offset,
+          [=](cl::sycl::id<1> id) { acc[0] = 10; });
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::no>>(
+            queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for<class other_kernel_name4>(
+          preBuiltKernel, range, [=](cl::sycl::item<1> item) { acc[0] = 10; });
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::yes>>(
+            queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for<class other_kernel_name5>(
+          preBuiltKernel, range, offset,
+          [=](cl::sycl::item<1> item) { acc[0] = 10; });
+    });
+  }
+
+  {
+    auto preBuiltKernel =
+        get_prebuilt_kernel<parallel_for_nd_range_functor>(queue);
+
+    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+      cgh.set_args(acc);
+      cgh.parallel_for<class other_kernel_name6>(
+          preBuiltKernel, ndRange,
+          [=](cl::sycl::nd_item<1> ndItem) { acc[0] = 10; });
+    });
+  }
+
+  return 0;
+}

From 805d8a352e3babbb210784e1416a7e526e3b08eb Mon Sep 17 00:00:00 2001
From: Alexander Flegontov <alexander.flegontov@intel.com>
Date: Wed, 29 Jul 2020 13:59:26 +0300
Subject: [PATCH 3/3] [SYCL] Fix coding style issue

Signed-off-by: Alexander Flegontov <alexander.flegontov@intel.com>
---
 .../basic_tests/handler/handler_set_args.cpp  | 189 ++++++++----------
 1 file changed, 82 insertions(+), 107 deletions(-)

diff --git a/sycl/test/basic_tests/handler/handler_set_args.cpp b/sycl/test/basic_tests/handler/handler_set_args.cpp
index 3589e36c0c27f..dcb613ec33627 100644
--- a/sycl/test/basic_tests/handler/handler_set_args.cpp
+++ b/sycl/test/basic_tests/handler/handler_set_args.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
@@ -12,69 +12,57 @@
 #include <CL/sycl.hpp>
 #include <cassert>
 
-struct use_offset {
-  static const int no = 0;
-  static const int yes = 1;
-};
-
-using accessor_t =
-    cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
-                       cl::sycl::access::target::global_buffer>;
-
-struct single_task_functor {
-  single_task_functor(accessor_t acc) : acc(acc) {}
+constexpr bool UseOffset = true;
+constexpr bool NoOffset = false;
+const cl::sycl::range<1> Range = 1;
 
-  void operator()() { acc[0] = 10; }
-
-  accessor_t acc;
-};
+using AccessorT = cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                                     cl::sycl::access::target::global_buffer>;
 
-struct single_task_new_functor {
-  single_task_new_functor(accessor_t acc) : acc(acc) {}
+struct SingleTaskFunctor {
+  SingleTaskFunctor(AccessorT acc) : MAcc(acc) {}
 
-  void operator()() { acc[0] = 10; }
+  void operator()() { MAcc[0] = 10; }
 
-  accessor_t acc;
+  AccessorT MAcc;
 };
 
-template <int useOffset> struct parallel_for_range_id_functor {
-  parallel_for_range_id_functor(accessor_t acc) : acc(acc) {}
+template <bool useOffset> struct ParallelForRangeIdFunctor {
+  ParallelForRangeIdFunctor(AccessorT acc) : MAcc(acc) {}
 
-  void operator()(cl::sycl::id<1> id) { acc[0] = 10; }
+  void operator()(cl::sycl::id<1> id) { MAcc[0] = 10; }
 
-  accessor_t acc;
+  AccessorT MAcc;
 };
 
-template <int useOffset> struct parallel_for_range_item_functor {
-  parallel_for_range_item_functor(accessor_t acc) : acc(acc) {}
+template <bool useOffset> struct ParallelForRangeItemFunctor {
+  ParallelForRangeItemFunctor(AccessorT acc) : MAcc(acc) {}
 
-  void operator()(cl::sycl::item<1> item) { acc[0] = 10; }
+  void operator()(cl::sycl::item<1> item) { MAcc[0] = 10; }
 
-  accessor_t acc;
+  AccessorT MAcc;
 };
 
-struct parallel_for_nd_range_functor {
-  parallel_for_nd_range_functor(accessor_t acc) : acc(acc) {}
+struct ParallelForNdRangeFunctor {
+  ParallelForNdRangeFunctor(AccessorT acc) : MAcc(acc) {}
 
-  void operator()(cl::sycl::nd_item<1> ndItem) { acc[0] = 10; }
+  void operator()(cl::sycl::nd_item<1> ndItem) { MAcc[0] = 10; }
 
-  accessor_t acc;
+  AccessorT MAcc;
 };
 
 template <class kernel_name>
-cl::sycl::kernel get_prebuilt_kernel(cl::sycl::queue &queue) {
+cl::sycl::kernel getPrebuiltKernel(cl::sycl::queue &queue) {
   cl::sycl::program program(queue.get_context());
   program.build_with_kernel_type<kernel_name>();
   return program.get_kernel<kernel_name>();
 }
 
-const cl::sycl::range<1> range = 1;
-
 template <class kernel_wrapper>
-void check_api_call(cl::sycl::queue &queue, kernel_wrapper &&kernelWrapper) {
+void checkApiCall(cl::sycl::queue &queue, kernel_wrapper &&kernelWrapper) {
   int result = 0;
   {
-    auto buf = cl::sycl::buffer<int, 1>(&result, range);
+    auto buf = cl::sycl::buffer<int, 1>(&result, Range);
     queue.submit([&](cl::sycl::handler &cgh) {
       auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
       kernelWrapper(cgh, acc);
@@ -84,42 +72,39 @@ void check_api_call(cl::sycl::queue &queue, kernel_wrapper &&kernelWrapper) {
 }
 
 int main() {
-  cl::sycl::queue queue;
-  const cl::sycl::id<1> offset(0);
-
-  const cl::sycl::nd_range<1> ndRange(range, range);
+  cl::sycl::queue Queue;
+  const cl::sycl::id<1> Offset(0);
+  const cl::sycl::nd_range<1> NdRange(Range, Range);
 
-  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
-    cgh.single_task(single_task_functor(acc));
+  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
+    cgh.single_task(SingleTaskFunctor(acc));
   });
 
-  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
-    cgh.parallel_for(range, parallel_for_range_id_functor<use_offset::no>(acc));
+  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
+    cgh.parallel_for(Range, ParallelForRangeIdFunctor<NoOffset>(acc));
   });
 
-  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
-    cgh.parallel_for(range, offset,
-                     parallel_for_range_id_functor<use_offset::yes>(acc));
+  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
+    cgh.parallel_for(Range, Offset, ParallelForRangeIdFunctor<UseOffset>(acc));
   });
 
-  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
-    cgh.parallel_for(range,
-                     parallel_for_range_item_functor<use_offset::no>(acc));
+  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
+    cgh.parallel_for(Range, ParallelForRangeItemFunctor<NoOffset>(acc));
   });
 
-  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
-    cgh.parallel_for(range, offset,
-                     parallel_for_range_item_functor<use_offset::yes>(acc));
+  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
+    cgh.parallel_for(Range, Offset,
+                     ParallelForRangeItemFunctor<UseOffset>(acc));
   });
 
-  check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
-    cgh.parallel_for(ndRange, parallel_for_nd_range_functor(acc));
+  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
+    cgh.parallel_for(NdRange, ParallelForNdRangeFunctor(acc));
   });
 
   {
-    auto preBuiltKernel = get_prebuilt_kernel<single_task_functor>(queue);
+    auto preBuiltKernel = getPrebuiltKernel<SingleTaskFunctor>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
       cgh.single_task(preBuiltKernel);
     });
@@ -127,126 +112,116 @@ int main() {
 
   {
     auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::no>>(
-            queue);
+        getPrebuiltKernel<ParallelForRangeIdFunctor<NoOffset>>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for(range, preBuiltKernel);
+      cgh.parallel_for(Range, preBuiltKernel);
     });
   }
 
   {
     auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::yes>>(
-            queue);
+        getPrebuiltKernel<ParallelForRangeIdFunctor<UseOffset>>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for(range, offset, preBuiltKernel);
+      cgh.parallel_for(Range, Offset, preBuiltKernel);
     });
   }
 
   {
     auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::no>>(
-            queue);
+        getPrebuiltKernel<ParallelForRangeItemFunctor<NoOffset>>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for(range, preBuiltKernel);
+      cgh.parallel_for(Range, preBuiltKernel);
     });
   }
 
   {
     auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::yes>>(
-            queue);
+        getPrebuiltKernel<ParallelForRangeItemFunctor<UseOffset>>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for(range, offset, preBuiltKernel);
+      cgh.parallel_for(Range, Offset, preBuiltKernel);
     });
   }
 
   {
-    auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_nd_range_functor>(queue);
+    auto preBuiltKernel = getPrebuiltKernel<ParallelForNdRangeFunctor>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for(ndRange, preBuiltKernel);
+      cgh.parallel_for(NdRange, preBuiltKernel);
     });
   }
 
   {
-    auto preBuiltKernel = get_prebuilt_kernel<single_task_functor>(queue);
+    auto preBuiltKernel = getPrebuiltKernel<SingleTaskFunctor>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.single_task<class other_kernel_name1>(preBuiltKernel,
-                                                [=]() { acc[0] = 10; });
+      cgh.single_task<class OtherKernelName1>(preBuiltKernel,
+                                              [=]() { acc[0] = 10; });
     });
   }
 
   {
     auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::no>>(
-            queue);
+        getPrebuiltKernel<ParallelForRangeIdFunctor<NoOffset>>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for<class other_kernel_name2>(
-          preBuiltKernel, range, [=](cl::sycl::id<1> id) { acc[0] = 10; });
+      cgh.parallel_for<class OtherKernelName2>(
+          preBuiltKernel, Range, [=](cl::sycl::id<1> id) { acc[0] = 10; });
     });
   }
 
   {
     auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_range_id_functor<use_offset::yes>>(
-            queue);
+        getPrebuiltKernel<ParallelForRangeIdFunctor<UseOffset>>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for<class other_kernel_name3>(
-          preBuiltKernel, range, offset,
+      cgh.parallel_for<class OtherKernelName3>(
+          preBuiltKernel, Range, Offset,
           [=](cl::sycl::id<1> id) { acc[0] = 10; });
     });
   }
 
   {
     auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::no>>(
-            queue);
+        getPrebuiltKernel<ParallelForRangeItemFunctor<NoOffset>>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for<class other_kernel_name4>(
-          preBuiltKernel, range, [=](cl::sycl::item<1> item) { acc[0] = 10; });
+      cgh.parallel_for<class OtherKernelName4>(
+          preBuiltKernel, Range, [=](cl::sycl::item<1> item) { acc[0] = 10; });
     });
   }
 
   {
     auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_range_item_functor<use_offset::yes>>(
-            queue);
+        getPrebuiltKernel<ParallelForRangeItemFunctor<UseOffset>>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for<class other_kernel_name5>(
-          preBuiltKernel, range, offset,
+      cgh.parallel_for<class OtherKernelName5>(
+          preBuiltKernel, Range, Offset,
           [=](cl::sycl::item<1> item) { acc[0] = 10; });
     });
   }
 
   {
-    auto preBuiltKernel =
-        get_prebuilt_kernel<parallel_for_nd_range_functor>(queue);
+    auto preBuiltKernel = getPrebuiltKernel<ParallelForNdRangeFunctor>(Queue);
 
-    check_api_call(queue, [&](cl::sycl::handler &cgh, accessor_t acc) {
+    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
       cgh.set_args(acc);
-      cgh.parallel_for<class other_kernel_name6>(
-          preBuiltKernel, ndRange,
+      cgh.parallel_for<class OtherKernelName6>(
+          preBuiltKernel, NdRange,
           [=](cl::sycl::nd_item<1> ndItem) { acc[0] = 10; });
     });
   }