[SYCL][Reduction] Fix identityless span reductions in variadic (#8655)

Using a span reduction without an identity in a parallel_for with
multiple reductions currently fail due to a missed use of identity. This
commit addresses this missed case.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/include/sycl/reduction.hpp

@@ -2151,11 +2151,15 @@ void reduAuxCGFuncImplArrayHelper(bool IsOneWG, nd_item<Dims> NDIt, size_t LID,
     // Add the initial value of user's variable to the final result.
     if (LID == 0) {
       size_t GrID = NDIt.get_group_linear_id();
-      Out[GrID * NElements + E] =
-          IsOneWG ? BOp(LocalReds[0], IsInitializeToIdentity
-                                          ? IdentityContainer.getIdentity()
-                                          : Out[E])
-                  : LocalReds[0];
+      if constexpr (Reduction::has_identity) {
+        Out[GrID * NElements + E] =
+            IsOneWG ? BOp(LocalReds[0], IsInitializeToIdentity
+                                            ? IdentityContainer.getIdentity()
+                                            : Out[E])
+                    : LocalReds[0];
+      } else {
+        Out[GrID * NElements + E] = LocalReds[0];
+      }
     }
 
     // Ensure item 0 is finished with LocalReds before next iteration

sycl/test/regression/reduction_identityless_span.cpp

@@ -0,0 +1,25 @@
+// RUN: %clangxx -fsycl -fsyntax-only %s
+
+// Tests that identityless reductions compile when applied to a span.
+
+#include <sycl/sycl.hpp>
+
+template <class T> struct PlusWithoutIdentity {
+  T operator()(const T &A, const T &B) const { return A + B; }
+};
+
+int main() {
+  sycl::queue Q;
+
+  int *ScalarMem = sycl::malloc_shared<int>(1, Q);
+  int *SpanMem = sycl::malloc_shared<int>(8, Q);
+  auto ScalarRed = sycl::reduction(ScalarMem, PlusWithoutIdentity<int>{});
+  auto SpanRed = sycl::reduction(sycl::span<int, 8>{SpanMem, 8},
+                                 PlusWithoutIdentity<int>{});
+  Q.parallel_for(sycl::range<1>{1024}, ScalarRed, SpanRed,
+                 [=](sycl::item<1>, auto &, auto &) {});
+  Q.parallel_for(sycl::nd_range<1>{1024, 1024}, ScalarRed, SpanRed,
+                 [=](sycl::nd_item<1>, auto &, auto &) {});
+
+  return 0;
+}