Add shorthand aliases for address spaces

[Pennycook]

Working with address spaces (e.g., when using address space casts) was previously very verbose. Defining shorthands allows long values like sycl::access::address_space::global_space to be replaced by much shorter values like sycl::global_space.

[Pennycook]

To kickstart the bikeshedding... @AerialMantis had suggested that we adopt different names here, that would be more aligned with names like memory_order_relaxed and memory_scope_group, something like address_space_global.

One way we could split the difference between sycl::global_space (19 characters) and sycl::address_space_global (27 characters) would be to define a new alias for access::address_space as well:

namespace sycl {
namespace access {
enum class address_space : int {
  global_space,
  local_space,
  private_space,
  generic_space
};
} // namespace access

using addrspace = access::address_space;
static constexpr addrspace addrspace_global = addrspace::global_space;
static constexpr addrspace addrspace_local = addrspace::local_space;
static constexpr addrspace addrspace_private = addrspace::private_space;
static constexpr addrspace addrspace_generic = addrspace::generic_space;

} // namespace sycl

...which would give us a shorthand of sycl::addrspace_global (23 characters).

Setting a precedent for use of addrspace instead of address_space might help with naming future features, too. The proposed static_address_cast and dynamic_address_cast functions could be renamed static_addrspace_cast and dynamic_addrspace_cast, maintaining the link to address spaces (and the enum) while being only 2 characters longer.

[TApplencourt]

I Agree with Gordon, I would like to keep the consistency of our mangling. I personally have no problem with address_space_global all fancy IDE will do the auto-complete (the problem was the :: IMO always painful to navigate namespace). Also we keep the 3 _ for regularity with other API as extra bonus :p

But I have no problem with addrs either.

(even if I'm still confused between Q.copy, Q.cpy, Q.memcopy, Q.memcpy and never know which one exists and the argument orders. End of my rant about alias)

[Pennycook]

(even if I'm still confused between Q.copy, Q.cpy, Q.memcopy, Q.memcpy and never know which one exists and the argument orders. End of my rant about alias)

I agree that having too many ways to do the same thing is bad. The only reason I'm suggesting to make an exception here is simply to avoid breakage between SYCL 2020 and SYCL-Next.

Between SYCL 1.2.1 and SYCL 2020 we renamed sycl::access::mode to sycl::access_mode and deprecated the old name. I think keeping the nested namespace sycl::access::address_space was a mistake, and if we'd caught this earlier we could have just had sycl::addrspace in SYCL 2020. My personal preference would still be to deprecate the old name and move everything to the new one... But I understand that this is an ABI/API concern, and so I could live with an alias. 😄

[TApplencourt]

My personal preference would still be to deprecate the old name and move everything to the new one... But I understand that this is an ABI/API concern, and so I could live with an alias

Agree! Should we deprecate it now, so we can remove it in SYCL Next?

[Pennycook]

Agree! Should we deprecate it now, so we can remove it in SYCL Next?

We should talk about this in a WG meeting.

I don't think we can deprecate things if we go the alias route. Adding [[deprecated]] to address_space would cause a deprecation warning to trigger every time somebody used addrspace (because it's only an alias).

I believe our options are:

1. Define addrspace as an alias of address_space, but keep address_space around forever for legacy code.
2. Re-define address_space as an alias of addrspace, and deprecate address_space (ABI break)
3. Define addrspace as an alias of address_space for now, planning to swap them some day far in the future.

[gmlueck]

There also a 4th option, which I'll throw out:

4. Define addrspace as a separate enumeration and add new APIs for functions that currently take address_space.

This would allow us to deprecate and eventually remove address_space and any API that uses it without incurring an ABI break.

Many of the uses of address_space are from multi_ptr, which we may decide to deprecate and replace anyways. When I noticed this, I was thinking that option (4) would be a good path forward. However, then I realized that we also use address_space extensively in atomic_ref, so option (4) seemed less ideal.

[keryell]

I am thinking to local_as, constant_as, private_as, generic_as.

as_local sounds more like a cast in LLVM/MLIR style.

[Pennycook]

I am thinking to local_as, constant_as, private_as, generic_as.

as_local sounds more like a cast in LLVM/MLIR style.

I do like that these are much shorter, but I think that "as" might be too short. Like you pointed out, it's very easy to read it as the English word "as" instead of as an initialism. Also, because we don't have any precedent for this elsewhere in SYCL, a new developer probably won't be able to guess what it means without looking at the specification.

[keryell]

I am thinking to local_as, constant_as, private_as, generic_as.
as_local sounds more like a cast in LLVM/MLIR style.

I do like that these are much shorter, but I think that "as" might be too short. Like you pointed out, it's very easy to read it as the English word "as" instead of as an initialism. Also, because we don't have any precedent for this elsewhere in SYCL, a new developer probably won't be able to guess what it means without looking at the specification.

Any reasonable developer not using punch cards will move the mouse on the word and the explanation on what is the meaning of private_as will pop up on screen thanks to LSP mode.
Is there any real usage of this with a code sample to see how confusing it would be?

[Pennycook]

Any reasonable developer not using punch cards will move the mouse on the word and the explanation on what is the meaning of private_as will pop up on screen thanks to LSP mode.

This is true during development, but the first time somebody sees these shorthands might be in a printed book, on a website, or some sort of training presentation. That's the main reason I'm leaning towards readable names, even if they're longer.

Is there any real usage of this with a code sample to see how confusing it would be?

I couldn't find a good real usage showing all the address-space functionality together, so I wrote something I think is fairly representative.

Original Proposal

void increment(sycl::multi_ptr<float, sycl::generic_space> generic_ptr) {
  auto global_ptr = sycl::static_address_cast<sycl::global_space>(generic_ptr);
  sycl::atomic_ref<float,
                   sycl::memory_order::relaxed,
                   sycl::memory_scope::device,
                   sycl::global_space>(*global_ptr)++;
}

Gordon's Proposal

void increment(sycl::multi_ptr<float, sycl::address_space_generic> generic_ptr) {
  auto global_ptr = sycl::static_address_cast<sycl::address_space_global>(generic_ptr);
  sycl::atomic_ref<float,
                   sycl::memory_order::relaxed,
                   sycl::memory_scope::device,
                   sycl::address_space_global>(*global_ptr)++;
}

John's Modified Proposal

void increment(sycl::multi_ptr<float, sycl::addrspace_generic> generic_ptr) {
  auto global_ptr = sycl::static_addrspace_cast<sycl::addrspace_global>(generic_ptr);
  sycl::atomic_ref<float,
                   sycl::memory_order::relaxed,
                   sycl::memory_scope::device,
                   sycl::addrspace_global>(*global_ptr)++;
}

Note that my modified proposal is to rename address_cast to addrspace_cast in line with the new enum name.

Ronan's Proposal

void increment(sycl::multi_ptr<float, sycl::generic_as> generic_ptr) {
  auto global_ptr = sycl::static_as_cast<sycl::global_as>(generic_ptr);
  sycl::atomic_ref<float,
                   sycl::memory_order::relaxed,
                   sycl::memory_scope::device,
                   sycl::global_as>(*global_ptr)++;
}

I know you didn't actually suggest static_as_cast, but I wanted to show what it would look like if we decided to adopt the same convention for "address space" everywhere that it appears.

[Pennycook]

I've updated the PR to reflect the discussion in today's WG meeting.

[keryell]

Since this is no longer constrained by namespace hierarchy, why not following English like
constexpr inline addrspace generic_addrspace = addrspace::generic_space;
?

[Pennycook]

@AerialMantis' suggestion was that we should always use enum_value, because that's the precedent elsewhere (e.g., memory_order_relaxed, memory_scope_device).

[keryell]

But at the end I would prefer

constexpr inline addrspace global_as = addrspace::global_space;
constexpr inline addrspace local_as = addrspace::local_space;
constexpr inline addrspace private_as = addrspace::private_space;
constexpr inline addrspace generic_as = addrspace::generic_space;

to have shorthand alias to be short.