[SYCL][Doc] Extended group load/store APIs proposal

[aelizaro]

An initial draft of extended group load/store APIs to provide capabilities to work with temporary memory buffers and load/store multiple elements per work item.

[aelizaro]

@andreyfe1, @Pennycook, could you please take a look at the initial proposal?

[Pennycook]

Please also wrap things at 80 columns; it makes it much easier to review.

[Pennycook]

Instead of using extra options in the GroupMemoryHelper, could we make these two new properties? Then a developer could use them with the overload that doesn't require a GroupMemoryHelper.

[aelizaro]

OK, removed (2) and (4) and added a TODO for other properties, wanna do some prototyping for it

[dkhaldi]

It will be good to have complete kernels for this extension in tests directory.
Specifically, I would like to see the usage of vec, marray, and simd as the span memory storage.

[aelizaro]

@dkhaldi, do you mean to add full tests without implementation, to illustrate how it supposes to work with sycl's data types?

[dkhaldi]

@dkhaldi, do you mean to add full tests without implementation, to illustrate how it supposes to work with sycl's data types?

Right.

[aelizaro]

Test cases are added to illustrate how APIs should work.
For sycl::vec case, we can use both span approaches (as it is shown in the test) or have a specialization for a single value case - it will look cleaner.
For sycl::marray it works nicely as its internal representation guarantees contiguous storage of memory, so we can use:

sycl::marray<InputT, items_per_thread> data;
sycl_exp::joint_load(item.get_group(), in.get_pointer(), sycl::span<InputT, items_per_thread>{ data.begin(), data.end() });

[Pennycook]

Spotted a few minor issues during review of the latest changes, but I think the fixes are straightforward. Let me know when they're applied and I'll approve.

[Pennycook]

I think we should combine these into one hint. We can call it full_group for now.

We already know whether we're using a group or a sub-group from the first argument. There are also going to be other group types in the future, and it would be convenient to be able to re-use the same property across all of those group types as well.

[aelovikov-intel]

I'm not sure full_group is a good name. parallel_for requires that global size is a multiple of a work group size, so WG is always "full" in some sense. It might still be that the last SG of each WG isn't full.

[Pennycook]

Right, but that just means that full_group doesn't do anything when the algorithm is invoked with a sycl::group, but improves performance when the algorithm is invoked with a sycl::sub_group. The "group" in "full group" is intended to mean "anything satisfying the group concept" and not "sycl::group".

I want to push for a renaming of sycl::group to sycl::work_group to avoid this sort of confusion.

[aelovikov-intel]

Then we need another property for the WG case, I think.

[Pennycook]

Why? Maybe I'm misunderstanding what you're expecting to happen here. I thought your earlier message was saying that work-groups were always "full"?

[aelovikov-intel]

They are, but the block-read is SG-level intrinsic. We need SGs to be "full" to use them.

[Pennycook]

I think we're talking past each other.

Consider the below:

auto sg = it.get_sub_group();
group_load(sg, input, output_span, properties{contiguous_memory, full_group});

The interpretation here is that sg represents a full sub-group.

[aelovikov-intel]

If the user writes

  auto g = it.get_group();
  group_load(g, input, output_span, properties{contiguous_memory, full_group});

we still can't use block loads without runtime checks. WG is full, but its last SG might not be. Also, full_group is meaningless for WG because it always full per SYCL spec parallel_for/nd_range restrictions.

[Pennycook]

We couldn't use block loads in this case anyway, because the stride in this case is the size of the work-group and not the sub-group. You're right that we could generate block loads in some cases (like a work-group + blocked data placement) but I think we'd still need a runtime check because different devices carve up work-groups into sub-groups differently.

full_group is currently meaningless for work-groups, but I think it makes more sense to have a generic full_group than to make the property sub-group specific. That allows a developer to write generic code that uses full_group and have it work in both cases, and it may be meaningful for future group types.

[Pennycook]

Users might not understand why this is helpful for implementations, so I think we should be a little more explicit here. Maybe something like:

Suggested change

	The following properties is introduced to be used
	as a hint that implementation can use get_max_local_range():
	The following property can be used as a hint that
	`get_local_range()` is equal to `get_max_local_range()`,
	which may enable more aggressive optimizations for some
	implementations.
	[NOTE]
	====
	Using `full_group` is necessary to generate SPIR-V block read
	and block write instructions, because these instructions are
	defined to use the maximum group size as the stride.
	====

[aelizaro]

@gmlueck, @Pennycook, @aelovikov-intel, could you please, double-check the PR? (and huge thanks to you for contribution!)

I think I addressed everything except one ongoing discussion about full_group hint naming #7593 (comment)

[Pennycook]

I spotted one minor potential issue, but apart from this and the full_group discussion, I think we're good.

[Pennycook]

Same question as above. Can we use contiguous here?

[gmlueck]

I think this looks good, but I have some comments about the use of the word "hint". See below.

[gmlueck]

LGTM!

[aelizaro]

@Pennycook, can we merge it now?

[Pennycook]

@Pennycook, can we merge it now?

I think so! @intel/llvm-gatekeepers, please merge.