Hierarchical: LowerWG pass ignores convergent functions

[Naghasan]

When I compile this code:

#include <CL/sycl.hpp>

using namespace cl::sycl;

int main() {
  queue q(default_selector().select_device());
  auto buf = cl::sycl::buffer<int, 1>(cl::sycl::range<1>(1));
  q.submit([&](handler &cgh) {
    auto globalAcc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
    auto sizeAcc = buf.get_access<cl::sycl::access::mode::read>(cgh);
    auto localAcc =
        cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
                            cl::sycl::access::target::local>(
            cl::sycl::range<1>(1), cgh);
    cgh.parallel_for_work_group<class example_kernel>(
        range<1>(1), range<1>(1),
        [=](group<1> g) {
            cl::sycl::device_event device_event_1 =
                g.async_work_group_copy(globalAcc.get_pointer(),
                                        localAcc.get_pointer(), 1);
        });
  });
  return 0;
}

The LowerWG pass wraps the call to async_work_group_copy so it gets only executed by the workgroup leader. This disregards the call to __spirv_GroupAsyncCopy made by async_work_group_copy which a convergent function, hence must be reachable by all work items.

[KornevNikita]

Hi! There have been no updates for at least the last 60 days, though the ticket has assignee(s).

@againull, could I ask you to take one of the following actions? :)

• Please provide an update if you have any (or just a small comment if you don't have any yet).
• OR mark this issue with the 'confirmed' label if you have confirmed the problem/request and our team should work on it.
• OR close the issue if it has been resolved.
• OR take any other suitable action.

Thanks!

[KornevNikita]

Though the issue was created 4+ years ago, but looks like it's still an issue. I updated the source as some interfaces were deprecated:

#include <sycl/sycl.hpp>

using namespace sycl;

int main() {
  queue q;
  auto buf = sycl::buffer<int, 1>(sycl::range<1>(1));
  q.submit([&](handler &cgh) {
    auto globalAcc = buf.get_access<sycl::access::mode::read_write>(cgh);
    auto sizeAcc = buf.get_access<sycl::access::mode::read>(cgh);
    auto localAcc = sycl::local_accessor<int>(sycl::range<1>(1), cgh);
    cgh.parallel_for_work_group<class example_kernel>(
        range<1>(1), range<1>(32),
        [=](group<1> g) {
            sycl::device_event device_event_1 =
                g.async_work_group_copy(globalAcc.get_multi_ptr<access::decorated::yes>(),
                                        localAcc.get_multi_ptr<access::decorated::yes>(), 1);
        });
  });
  return 0;
}

And @_Z22__spirv_GroupAsyncCopy is getting called under wg_leader.i entry: main-sycl-spir64-unknown-unknown.ll.txt

[againull]

SYCL 2020 recommend that new codes refrain from using this feature, so it is not planned to fix this bug.