search

KhronosGroup / SYCL-Docs

SYCL Open Source Specification
Other
108 stars 67 forks source link

Reductions and scans: types allowed #340

Open Melirius opened 1 year ago

Melirius commented 1 year ago

In the spec 4.17.4.5 and 4.17.4.6 all the functions are defined for "fundamental types" that are defined in 5.5, Table 183 and only mentioned in the spec once outside of these sections. They do not include sycl::half, however, that is an oversight for the functions supposed to support reduction/scan also of half precision numbers. Two solutions are

  1. Add sycl::half to the table of SYCL supported fundamental types like double (that is also optional on device side). Maybe it is not optimal, taking into account what traits should be guaranteed for C++ fundamental type. Then should SYCL fundamental type imply all the properties of C++ fundamental type? The SYCL supported fundamental types are also called "scalar types" and additional "scalar types" are listed in Table 142. Maybe "scalar types" can be used instead of "fundamental types"? sycl::byte is deprecated, however.
  2. Add sycl::half to every mention of "fundamental types" in reduce/scan functions (or introduce something like "reducible types" = "fundamental types" + sycl::half, and use this term).
nliber commented 1 year ago

I'm not against (2) per se. However, I don't like the direction that every feature has a slightly different subset of types it can be used on.

For instance: we can't have a sycl::atomic_ref<half>, so unless that is just an oversight, that is yet another slightly different subset.

Pennycook commented 1 year ago

For instance: we can't have a sycl::atomic_ref, so unless that is just an oversight, that is yet another slightly different subset.

The reason we can't have sycl::atomic_ref<half> is that atomic capabilities are currently tied to hardware support for 32- and 64-bit atomics. All SYCL devices support 32-bit atomics, but only some support 64-bit atomics (see aspect::atomic64).

If there is a desire to add sycl::half to a list of fundamental types and use that everywhere, I think we'd just need a new way to describe the limits of sycl::atomic_ref<>. For example, we could say that sycl::atomic_ref<> can always be constructed for a sycl::half, but that it may not work on some devices (similar to how we handle support for double and long).

fraggamuffin commented 1 year ago

Test already has sycl half does not affect test progress is it optional or fundamental? definitely optional, but not fundamental according to Table 183 also scalar types should use sycl::byte also what is the notion of sycl::byte, because it is implicitly replaced by std::byte? It is stated in what's changed section but we are still missing std::byte?