Skip to content

[SYCL][Docs] Update sycl_ext_intel_usm_address_spaces and fix ctors #7680

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
@@ -0,0 +1,169 @@
= sycl_ext_intel_usm_address_spaces

:source-highlighter: coderay
:coderay-linenums-mode: table

// This section needs to be after the document title.
:doctype: book
:toc2:
:toc: left
:encoding: utf-8
:lang: en
:dpcpp: pass:[DPC++]

// Set the default source code type in this document to C++,
// for syntax highlighting purposes. This is needed because
// docbook uses c++ and html5 uses cpp.
:language: {basebackend@docbook:c++:cpp}


== Notice

[%hardbreaks]
Copyright (C) 2022 Intel Corporation. All rights reserved.

Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by
permission by Khronos.


== Contact

To report problems with this extension, please open a new issue at:

https://github.com/intel/llvm/issues


== Dependencies

This extension is written against the SYCL 2020 revision 6 specification. All
references below to the "core SYCL specification" or to section numbers in the
SYCL specification refer to that revision.

If SPIR-V is used by the implementation, this extension also requires support
for the SPV_INTEL_usm_storage_classes SPIR-V extension.


== Status

This is an experimental extension specification, intended to provide early
access to features and gather community feedback. Interfaces defined in this
specification are implemented in {dpcpp}, but they are not finalized and may
change incompatibly in future versions of {dpcpp} without prior notice.
*Shipping software products should not rely on APIs defined in this
specification.*


== Overview

This extension adds two new address spaces: device and host that are subsets of
the global address space.
New interfaces for `multi_ptr` are added for each of these address spaces.

The goal of this division of the global address space is to enable users to
explicitly tell the compiler which address space a pointer resides in for the
purposes of enabling optimization.
While automatic address space inference is often possible for accessors, it is
harder for USM pointers as it requires inter-procedural optimization with the
host code.
This additional information can be particularly beneficial on FPGA targets where
knowing that a pointer only ever accesses host or device memory can allow
compilers to produce more area efficient memory-accessing hardware.


== Specification

=== Feature test macro

This extension provides a feature-test macro as described in the core SYCL
specification. An implementation supporting this extension must predefine the
macro `SYCL_EXT_INTEL_USM_ADDRESS_SPACES` to one of the values defined in the table
below. Applications can test for the existence of this macro to determine if
the implementation supports this feature, or applications can test the macro's
value to determine which of the extension's features the implementation
supports.

[%header,cols="1,5"]
|===
|Value
|Description

|1
|The APIs of this experimental extension are not versioned, so the
feature-test macro always has this value.
|===

== Modifications to SYCL 2020

The following sections contain the related changes and additions to the SYCL
2020 specification relating to this extension.

=== SYCL Device Memory Model

Add to the end of the definition of global memory:
Global memory is a virtual address space which overlaps the device and host
address spaces.

Add two new memory regions as follows:

*Device memory* is a sub-region of global memory that is not directly accessible
by the host. Global accessors and USM allocations of the device alloc type
reside in this address space.

*Host memory* is a sub-region of global memory. USM pointers allocated with the
host alloc type reside in this address space.


=== Multi-pointer Class

Add the following enumerations to the `access::address_space` enum:
```c++
enum class address_space : /* unspecified */ {
...
ext_intel_global_device_space,
ext_intel_global_host_space
};
```

Add the following new conversion operator to the `multi_ptr` class:
```c++
// Explicit conversion to global_space
// Only available if Space == address_space::ext_intel_global_device_space || Space == address_space::ext_intel_global_host_space
explicit operator multi_ptr<ElementType, access::address_space::global_space, DecorateAddress>() const;
```

Change the `multi_ptr` constructor taking an accessor with `target::device` to
also allow `access::address_space::ext_intel_global_device_space` as follows:

--
[options="header"]
|===
| Constructor | Description
a|
```c++
template <int Dimensions, access_mode Mode, access::placeholder IsPlaceholder>
multi_ptr(
accessor<ElementType, Dimensions, Mode, target::device, IsPlaceholder>);
```
| Available only when:
`Space == access::address_space::global_space \|\| Space == access::address_space::ext_intel_global_device_space \|\| Space == access::address_space::generic_space`.

Constructs a `multi_ptr` from an accessor of `target::device`.

This constructor may only be called from within a command.
|===
--


=== Explicit Pointer Aliases

Add `device_ptr` and `host_ptr` aliases to the list of `multi_ptr` aliases as
follows:
```c++
template<typename ElementType, access::decorated IsDecorated = access::decorated::legacy>
using device_ptr = multi_ptr<ElementType, access::address_space::ext_intel_global_device_space, IsDecorated>

template<typename ElementType, access::decorated IsDecorated = access::decorated::legacy>
using host_ptr = multi_ptr<ElementType, access::address_space::ext_intel_global_host_space, IsDecorated>
```
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Clarify which namespace these are in. I think they are currently in sycl, but that's not the right place. They should either be in sycl::ext::intel or sycl::ext::intel::experimental, depending on the status of this extension.

I think we should deprecate the existing aliases and add new ones in the correct namespace.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Currently it is in the sycl namespace, but I agree they should be moved. Should we also add raw_ and decorated_ variants while we are in here? @GarveyJoe - Are you okay if we just move them or would you prefer we keep the old ones in the sycl namespace as deprecated for now?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please keep the old ones working with deprecation warnings for now so that we can migrate gradually. FYI, @ajaykumarkannan

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I have added a deprecation warning to the sycl::{device|host}_ptr aliases, but it seems that clang isn't handling deprecations of templated aliases correctly at the moment, so users won't be told until that is fixed.

The new pointer aliases are:

  • sycl::ext::intel::device_ptr
  • sycl::ext::intel::raw_device_ptr
  • sycl::ext::intel::decorated_device_ptr
  • sycl::ext::intel::host_ptr
  • sycl::ext::intel::raw_host_ptr
  • sycl::ext::intel::decorated_host_ptr

following the same definitions as other pointer aliases with similar naming in SYCL 2020.


This file was deleted.

11 changes: 11 additions & 0 deletions sycl/include/sycl/access/access.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -356,6 +356,17 @@ template <typename ToT, typename FromT> inline ToT cast_AS(FromT from) {
return reinterpret_cast<ToT>(from);
#endif // defined(__NVPTX__) || defined(__AMDGCN__)
} else
#ifdef __ENABLE_USM_ADDR_SPACE__
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this extension is only enabled if the user defines this macro? That isn't consistent with other extensions and it isn't documented in the extension specification.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

My understanding is that if this isn't defined it will just use the same address space qualifier as for global_space, so the extension is still supported but does not have much of an effect.

if constexpr (FromAS == access::address_space::global_space &&
(ToAS ==
access::address_space::ext_intel_global_device_space ||
ToAS ==
access::address_space::ext_intel_global_host_space)) {
// Casting from global address space to the global device and host address
// spaces is allowed.
return (ToT)from;
} else
#endif // __ENABLE_USM_ADDR_SPACE__
#endif // __SYCL_DEVICE_ONLY__
{
return reinterpret_cast<ToT>(from);
Expand Down
12 changes: 8 additions & 4 deletions sycl/include/sycl/multi_ptr.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -126,7 +126,8 @@ class multi_ptr {
multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::device,
isPlaceholder, PropertyListT>
Accessor)
: multi_ptr(Accessor.get_pointer().get()) {}
: multi_ptr(
detail::cast_AS<decorated_type *>(Accessor.get_pointer().get())) {}

// Only if Space == local_space || generic_space
template <int Dimensions, access::mode Mode,
Expand Down Expand Up @@ -177,7 +178,8 @@ class multi_ptr {
accessor<typename detail::remove_const_t<RelayElementType>, Dimensions,
Mode, access::target::device, isPlaceholder, PropertyListT>
Accessor)
: multi_ptr(Accessor.get_pointer().get()) {}
: multi_ptr(
detail::cast_AS<decorated_type *>(Accessor.get_pointer().get())) {}

// Only if Space == local_space || generic_space and element type is const
template <int Dimensions, access::mode Mode,
Expand Down Expand Up @@ -441,7 +443,8 @@ class multi_ptr<const void, Space, DecorateAddress> {
multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::device,
isPlaceholder, PropertyListT>
Accessor)
: multi_ptr(Accessor.get_pointer().get()) {}
: multi_ptr(
detail::cast_AS<decorated_type *>(Accessor.get_pointer().get())) {}

// Only if Space == local_space
template <
Expand Down Expand Up @@ -566,7 +569,8 @@ class multi_ptr<void, Space, DecorateAddress> {
multi_ptr(accessor<ElementType, Dimensions, Mode, access::target::device,
isPlaceholder, PropertyListT>
Accessor)
: multi_ptr(Accessor.get_pointer().get()) {}
: multi_ptr(
detail::cast_AS<decorated_type *>(Accessor.get_pointer().get())) {}

// Only if Space == local_space
template <
Expand Down