[SYCL] Add support for SYCL 2020 in class group

[aobolensk]

Implement missing methods for class group according to SYCL 2020 4.9.1.7.

[Pennycook]

This change is good, but I think if we're going to add some SYCL 2020 functionality to the group class we should go all the way. If we only partially implement the SYCL 2020 interface, developers might be confused.

The SYCL 2020 group class is documented here: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#group-class, and I don't think many extra changes are required:

• Deprecate get_id()
• Add get_local_id()
• Add get_max_local_range()
• Add get_group_linear_id()
• Add get_local_linear_id()
• Add get_group_linear_range()
• Add get_local_linear_range()
• Add leader()

[aobolensk]

@Pennycook added functions you asked for in your comment. Also extended tests for group and added test for local_id on device (intel/llvm-test-suite#821).
Limitation: get_local_id() host implementation requires ABI breaking change, so it was stubbed with an exception. On device it works well.

[Pennycook]

I'd like @gmlueck to double-check this. Greg, can this be declared in our headers as a template function if it isn't a template function in the specification? My concern here is that this might be user-observable: a user could call get_local_linear_id<Dimensions + 1>, for example, or might have to insert extra template keywords in some situations.

If I'm right to be concerned, I think we can work around this by moving the specialization to a helper function, something like the following:

namespace sycl {
namespace detail {
  size_t get_local_linear_id(sycl::group<0>& g) const {
    return g.get_local_id();
  }

  // + Specializations for group<1> and group<2>

}
}

size_t get_local_linear_id() const {
  return sycl::detail::get_local_linear_id(*this);
}

[gmlueck]

Yes, it does seem weird that we would implement this function as a template when it is not that way in the spec. It seems like it would be easy to fix by calling some separate templated function as John proposes. I assume the separate function could also be a private member function of group?

class group {
 public:
  size_t get_local_linear_id() const {
    return get_local_linear_id_impl<Dimensions>();
  }

private:
  template <int dims>
  typename detail::enable_if_t<(dims == 1), size_t>
  get_local_linear_id_impl() const {
    id<Dimensions> localId = get_local_id();
    return localId[0];
  }
}

Or, you could use constexpr-if, but that would require the compiler to be C++17.

[aobolensk]

Done

[Pennycook]

I appreciate the specification could be clearer here, but the intent of this function is to return the maximum number of work-items in any work-group for the kernel that is currently executing, rather than the maximum number of work-items supported by the device.

Because we don't support non-uniform work-group sizes, I think you can replace all of this with:

size_t get_max_local_range() const {
  return get_local_range();
}

The reason both functions exist is that they return different values for the sub_group class. Supporting both in group allows developers to write generic code accepting either a sub_group or group.

[aobolensk]

Done

[Pennycook]

The specification for this function says:

Return true for exactly one work-item in the work-group, if the calling work-item is the leader of the work-group, and false for all other work-items in the work-group.

The leader of the work-group is determined during construction of the work-group, and is invariant for the lifetime of the work-group. The leader of the work-group is guaranteed to be the work-item with a local id of 0.

Because of the last sentence, I think the correct implementation is:

bool leader() const {
  return (get_local_linear_id() == 0);
}

[aobolensk]

Done

[bader]

@alexbatashev, take a look at this failure - https://github.com/intel/llvm/runs/5155471215?check_suite_focus=true. I think we should improve the CI scripts to avoid this.

[Pennycook]

Sorry for not being clear in my initial review -- my comment about get_local_linear_id() applies to all of the new template functions you've introduced here. I think there are three such functions left, and I've called out the first occurrence of each.

[Pennycook]

This should call a get_local_linear_range_impl().

[Pennycook]

This should call a get_group_linear_range_impl().

[Pennycook]

This should call a get_group_linear_id_impl().

[aobolensk]

Sorry for not being clear in my initial review -- my comment about get_local_linear_id() applies to all of the new template functions you've introduced here. I think there are three such functions left, and I've called out the first occurrence of each.

Done for rest of the functions

[againull]

@aobolensk Could you please check if this failure is related to your PR:

---

Failed Tests (1):
SYCL :: Basic/handler/handler_mem_op.cpp

[againull]

@aobolensk Could you please check if this failure is related to your PR:

Failed Tests (1): SYCL :: Basic/handler/handler_mem_op.cpp

Never mind, it is unrelated. Found a PR isolated to L0 plugin which has the same failure: #5541