Improve test and add comments

Author: lbushi25

clang/lib/Sema/SemaSYCL.cpp

@@ -6284,7 +6284,13 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "#include <sycl/detail/defines_elementary.hpp>\n";
   O << "#include <sycl/detail/kernel_desc.hpp>\n";
   O << "#include <sycl/ext/oneapi/experimental/free_function_traits.hpp>\n";
-  O << "#include <sycl/ext/oneapi/properties/properties.hpp>\n";
+  // When using work group memory parameters in free kernel functions, the
+  // integration header emits incorrect forward declarations for the work group
+  // memory type because it ignores default arguments. This means the user
+  // cannot use work group memory types with parameters omitted such as
+  // work_group_memory<int> where the hidden second parameter has a default
+  // value. To circumvent this, we include the correct forward declaration
+  // ourselves.
   O << "#include <sycl/ext/oneapi/experimental/work_group_memory_forward_decl.hpp>\n";
   O << "\n";
 

sycl/include/sycl/handler.hpp

@@ -7,7 +7,6 @@
 //===----------------------------------------------------------------------===//
 
 #pragma once
-#include <iostream>
 #include <sycl/access/access.hpp>
 #include <sycl/accessor.hpp>
 #include <sycl/context.hpp>

sycl/test-e2e/WorkGroupMemory/free_function_kernel.cpp

@@ -14,13 +14,13 @@ using namespace sycl;
 queue q;
 context ctx = q.get_context();
 
-int sum_helper(sycl::ext::oneapi::experimental::work_group_memory<int[]> mem,
-               size_t WGSIZE) {
-  int ret = 0;
+void sum_helper(sycl::ext::oneapi::experimental::work_group_memory<int[]> mem,
+                sycl::ext::oneapi::experimental::work_group_memory<int> ret,
+                size_t WGSIZE) {
+  ret = 0;
   for (int i = 0; i < WGSIZE; ++i) {
     ret += mem[i];
   }
-  return ret;
 }
 
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
@@ -38,7 +38,9 @@ void sum(sycl::ext::oneapi::experimental::work_group_memory<int[]> mem,
         *Result += mem[i];
       }
     } else {
-      *Result = sum_helper(mem, WGSIZE);
+      sycl::ext::oneapi::experimental::work_group_memory<int> ret;
+      sum_helper(mem, ret, WGSIZE);
+      *Result = ret;
     }
   }
 }