Remove restrictions on sycl::group_broadcast

[masterleinad]

Is your feature request related to a problem? Please describe
Related to kokkos/kokkos#4388.
Despite SYCL2020 not constraining sycl::group_broadcast (https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#_group_broadcast), the oneAPI implementation does, see

llvm/sycl/include/CL/sycl/group_algorithm.hpp

Lines 452 to 499 in f844f70

	// ---- group_broadcast
	template <typename Group, typename T>
	detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
	detail::is_scalar_arithmetic<T>::value),
	T>
	group_broadcast(Group, T x, typename Group::id_type local_id) {
	#ifdef __SYCL_DEVICE_ONLY__
	return sycl::detail::spirv::GroupBroadcast<Group>(x, local_id);
	#else
	(void)x;
	(void)local_id;
	throw runtime_error("Group algorithms are not supported on host device.",
	PI_INVALID_DEVICE);
	#endif
	}
	template <typename Group, typename T>
	detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
	detail::is_scalar_arithmetic<T>::value),
	T>
	group_broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
	#ifdef __SYCL_DEVICE_ONLY__
	return group_broadcast(
	g, x,
	sycl::detail::linear_id_to_id(g.get_local_range(), linear_local_id));
	#else
	(void)g;
	(void)x;
	(void)linear_local_id;
	throw runtime_error("Group algorithms are not supported on host device.",
	PI_INVALID_DEVICE);
	#endif
	}
	template <typename Group, typename T>
	detail::enable_if_t<(is_group_v<std::decay_t<Group>> &&
	detail::is_scalar_arithmetic<T>::value),
	T>
	group_broadcast(Group g, T x) {
	#ifdef __SYCL_DEVICE_ONLY__
	return group_broadcast(g, x, 0);
	#else
	(void)g;
	(void)x;
	throw runtime_error("Group algorithms are not supported on host device.",
	PI_INVALID_DEVICE);
	#endif
	}

.

Describe the solution you would like
Remove the restriction to arithmetic types.

Describe alternatives you have considered
There is a workaround in kokkos/kokkos#4388 that should also work as a general fallback option if a more efficient implementation for non-arithmetic types cannot be provided.

[Pennycook]

DPC++ used to support group_broadcast for trivially copyable types, and the deprecated extension still does:

llvm/sycl/include/sycl/ext/oneapi/group_algorithm.hpp

Lines 145 to 160 in f844f70

	template <typename Group, typename T>
	__SYCL2020_DEPRECATED(
	"ext::oneapi::broadcast is deprecated. Use group_broadcast instead.")
	detail::enable_if_t<(detail::is_generic_group<Group>::value &&
	std::is_trivially_copyable<T>::value &&
	!detail::is_vector_arithmetic<T>::value),
	T> broadcast(Group, T x, typename Group::id_type local_id) {
	#ifdef __SYCL_DEVICE_ONLY__
	return sycl::detail::spirv::GroupBroadcast<Group>(x, local_id);
	#else
	(void)x;
	(void)local_id;
	throw runtime_error("Group algorithms are not supported on host device.",
	PI_INVALID_DEVICE);
	#endif
	}

SYCL 2020 used to say very little about the types supported by group_broadcast, but the section introducing the group functions (https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:group-functions) says:

All group functions are supported for the fundamental scalar types supported by SYCL (see Table 179) and instances of the SYCL vec and marray classes.

We decided to drop to arithmetic types pending discussing at Khronos, see discussion here: #3786 (comment).

This was eventually clarified in KhronosGroup/SYCL-Docs#146, but group_broadcast is now explicitly constrained to trivially copyable types. Is a trivially copyable restriction acceptable for your use-case? If not, please open an issue on https://github.com/KhronosGroup/SYCL-Docs/ detailing what you think the restriction should be.

[Pennycook]

#4910 restored support for trivially copyable types. @masterleinad, please close this issue if that is sufficient.