Implement wrapper type for warp masks

[sbaldu]

This PR addresses issue #2615 from @fwyzard, implementing a wrapper type alpaka::warp::Mask around the type returned by activeMask and ballot, providing the same interface for all the backends.

[fwyzard]

can this be

Suggested change

	std::int32_t predicate) -> MaskType<WarpUniformCudaHipBuiltIn>::type
	std::int32_t predicate) -> Mask<WarpUniformCudaHipBuiltIn>

?

[fwyzard]

I am mildly concerned that this may introduce an ODR violation.
Do you think the warp type could be templated on TApi, and that could be used to distinguish the warp size ?

[fwyzard]

Can you make the mask type a nested type of the warp itself ?

Being able to write Warp::mask_t is simpler than alpaka::warp::Mask<Warp>.

[fwyzard]

Also, could you update the tests and example to make use of this type ?

[psychocoderHPC]

The mask type depends on the architecture you are targeting. If you target AMD it is 64bit because the warp/wave size is 64. The reason why it is 64 is that the OneApi plugin is using ROCm as connector to AMD hardware and in ROCm the mask size is devied as 64bit data type, IMO even if the warp size is 32.

[psychocoderHPC]

In alpaka3 we depend on the warp size but in alpaka mainline the warp size is not compile-time.
Maybe we can specialize the type on the accelerator type.

[fwyzard]

Yes, oneAPI uses uint64_t as the underlying type for the mask:

struct sub_group_mask {
  friend class sycl::detail::Builder;
  using BitsType = uint64_t;
...

As you suggest we could use different mask types for different oneAPI back-ends:

• 32 bit for NVIDIA GPUs
• 64 bit for AMD GPUs
• 32 bit for Intel GPUs (?)
• 64 bit for CPUs (though last time I checked a subgroup size of 64 was broken, but it's been a while)
• no idea for Intel/Altera FPGAs, probably 32 bit

or just 64 bits for all oneAPI back-ends if it's simpler and does not have any significant impact on the performance.

[psychocoderHPC]

IMO the version which uses different sizes per device type makes sense. Using always 64bit means one register overhead on systems with 32 as warp size.

I am sick and will not block the pr

[SimeonEhrig]

@sbaldu There is a typename missing in the code: https://gitlab.com/hzdr/crp/alpaka/-/jobs/13914947238