diff --git a/sycl/doc/extensions/supported/sycl_ext_intel_usm_address_spaces.asciidoc b/sycl/doc/extensions/supported/sycl_ext_intel_usm_address_spaces.asciidoc index 43067e1e115f9..04c21bda89a78 100644 --- a/sycl/doc/extensions/supported/sycl_ext_intel_usm_address_spaces.asciidoc +++ b/sycl/doc/extensions/supported/sycl_ext_intel_usm_address_spaces.asciidoc @@ -1,103 +1,136 @@ = sycl_ext_intel_usm_address_spaces -== Introduction -This extension introduces two new address spaces and their corresponding multi_ptr specializations. -These address spaces are subsets of the global address space and are added to enable users to provide more optimization information to their compiler. +:source-highlighter: coderay +:coderay-linenums-mode: table -IMPORTANT: This specification is a draft. +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en +:dpcpp: pass:[DPC++] -NOTE: 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. +// 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} -NOTE: This document is better viewed when rendered as html with asciidoctor. GitHub does not render image icons. -This document describes an extension to the SYCL USM extension that adds new explicit address spaces for the possible locations that USM pointers can be allocated. Users can create pointers that point into these address spaces explicitly in order to pass additional information to their compiler so as to enable optimizations. == Notice -Copyright (c) 2020 Intel Corporation. All rights reserved. -== Status +[%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. -Draft -This is a preview extension specification, intended to provide early access to a feature for review and community feedback. When the feature matures, this specification may be released as a formal extension. +== Contact -Because the interfaces defined by this specification are not final and are subject to change they are not intended to be used by shipping software products. +To report problems with this extension, please open a new issue at: -== Version +https://github.com/intel/llvm/issues -Built On: {docdate} + -Revision: 2 == Dependencies -This extension is written against the SYCL 2020 specification, Revision 3. +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. + + +== Status + +This extension is implemented and fully supported by {dpcpp}. + + +== 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. + -If SPIR-V is used by the implementation, this extension also requires support for the SPV_INTEL_usm_storage_classes SPIR-V extension. +== Specification -== Feature Test Macro +=== Feature test macro This extension provides a feature-test macro as described in the core SYCL -specification section 6.3.3 "Feature test macros". Therefore, 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 APIs the implementation supports. +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 |Initial extension version. Base features are supported. -|=== +|Value +|Description -== Overview +|1 +|Initial version of this extension. -This extension adds two new address spaces: device and host that are subsets of the global address space. -New specializations of multi_ptr are added for each of these address spaces. +|2 +|Adds `sycl::ext::intel::host_ptr`, `sycl::ext::intel::raw_host_ptr`, +`sycl::ext::intel::decorated_host_ptr`, `sycl::ext::intel::device_ptr`, +`sycl::ext::intel::raw_device_ptr` and `sycl::ext::intel::decorated_device_ptr`. +`sycl::host_ptr` and `sycl::device_ptr` are deprecated. +|=== -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. +== Modifications to SYCL 2020 -== Modifications to the SYCL Specification, Version 2020 revision 3 +The following sections contain the related changes and additions to the SYCL +2020 specification relating to this extension. -=== Section 3.8.2 SYCL Device Memory Model +=== 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. +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. +*Device memory* is a sub-region of global memory that is not directly accessible +by the host. Buffer accessors and USM allocations whose kind is +`usm::alloc::device` reside in this address space. -=== Section 3.8.2.1 Access to memory +*Host memory* is a sub-region of global memory. USM allocations whose kind is +`usm::alloc::host` reside in this address space. -In the second last paragraph, add sycl::device_ptr and sycl::host_ptr to the list of explicit pointer classes. -=== Section 4.7.7.1 Multi-pointer Class +=== Multi-pointer Class -In the overview of the multi_ptr class replace the address_space enum with the following: +Add the following enumerations to the `access::address_space` enum: ```c++ -enum class address_space : int { - global_space, - local_space, - constant_space, // Deprecated in SYCL 2020 - private_space, - generic_space, +enum class address_space : /* unspecified */ { + ... ext_intel_global_device_space, ext_intel_global_host_space }; ``` -Add the following new conversion operator: +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() const; +explicit operator multi_ptr() const; ``` -Add a new row to Table 91: Constructors of the SYCL multi_ptr class template, as follows: +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"] @@ -105,35 +138,66 @@ Add a new row to Table 91: Constructors of the SYCL multi_ptr class template, as | Constructor | Description a| ```c++ -template - template - multi_ptr( - accessor) -``` | Constructs a multi_ptr from an accessor of access::target::global_buffer. +template +multi_ptr( + accessor); +``` +| 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. |=== -- -=== Section 4.7.7.2 Explicit Pointer Aliases -Add device_ptr and host_ptr aliases to the list of multi_ptr aliases as follows: +=== Explicit Pointer Aliases + +Add `device_ptr` and `host_ptr` aliases to the list of `multi_ptr` aliases as +follows: ```c++ +namespace sycl { + +// Deprecated. +template +using device_ptr = + multi_ptr + +// Deprecated. +template +using host_ptr = + multi_ptr + +namespace ext { +namespace intel { + template -using device_ptr = multi_ptr +using raw_device_ptr = + multi_ptr template -using host_ptr = multi_ptr -``` +using raw_host_ptr = + multi_ptr -== Revision History +template +using decorated_device_ptr = + multi_ptr + +template +using decorated_host_ptr = + multi_ptr + +} // namespace intel +} // namespace ext +} // namespace sycl +``` -[cols="5,15,15,70"] -[grid="rows"] -[options="header"] -|======================================== -|Rev|Date|Author|Changes -|1|2020-06-18|Joe Garvey|Initial public draft -|2|2021-08-30|Dmitry Vodopyanov|Updated according to SYCL 2020 reqs for extensions -|======================================== diff --git a/sycl/include/sycl/access/access.hpp b/sycl/include/sycl/access/access.hpp index 56943580ad81f..c0c605558b282 100644 --- a/sycl/include/sycl/access/access.hpp +++ b/sycl/include/sycl/access/access.hpp @@ -337,6 +337,17 @@ template inline ToT cast_AS(FromT from) { return reinterpret_cast(from); #endif // defined(__NVPTX__) || defined(__AMDGCN__) } else +#ifdef __ENABLE_USM_ADDR_SPACE__ + 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(from); diff --git a/sycl/include/sycl/ext/intel/usm_pointers.hpp b/sycl/include/sycl/ext/intel/usm_pointers.hpp new file mode 100644 index 0000000000000..7c4243e4e7015 --- /dev/null +++ b/sycl/include/sycl/ext/intel/usm_pointers.hpp @@ -0,0 +1,64 @@ +//==-------- usm_pointers.hpp - Extended SYCL pointers classes -------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once +#include + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { + +template +class multi_ptr; + +namespace ext { +namespace intel { + +template +using device_ptr = + multi_ptr; + +template +using host_ptr = + multi_ptr; + +// Template specialization aliases for different pointer address spaces. +// The interface exposes non-decorated pointer while keeping the +// address space information internally. + +template +using raw_device_ptr = + multi_ptr; + +template +using raw_host_ptr = + multi_ptr; + +// Template specialization aliases for different pointer address spaces. +// The interface exposes decorated pointer. + +template +using decorated_device_ptr = + multi_ptr; + +template +using decorated_host_ptr = + multi_ptr; + +} // namespace intel +} // namespace ext +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/include/sycl/feature_test.hpp.in b/sycl/include/sycl/feature_test.hpp.in index a45ce5fac7cf2..65766ddd35d3b 100755 --- a/sycl/include/sycl/feature_test.hpp.in +++ b/sycl/include/sycl/feature_test.hpp.in @@ -62,7 +62,7 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) { #define SYCL_EXT_INTEL_FPGA_REG 1 #define SYCL_EXT_INTEL_KERNEL_ARGS_RESTRICT 1 #define SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY 1 -#define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 1 +#define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 2 #define SYCL_EXT_INTEL_RUNTIME_BUFFER_LOCATION 1 #define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 3 #define SYCL_EXT_ONEAPI_USM_DEVICE_READ_ONLY 1 diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp index 83d479b7e26fa..ea9fa902acd10 100644 --- a/sycl/include/sycl/multi_ptr.hpp +++ b/sycl/include/sycl/multi_ptr.hpp @@ -126,7 +126,8 @@ class multi_ptr { multi_ptr(accessor Accessor) - : multi_ptr(Accessor.get_pointer().get()) {} + : multi_ptr( + detail::cast_AS(Accessor.get_pointer().get())) {} // Only if Space == local_space || generic_space template , Dimensions, Mode, access::target::device, isPlaceholder, PropertyListT> Accessor) - : multi_ptr(Accessor.get_pointer().get()) {} + : multi_ptr( + detail::cast_AS(Accessor.get_pointer().get())) {} // Only if Space == local_space || generic_space and element type is const template { multi_ptr(accessor Accessor) - : multi_ptr(Accessor.get_pointer().get()) {} + : multi_ptr( + detail::cast_AS(Accessor.get_pointer().get())) {} // Only if Space == local_space template < @@ -564,7 +567,8 @@ class multi_ptr { multi_ptr(accessor Accessor) - : multi_ptr(Accessor.get_pointer().get()) {} + : multi_ptr( + detail::cast_AS(Accessor.get_pointer().get())) {} // Only if Space == local_space template < diff --git a/sycl/include/sycl/pointers.hpp b/sycl/include/sycl/pointers.hpp index 7528f813f08b1..2872e0703312d 100644 --- a/sycl/include/sycl/pointers.hpp +++ b/sycl/include/sycl/pointers.hpp @@ -27,15 +27,20 @@ template ; +// Note: Templated alias deprecation is not currently working in clang. See +// https://github.com/llvm/llvm-project/issues/18236. template -using device_ptr = +using device_ptr __SYCL_DEPRECATED( + "Use 'sycl::ext::intel::device_ptr' instead.") = multi_ptr; +// Note: Templated alias deprecation is not currently working in clang. See +// https://github.com/llvm/llvm-project/issues/18236. template -using host_ptr = +using host_ptr __SYCL_DEPRECATED("Use 'sycl::ext::intel::host_ptr' instead.") = multi_ptr; diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index 5664f7270c7d4..48756187350d8 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -60,6 +60,7 @@ #include #endif #include +#include #include #include #include diff --git a/sycl/test/check_device_code/usm_pointers.cpp b/sycl/test/check_device_code/usm_pointers.cpp index fbd005f47afb1..c9d339d2feb88 100644 --- a/sycl/test/check_device_code/usm_pointers.cpp +++ b/sycl/test/check_device_code/usm_pointers.cpp @@ -36,8 +36,8 @@ int main() { queue.submit([&](sycl::handler &cgh) { cgh.single_task([=]() { void *Ptr = nullptr; - device_ptr DevPtr(Ptr); - host_ptr HostPtr(Ptr); + ext::intel::device_ptr DevPtr(Ptr); + ext::intel::host_ptr HostPtr(Ptr); global_ptr GlobPtr = global_ptr(DevPtr); GlobPtr = global_ptr(HostPtr); }); diff --git a/sycl/test/extensions/fpga.cpp b/sycl/test/extensions/fpga.cpp index 7f3a126d49859..48a366e2f8ff1 100644 --- a/sycl/test/extensions/fpga.cpp +++ b/sycl/test/extensions/fpga.cpp @@ -112,8 +112,8 @@ int main() { auto *in_ptr = sycl::malloc_device(1, Queue); Queue.submit([&](sycl::handler &cgh) { cgh.single_task([=]() { - sycl::device_ptr input_ptr(in_ptr); - sycl::device_ptr output_ptr(out_ptr); + sycl::ext::intel::device_ptr input_ptr(in_ptr); + sycl::ext::intel::device_ptr output_ptr(out_ptr); intelfpga::lsu_body< int, sycl::access::address_space::ext_intel_global_device_space>( input_ptr, output_ptr); diff --git a/sycl/test/extensions/usm_pointers_aliases.cpp b/sycl/test/extensions/usm_pointers_aliases.cpp new file mode 100644 index 0000000000000..2a1ec8b4bd0e9 --- /dev/null +++ b/sycl/test/extensions/usm_pointers_aliases.cpp @@ -0,0 +1,66 @@ +// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify %s -o %t.out +// expected-no-diagnostics + +#include + +using namespace sycl; + +int main() { + // Check device_ptr types. + static_assert( + std::is_same_v< + ext::intel::device_ptr, + multi_ptr>, + "Unexpected type for device_ptr"); + static_assert( + std::is_same_v< + ext::intel::device_ptr, + multi_ptr>, + "Unexpected type for device_ptr"); + static_assert( + std::is_same_v< + ext::intel::device_ptr, + multi_ptr>, + "Unexpected type for device_ptr"); + static_assert( + std::is_same_v, + ext::intel::device_ptr>, + "Unexpected type for decorated_device_ptr"); + static_assert( + std::is_same_v, + ext::intel::device_ptr>, + "Unexpected type for raw_device_ptr"); + + // Check host_ptr types. + static_assert( + std::is_same_v< + ext::intel::host_ptr, + multi_ptr>, + "Unexpected type for host_ptr"); + static_assert( + std::is_same_v< + ext::intel::host_ptr, + multi_ptr>, + "Unexpected type for host_ptr"); + static_assert( + std::is_same_v< + ext::intel::host_ptr, + multi_ptr>, + "Unexpected type for host_ptr"); + static_assert( + std::is_same_v, + ext::intel::host_ptr>, + "Unexpected type for decorated_host_ptr"); + static_assert( + std::is_same_v, + ext::intel::host_ptr>, + "Unexpected type for raw_host_ptr"); + + return 0; +}