Include initial value of reduction variable in reduction

[jolly-chen]

Hi all,

I have a situation where I need to do a reduction in multiple phases, so the result of the current reduction needs to be combined with the result of the previous reduction. Under section 4.9.2.1. reduction interface, the standard states that

The initial value of the reduction variable is included in the reduction operation, unless the property::reduction::initialize_to_identity property was specified when the reduction interface was invoked.

Unless I am misinterpreting this sentence, I believe that I can achieve the above by simply defining a reduction again with the same output location.

However, when I run the code example below with AdaptiveCpp the original value seems to be overwritten and I only get 512. When I compile the code snippet below with icpx I get the desired result 1024. Is this not implemented in AdaptiveCpp yet or am I missing something?

Thanks!
-- Jolly

#include <algorithm>
#include <sycl/sycl.hpp>

int main(int argc, char const *argv[]) {
  size_t n = 512;
  sycl::queue q;
  sycl::buffer<int> values_buff{n};
  {
    sycl::host_accessor a{values_buff};
    std::fill(a.begin(), a.end(), 1);
  }

  int *sum_ptr = sycl::malloc_device<int>(1, q);
  q.memset(sum_ptr, 0, sizeof(int));

  q.submit([&](sycl::handler &cgh) {
     auto values_acc = values_buff.get_access<sycl::access_mode::read>(cgh);
     auto sumReduction = sycl::reduction(sum_ptr, sycl::plus<int>());

     cgh.parallel_for(
         sycl::range<1>{n}, sumReduction,
         [=](sycl::id<1> idx, auto &sum) { sum += values_acc[idx]; });
   }).wait();
  q.submit([&](sycl::handler &cgh) {
     auto values_acc = values_buff.get_access<sycl::access_mode::read>(cgh);
     auto sumReduction = sycl::reduction(sum_ptr, sycl::plus<int>());

     cgh.parallel_for(
         sycl::range<1>{n}, sumReduction,
         [=](sycl::id<1> idx, auto &sum) { sum += values_acc[idx]; });
   }).wait();

  int result;
  q.memcpy(&result, sum_ptr, sizeof(int)).wait();
  printf("Result: %d\n", result);
}

[illuhad]

Your are correct with everything you say. Our reduction implementation is not yet complete. You could work around the issue by having e.g. work item 0 fetch the result of the previous reduction, and add it to the current reduction.