Simple reduction example with buffers fails with an assert when AdaptiveCpp is compiled in debug mode.

[nilsfriess]

To Reproduce
Build AdaptiveCpp from current develop in debug mode (i.e. pass -DCMAKE_BUILD_TYPE=Debug) and then compile the following example:

#include <sycl/sycl.hpp>
#include <numeric>
#include <cassert>

int main() {
  sycl::buffer<int> vals{1024};

  {
    sycl::host_accessor a{vals};
    std::iota(a.begin(), a.end(), 0);
  }

  int sum_res = 0;
  sycl::buffer<int> sum_buf{&sum_res, 1};

  sycl::queue{}.submit([&](auto &cgh) {
    auto input = vals.get_access<sycl::access_mode::read>(cgh);

    auto sum_red = sycl::reduction(sum_buf, cgh, sycl::plus<int>());

    cgh.parallel_for(sycl::range<1>{1024}, sum_red, [=](sycl::id<1> idx, auto &sum) {
      sum += input[idx];
    });
  });

  assert(sum_buf.get_host_access()[0] == 523776);
}

This compiles fine but running the example fails with the message:

/home/nfriess/AdaptiveCpp/src/runtime/multi_queue_executor.cpp:68: std::size_t hipsycl::rt::{anonymous}::determine_target_lane(hipsycl::rt::dag_node_ptr, const hipsycl::rt::node_list_t&, const hipsycl::rt::multi_queue_executor*, const hipsycl::rt::moving_statistics&, hipsycl::rt::backend_execution_lane_range): Assertion `req->is_submitted()' failed.

Happens both with --acpp-targets=cuda:sm_61 and generic. Here's the full output with ACPP_DEBUG_LEVEL=3: out.log

When AdaptiveCpp is compiled in release mode, the example runs fine and computes the correct result.

[FattiMei]

I have compiled your program with a default adaptiveCpp config and it gives the runtime error:

[AdaptiveCpp Warning] dag_direct_scheduler: Detected a requirement that is neither of discard access mode (SYCL 1.2.1) nor no_init property (SYCL 2020) that accesses uninitialized data. Consider changing to discard/no_init. Optimizing potential data transfers away.

And then segfaults, not sure if it contributes to the discussion

[illuhad]

[AdaptiveCpp Warning] dag_direct_scheduler: Detected a requirement that is neither of discard access mode (SYCL 1.2.1) nor no_init property (SYCL 2020) that accesses uninitialized data. Consider changing to discard/no_init. Optimizing potential data transfers away.

This is not an error, just a warning that points out that you could have created the accessor with the no_init property as an optimization. If you get a segfault, it is not related to this warning. Without stacktrace it is unclear whether your segfault is related to this issue.

[FattiMei]

Are there adaptive tools to produce such stacktrace?

acpp --acpp-targets="cuda:sm_60" works but acpp --acpp-targets="cuda:sm_80" doesn't

[illuhad]

@FattiMei Just use gdb?

[FattiMei]

Thank you for the suggestions, unfortunately I can't reproduce the segfault anymore, but I solved my problems

[illuhad]

It seems that the problem described by @nilsfriess is limited to the buffer-accessor model. With USM I don't see the issue.