DPCPP cooperative group

[yhmtsai]

This PR adds the cooperative group in dpcpp to keep the same interface as cuda/hip.
Note. the subgroup seems to be warp in cuda not subwarp.
Reference: https://intel.github.io/llvm-docs/cuda/opencl-subgroup-vs-cuda-crosslane-op.html

This PR is WIP but I would like to bring it to review.
to check the config selection can be acceptable from ginkgo
for config selection, please check the cg_shuffle_config_call in dpcpp/test/components/cooperative_groups_kernels.dp.cpp

Summary:

• dim3 and sycl_nd_range: a cuda-like usage for sycl range and nd_range with tests
• helper gives default implementation macro for simple kernel cases (no explicit template parameter and 1d block)
• __WG_BOUND__ gives something like __launch_bound__ but it needs the 3d information not the product
  __WG_BOUND_CONFIG__ can use ConfigSet for easy unpack
• cooperative group implementation and set the test result individually
• another selection for config (it allows bool, int, size_type template by roughly go through all kernel template)
• update format_header such that it can handle the generated dpcpp file (the script is not yet here)
• ConfigSet: @tcojean provides a general way to embed information into encoded integer

TODO:
also set the cooperative group test result individually in cuda/hip?
I decide to move another PR to handle the cuda/hip cooperative group because it contains another subwarp test

[tcojean]

Thanks Mike for creating the PR. I'm trying to see what I can do with a templated configuration type as possible improvement, I'll tell you whether it works.

[upsj]

I didn't go into the config right now, but the rest looks good

[codecov]

Codecov Report

Merging #757 (7f72418) into develop (95c7652) will increase coverage by 0.00%.
The diff coverage is 100.00%.

@@           Coverage Diff            @@
##           develop     #757   +/-   ##
========================================
  Coverage    94.17%   94.17%           
========================================
  Files          400      400           
  Lines        31051    31080   +29     
========================================
+ Hits         29241    29270   +29     
  Misses        1810     1810           

Impacted Files	Coverage Δ
include/ginkgo/core/base/types.hpp	92.59% <ø> (ø)
core/test/base/types.cpp	100.00% <100.00%> (ø)

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 95c7652...7f72418. Read the comment docs.

[yhmtsai]

test_all_subgroup will run 4, 8, 16 for cpu and 8, 16, 32 for intel gpu
search dpcpp/test/components/cooperative_groups_kernels on dgpu-job, or cpu-job

[Slaedr]

Overall, this looks good!

The main point of interest for me is the choice of 32 for the warp size. Could you document somewhere, maybe just in this PR, why you chose 32? Will shuffles be expected to work for 32 work-items independent of vector architecture? Can we get away with a smaller warp size? (I think a smaller warp size, if it still fully utilizes the vector units, can enable more flexible parallelism.)

Apart from that I have a few other points. I guess the sync implementation and tests still remain, and documentation can be improved in several places.

[Slaedr]

Do we need this-> here? If not, it's probably better to remove it. Maybe

Suggested change

	const auto result = this->shuffle_up(var, selector);
	const auto result = sub_group::shuffle_up(var, selector);

[yhmtsai]

I think it needs this-> to use the generated subgroup information

[Slaedr]

Maybe just ShflOp is enough here?

Suggested change

	return this->ShflOp(var, selector); \
	return ShflOp(var, selector); \

[yhmtsai]

When I call the member function, I will use this->

[upsj]

member functions are sometimes not found in a templated member function unless you use this->, which is also why all of our templated tests need to use it.

[Slaedr]

Typically, that is not needed for inline definitions (in the class body rather than outside). I was just thinking of avoiding an explicit pointer indirection if we can help it. I understand it's probably fine because it should be optimized out.

[Slaedr]

Can you give an example of a function this will enable? In general it seems like a bad idea for us to depend on sycl's detail namespace - I guess they can change it any time.

[yhmtsai]

for example, we use the group algorithm the reduce in ballot.
If we do not do this way, when we need use oneAPI implemenation on gko's subgroup, we need to add static_cast to subgroup type.
I agree with you, they might change it at some point.

[Slaedr]

Yes, if it works with static casting to sycl::ONEAPI::sub_group* when you call reduce, perhaps we should just do that.

[yhmtsai]

I stay this version first because it still works and it makes ginkgo cooperative group like an extension not another object.
But if they change frequently and we are hard to stay with them, we can delete them and use static_cast.

[Slaedr]

Why is this specialization needed? Do you expect we'll need to use thread_block_tiles of size 1 in algorithms? Otherwise, if this case is correctly but inefficiently handled by the generic implementation, I guess we don't need this specialization. If it's needed, it would be nice to have unit tests for this too.

[yhmtsai]

it is for those kernels implemented by warp sense.
providing it such that we can use the same kernel on subgroup(1) to get the single thread implementation

[yhmtsai]

I add a few default helper to reduce the internal layer code.
GKO_ENABLE_DEFAULT_HOST_CONFIG(name, kernel) generates the default implementation for those kernels without expilict template parameter.
GKO_ENABLE_DEFAULT_CONFIG_CALL(name, callable, cfg, list)
cfg gives how to decode the int and list gives the available list for selection.
Also, @tcojean provides a general way, ConfigSet, to embed the several information into one int.
It can use not only in dpcpp but also in cuda/hip thing like ellpack we have 2 (3) parameters for selection.

TODO:

• add documentation

[yhmtsai]

@ginkgo-project/reviewers this PR can be reviewed. Also welcome any name suggestion of these things

[tcojean]

LGTM

[tcojean]

Will we ever need this? I think we cannot have masked operation anyway. Maybe all of the config can be removed to begin with?

[upsj]

I also think the full lane mask should have the same number of bits as a full lane has threads. Is that true here? Otherwise things related to popcnt and ballot might break.

[yhmtsai]

you need to use mask<subgroup_size, lane_mask_type> to get the last subgroup_size bits activated in lane_mask_type

[tcojean]

You can use validate_function instead?

Also, why are these protected? How do you use them?

[yhmtsai]

the protected structure is the original one.
I add a public function to get the const exec_info.
Do you have any comment about using function directly or structure function?

[upsj]

Looks mostly good to me, I would like to see if we can remove a few CUDA-isms though, since they only relate to kernel launch, not the actual kernel implementation.

[upsj]

I also think the full lane mask should have the same number of bits as a full lane has threads. Is that true here? Otherwise things related to popcnt and ballot might break.

[upsj]

Are we programming for DPC++ or SYCL? For the former, can't we include dpct.hpp directly? For the latter, is this portable?

[yhmtsai]

this is due to the issues from dpct before. for example, including dpct file gives error or their atomic_add has some issue, so we also have atomic_add implementation for real number in another pr

[upsj]

We are building a lot of complexity here (also the config selection dim3 integration), so I want to ask: Why do we need this additional wrapper? Can't we use SYCL primitives directly?

[yhmtsai]

Yes, we can use the SYCL native sycl::range<3>
but it will gives different view on kernel launch and build the sycl::range
for example, we use kernel<<<32, 32>>> and sycl will need to use
kernel(sycl::nd_range<3>(sycl::range<3>(1, 1, 32) * sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)))
but with dim3, we can still use kernel(32, 32)
and in the beginning, I would like to reduce the difference from cuda to dpcpp

[upsj]

In which places does that actually matter? Most of the time, we are using one-dimensional kernels, except for SpMV (multiple columns), some Dense kernels (2D) and Jacobi. Wouldn't you use sycl::ndrange<1>(...) directly?

[yhmtsai]

I think the mapping will be different. At least, it will require us to develop own dpct to avoid any conversion like threadIdx -> get_local_id(3)

[yhmtsai]

and it will gives inconsistent index sense between the kernel1D/2D/3D.
range<1>(x) get_local_id(0) is x and x is contiguous as cuda
range<2>(x, y) get_local_id(0) is still x but x is not contiguous
or range<2>(y, x) x is contiguous but x needs get_local_id(1)

[upsj]

But that only matters on the launch side, right? I would assume that the reason for this is that it maps more cleanly to the nested for-loop model

for (; i < range[0]; i++)
    for (; j < range[1]; j++)
        for (; k < range[2]; k++)

[yhmtsai]

it affects not only launch side but the kernel index for threads

[upsj]

member functions are sometimes not found in a templated member function unless you use this->, which is also why all of our templated tests need to use it.

[Slaedr]

Excellent work! It's pretty cool, eg. how you handle stuff like __WG_BOUNDS__. Guess you want to make DPC++ as close to CUDA as possible, which is what a lot of the code here seems to be doing.

Below, I have some concerns around the ConfigSet stuff and other comments / suggestions.

[Slaedr]

Suggested change

	template <int num_groups, int current_shift>
	template <int current_shift, int num_groups>

This way, I think num_groups can be inferred and we'll need to provide only current_shift while using this.

[yhmtsai]

I tried it before but it did not get the num_group information from array

[Slaedr]

Oh I see. Even with the switched order? Then never mind.

[yhmtsai]

it needs switched order to give the possibility for not explicitly setting.
I will put the my trying code later.
It is also not my expectation, so maybe I did something wrong there.

[upsj]

Did you try adding a deduction guide? Neat little C++17 feature that might help here.

[yhmtsai]

I think ConfigSet is also a general solution for others kernel (not dpcpp) so we still need to stay C++14.
I put the related code: https://godbolt.org/z/oa5d4arMs
it switch the order but can not miss the num_groups

[Slaedr]

Suggested change

	template <int num_groups, int current_shift>
	template <int current_shift, int num_groups>

[Slaedr]

Then this can simply be

Suggested change

	shift<num_groups, (current_shift + 1)>(bits);
	shift<current_shift + 1>(bits);

[Slaedr]

Suggested change

	* certain bits. The usage will be the following
	* certain bits.
	*
	* The usage will be the following:

[Slaedr]

You want to put ints into a char array?

[Slaedr]

Somewhat crazy suggestion but might be great to have: maybe you could include a mathematical proof that this coding scheme is indeed a unique map when the numbers to encode are small enough. You could precisely state (in the assumptions) the maximum values of the arguments of encode for the coding to work, and then the proof would be couple of lines to prove:
For integer vectors x and y with all components less than the respective maxima, x != y implies encode(x...) != encode(y...).

If you say that is unnecessary, that is completely fine.

[Slaedr]

Yes, if it works with static casting to sycl::ONEAPI::sub_group* when you call reduce, perhaps we should just do that.

[Slaedr]

I guess 12 and 7 are arbitrarily chosen, or do you like 12 and 7 for some specific reason?

[yhmtsai]

I use wrong number here. It should be 11, 7.
it is from log_2(1024) + 1 for workgoup_size and log_2(64) + 1 for subgroup_size

[Slaedr]

That appears to make sense, but I don't really see it. It seems like you want the bits array to contain the maximum number of bits needed by each position. But I guess it will work correctly for many combinations, depending on which numbers are given as parameters to the encode function.

[yhmtsai]

it is described in ConfigSet note. the #bit should be log_2(max)+1

[upsj]

LGTM - nice job! I am still not entirely happy about the dim3 thing, since it means adding CUDA-isms to our SYCL code, but I guess it would only be used in a small number of places anyways, so I guess it will be okay.

[upsj]

Suggested change

	ASSERT_EQ((std::is_same<decltype(mask3_u), const unsigned int>::value),
	true);
	ASSERT_EQ((std::is_same<decltype(fullmask_u), const unsigned int>::value),
	true);
	ASSERT_TRUE((std::is_same<decltype(mask3_u), const unsigned int>::value));
	ASSERT_TRUE((std::is_same<decltype(fullmask_u), const unsigned int>::value));

Not sure if you even need the additional parentheses then?

[yhmtsai]

It still needs the additional parentheses such that the macro unpacks parameter correctly.

[upsj]

same here and the following

[upsj]

What is the reason for these changes again? Do we no longer need it?

[yhmtsai]

the behavior still exists. DPCPP gives more additional header than <hip/runtime.h>, so I move the all additional before the LICENSE into header section not just hip/runtime.h

[upsj]

Did you try adding a deduction guide? Neat little C++17 feature that might help here.

[yhmtsai]

it depends on how much effort we would like to start DPCPP in the beginning.
If we wouldn't like Cuda-isms into dpcpp, we should reimplement all not porting in the begining.
Using dpct gives the index structure and also the global range is not possible be the same as grid.
without dim3, we still need to do the same thing manually and need to switch our brand between SYCL/(HIPCUDA)

[yhmtsai]

I swap the order of shift although we do not have auto deduction there.
But if someway to achieve it or use C++17, this order of template is unchanged.
I add some comments to show the encode/decode is bijection function.
Also the ConfigSet can not be compiled or can throw the error.

• the total bits is larger than the bits of ConfigSetType
• the value is larger than the certain bit representation (delete due to gcc 5.x does not support conditional throw in constexpr)
• encode uses wrong number of args
• decode out-of-range position

[sonarcloud]

Kudos, SonarCloud Quality Gate passed!

0 Bugs
0 Vulnerabilities
0 Security Hotspots
1 Code Smell

No Coverage information
3.4% Duplication