[SYCL] Add structure to pass references to ranges

[Alexandr-Konovalov]

This allows to have single entry point for all dimentions and keep reference to user-provided range data, not copy them.

[uditagarwal97]

Just thinking out load, would it be better to call this class ranges_ref_view? Since this class is just a "view" over the user-provided range and doesn't own the lifetime of the underlying range object. Similar to how string_view or ranges::ref_view (https://www.en.cppreference.com/w/cpp/ranges/ref_view.html) works.

[Alexandr-Konovalov]

Sounds good, done.

[Alexandr-Konovalov]

Issue probably related to the test fail: Reduction/reduction_internal_nd_range_1dim.cpp sporadically fails in pre-commit BMG&L0 #19767

[vinser52]

I am not sure that it is the best name. We have sycl::range and someone might assume that ranges_ref_view is a view to the sycl::range. Also it is not lear what _ref_ means in the name. Why not just nd_range_view?

[Alexandr-Konovalov]

View has strong connotation to views from C++20, that may be or may be not good.

Alternatively, it can be nd_range_ref. I don't know.

Also it is not lear what ref means in the name.

That's for "reference", I hope.

[vinser52]

I vote for nd_range_view.

@sergey-semenov, @aelovikov-intel, @slawekptak what do you think?

[aelovikov-intel]

I'll copy my sketch from the offline group chat:

struct range_storage { std::array<size_t, 3> r; };

struct range_view {
  range_storage &s;
  size_t Dims;

  range_view(const range_base&);
  range_view(const range_storage&, size_t);

  operator range_base();

  // define all range interfaces;
};

struct range_base : range_storage { size_t Dims; 
// All interfaces delegate to range_view(*this).interface(...);
}
template <size_t> range : range_base {};

struct nd_range_base {
  size_t Dims;
  range_storage global;
  ...
};

[slawekptak]

nd_range_view sounds good to me. Maybe we should change GlobalOffset to Offset, so the naming fully reflects the nd_range naming?

[Alexandr-Konovalov]

I'll copy my sketch from the offline group chat:

@aelovikov-intel , Could you please clarify what is the problem are you solving? All current implementation can do is to consume sycl::ranges or nd_range and to produce sycl::detail::NDRDescT. Probably, you are thinking about wider context.

[Alexandr-Konovalov]

ok, renamed to nd_range_view.

[aelovikov-intel]

More code unification, less than 4 pointer-sized data members for the "view", ability to pass existing sycl::[nd_]range<N> as const sycl::detail::[nd_]range_base & directly to all the helper function unifying all Dims handling.

Do you have subsequent refactoring changes available? POC/draft quality would be fine. Are you going to switch all ABI entry points to accept the view?

[uditagarwal97]

The testing logic for 1,2 and 3 dimensions ranges_ref_view is the same. Would it be possible to abstract out the testing logic in a function, templated over the dimension parameter?
For example:

template <int dims>
void TestNDRangesRefView(sycl::ranges<dims> global, sycl::ranges<dims> local, sycl::id<dims> offset) {
sycl::nd_range<dims> nd_range{global, local, offset};
  {
    sycl::detail::ranges_ref_view r{nd_range};
     ASSERT_EQ(r.Dims, size_t{dims}); // I guess this can be a compile-time static assert?
     for (int d = 0, d < dims, d++) {
        ASSERT_EQ(r.GlobalSize[d], global_range[d]);
        ASSERT_EQ(r.LocalSize[d], local_range[d]);
        ASSERT_EQ(r.GlobalOffset[d], offset[d]);
     }
  }
 ...
}

And then this function can be called thrice, once for each dimension, 1, 2, 3. That should reduce code duplication and would be easier to update it, if something changes in ranges_ref_view

[Alexandr-Konovalov]

Thanks, that's much better!

Unfortunately I can't use static_assert for the dimension check, constness should be improved.

[slawekptak]

NDRDescT has some additional logic in the constructors, which probably has to be included here. Would it make sense to use the NDRDescT constructors, depending on which fields of nd_range_view are set? Also, there is SetNumWorkGroups argument - it is set to false, when the NDRDescT(range) constructor is set, but in some cases it seems that it is set to true. Maybe it needs to be an argument of toNDRDescT function.

[Alexandr-Konovalov]

Are you talking about something like below, right?

  if (!LocalSize && !Offset)
    if (Dims == 1)
      return NDRDescT<1>(sycl::range<1>(GlobalSize[0]));
    else if (Dims == 2)
      return NDRDescT<2>(sycl::range<2>(GlobalSize[0], GlobalSize[1]));
    else if (Dims == 3)
      return NDRDescT<3>(sycl::range<3>(GlobalSize[0], GlobalSize[1], GlobalSize[2]));

sycl::range here owns values, so to call right constructor we must create a copy. That's what I trying not to do.

Probably smart compiler able to inline everything and eliminate creation of the intermittent copies. I don't know.

Wrt SetNumWorkGroups. According to my understanding, they may appear via parallel_for_work_group. But in sycl_khr_free_function_commands extension as well as between queue shortcut functions I don't see something like that. In what contexts can we expect to meet them?

[slawekptak]

I think you are right about SetNumWorkGroups. Maybe if we'd like to use nd_range_view in all of the paths, we would need it, but it seems like for the khr extension and queue shortcuts we don't need it.

I still feel, that NDRDescT should be allocated using a constructor, so all the required padding is applied. What do you think about adding a new NDRDescT constructor, which takes a pointer and dims, and can initialize the internal fields and apply the same logic as other constructors? Maybe the padding logic can be extracted to separate methods and called in both constructors.

[Alexandr-Konovalov]

What do you think about adding a new NDRDescT constructor, which takes a pointer and dims, and can initialize the internal fields and apply the same logic as other constructors?

Sure, this functionality is naturally belongs to NDRDescT. Will do.