Address more feedback

Author: lbushi25

sycl/test-e2e/FreeFunctionKernels/free_function_kernel_local_memory.cpp

@@ -10,12 +10,12 @@
 // either by way of the work group scratch memory extension or the work group
 // static memory extension.
 
-#include <sycl/ext/oneapi/work_group_static.hpp>
-
 #include "helpers.hpp"
+
 #include <cassert>
 #include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
 #include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/ext/oneapi/work_group_static.hpp>
 #include <sycl/group_barrier.hpp>
 #include <sycl/usm.hpp>
 
@@ -25,36 +25,38 @@ namespace syclexp = sycl::ext::oneapi::experimental;
 constexpr int SIZE = 16;
 
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
-void scratch_kernel(float *src, float *dst) {
+void scratchKernel(float *src, float *dst) {
   size_t lid = syclext::this_work_item::get_nd_item<1>().get_local_linear_id();
-  float *local_mem = (float *)syclexp::get_work_group_scratch_memory();
-  local_mem[lid] = 2 * src[lid];
-  dst[lid] = local_mem[lid];
+  float *localMem =
+      reinterpret_cast<float *>(syclexp::get_work_group_scratch_memory());
+  localMem[lid] = 2 * src[lid];
+  dst[lid] = localMem[lid];
 }
 
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
-void static_kernel(float *src, float *dst) {
+void staticKernel(float *src, float *dst) {
   sycl::nd_item<1> item = syclext::this_work_item::get_nd_item<1>();
   size_t lid = item.get_local_linear_id();
-  syclexp::work_group_static<float[SIZE]> local_mem;
-  local_mem[lid] = src[lid] * src[lid];
+  syclexp::work_group_static<float[SIZE]> localMem;
+  localMem[lid] = src[lid] * src[lid];
   sycl::group_barrier(item.get_group());
   if (item.get_group().leader()) { // Check that memory is indeed shared between
-                                   // the work group
+                                   // the work group.
     for (int i = 0; i < SIZE; ++i)
-      assert(local_mem[i] == src[i] * src[i]);
+      assert(localMem[i] == src[i] * src[i]);
   }
-  dst[lid] = local_mem[lid];
+  dst[lid] = localMem[lid];
 }
 
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
-void scratch_static_kernel(float *src, float *dst) {
+void scratchStaticKernel(float *src, float *dst) {
   size_t lid = syclext::this_work_item::get_nd_item<1>().get_local_linear_id();
-  float *scratch_mem = (float *)syclexp::get_work_group_scratch_memory();
-  syclexp::work_group_static<float[SIZE]> static_mem;
-  scratch_mem[lid] = src[lid];
-  static_mem[lid] = src[lid];
-  dst[lid] = scratch_mem[lid] + static_mem[lid];
+  float *scratchMem =
+      reinterpret_cast<float *>(syclexp::get_work_group_scratch_memory());
+  syclexp::work_group_static<float[SIZE]> staticMem;
+  scratchMem[lid] = src[lid];
+  staticMem[lid] = src[lid];
+  dst[lid] = scratchMem[lid] + staticMem[lid];
 }
 
 int main() {
@@ -66,41 +68,41 @@ int main() {
     src[i] = i;
   }
 
-  auto scratchbndl = syclexp::get_kernel_bundle<scratch_kernel,
-                                                sycl::bundle_state::executable>(
-      q.get_context());
-  auto staticbndl =
-      syclexp::get_kernel_bundle<static_kernel, sycl::bundle_state::executable>(
+  auto scratchBndl =
+      syclexp::get_kernel_bundle<scratchKernel, sycl::bundle_state::executable>(
+          q.get_context());
+  auto staticBndl =
+      syclexp::get_kernel_bundle<staticKernel, sycl::bundle_state::executable>(
           q.get_context());
-  auto scratchstaticbndl = syclexp::get_kernel_bundle<
-      scratch_static_kernel, sycl::bundle_state::executable>(q.get_context());
-
-  sycl::kernel ScratchKernel =
-      scratchbndl.template ext_oneapi_get_kernel<scratch_kernel>();
-  sycl::kernel StaticKernel =
-      staticbndl.template ext_oneapi_get_kernel<static_kernel>();
-  sycl::kernel ScratchStaticKernel =
-      scratchstaticbndl.template ext_oneapi_get_kernel<scratch_static_kernel>();
-  syclexp::launch_config ScratchKernelcfg{
+  auto scratchStaticBndl = syclexp::get_kernel_bundle<
+      scratchStaticKernel, sycl::bundle_state::executable>(q.get_context());
+
+  sycl::kernel scratchKrn =
+      scratchBndl.template ext_oneapi_get_kernel<scratchKernel>();
+  sycl::kernel staticKrn =
+      staticBndl.template ext_oneapi_get_kernel<staticKernel>();
+  sycl::kernel scratchStaticKrn =
+      scratchStaticBndl.template ext_oneapi_get_kernel<scratchStaticKernel>();
+  syclexp::launch_config scratchKernelcfg{
       ::sycl::nd_range<1>(::sycl::range<1>(SIZE), ::sycl::range<1>(SIZE)),
       syclexp::properties{
           syclexp::work_group_scratch_size(SIZE * sizeof(float))}};
-  syclexp::launch_config StaticKernelcfg{
+  syclexp::launch_config staticKernelcfg{
       ::sycl::nd_range<1>(::sycl::range<1>(SIZE), ::sycl::range<1>(SIZE))};
 
-  syclexp::nd_launch(q, ScratchKernelcfg, ScratchKernel, src, dst);
+  syclexp::nd_launch(q, scratchKernelcfg, scratchKrn, src, dst);
   q.wait();
   for (int i = 0; i < SIZE; i++) {
     assert(dst[i] == 2 * src[i]);
   }
 
-  syclexp::nd_launch(q, StaticKernelcfg, StaticKernel, src, dst);
+  syclexp::nd_launch(q, staticKernelcfg, staticKrn, src, dst);
   q.wait();
   for (int i = 0; i < SIZE; i++) {
     assert(dst[i] == src[i] * src[i]);
   }
 
-  syclexp::nd_launch(q, ScratchKernelcfg, ScratchStaticKernel, src, dst);
+  syclexp::nd_launch(q, scratchKernelcfg, scratchStaticKrn, src, dst);
   q.wait();
   for (int i = 0; i < SIZE; i++) {
     assert(dst[i] == 2 * src[i]);