diff --git a/sycl/doc/CompilerAndRuntimeDesign.md b/sycl/doc/CompilerAndRuntimeDesign.md index 9af90e248b6c..5a408e5b8b14 100644 --- a/sycl/doc/CompilerAndRuntimeDesign.md +++ b/sycl/doc/CompilerAndRuntimeDesign.md @@ -915,8 +915,8 @@ space attributes in SYCL mode: | Address space attribute | SYCL address_space enumeration | |-------------------------|--------------------------------| | `__attribute__((opencl_global))` | global_space, constant_space | -| `__attribute__((opencl_global_host))` | global_host_space | -| `__attribute__((opencl_global_device))` | global_device_space | +| `__attribute__((opencl_global_host))` | ext_intel_global_host_space | +| `__attribute__((opencl_global_device))` | ext_intel_global_device_space | | `__attribute__((opencl_local))` | local_space | | `__attribute__((opencl_private))` | private_space | | `__attribute__((opencl_constant))` | N/A diff --git a/sycl/doc/extensions/EnqueueBarrier/enqueue_barrier.asciidoc b/sycl/doc/extensions/EnqueueBarrier/enqueue_barrier.asciidoc index 76cd475c27bd..880313dae5af 100644 --- a/sycl/doc/extensions/EnqueueBarrier/enqueue_barrier.asciidoc +++ b/sycl/doc/extensions/EnqueueBarrier/enqueue_barrier.asciidoc @@ -1,4 +1,4 @@ -= SYCL_INTEL_enqueue_barrier += SYCL_EXT_ONEAPI_ENQUEUE_BARRIER :source-highlighter: coderay :coderay-linenums-mode: table @@ -25,11 +25,6 @@ NOTE: This document is better viewed when rendered as html with asciidoctor. Gi This document presents a series of changes proposed for a future version of the SYCL Specification. The goal of this proposal is to provide non-blocking APIs that provide synchronization on SYCL command queue for programmers. - -== Name Strings - -+SYCL_INTEL_enqueue_barrier+ - == Notice Copyright (c) 2019-2020 Intel Corporation. All rights reserved. @@ -45,19 +40,35 @@ Because the interfaces defined by this specification are not final and are subje == Version Built On: {docdate} + -Revision: 1 +Revision: 2 == Contact Please open an issue in the https://github.com/intel/llvm/tree/sycl/sycl/doc/extensions/[extensions repository] +== 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_ONEAPI_ENQUEUE_BARRIER` 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. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|=== + == Dependencies -This extension is written against the SYCL 1.2.1 specification, Revision v1.2.1-6. +This extension is written against the SYCL 2020 specification, revision 3. == Overview -SYCL 1.2.1 defines a graph-based task execution model, based on kernels or explicit memory operations submitted to out-of-order queues. Dependencies between these kernels are represented by -accessors that form data dependence edges in the execution graph. The USM extension <> doesn't have accessors, so instead solves +SYCL 2020 defines a graph-based task execution model, based on kernels or explicit memory operations submitted to out-of-order queues. Dependencies between these kernels are represented by +accessors that form data dependence edges in the execution graph. Unified Shared Memory (USM) doesn't have accessors, so instead solves this by defining `handler::depends_on` methods to specify event-based control dependencies between command groups. There are situations where defining dependencies based on events is more explicit than desired or required by an application. For instance, the user may know that a given task depends on all previously submitted tasks. Instead of explicitly adding all the required depends_on calls, the user could express this intent via a single call, making the program more concise and explicit. @@ -75,9 +86,9 @@ two new members to the `queue` class: [grid="rows"] [options="header"] |======================================== -|*handler::barrier*|*queue::submit_barrier* -|`void barrier()` | `event submit_barrier()` -|`void barrier( const vector_class &waitList )` | `event submit_barrier( const vector_class &waitList )` +|*handler::ext_oneapi_barrier*|*queue::ext_oneapi_submit_barrier* +|`void ext_oneapi_barrier()` | `event ext_oneapi_submit_barrier()` +|`void ext_oneapi_barrier( const vector_class &waitList )` | `event ext_oneapi_submit_barrier( const vector_class &waitList )` |======================================== The first variant of the barrier takes no parameters, and waits for all previously submitted commands to the queue to enter the `info::event_command_status::complete` state before any command later submitted to the same queue is allowed to execute. A second variant of the barrier accepts a list of events, with the behavior that no commands submitted to the same queue after barrier submission may execute until all events in the `waitList` have entered the `info::event_command_status::complete` state. Both variants are non-blocking from the host program perspective, in that they do not wait for the barrier conditions to have been met before returning. @@ -93,7 +104,7 @@ Some forms of the new barrier methods return an `event`, which can be used to pe CG4 doesn't execute until all previous command groups submitted to the same queue (CG1, CG2, CG3) have entered the completed state. -==== 1. Using `handler::barrier()`: +==== 1. Using `handler::ext_oneapi_barrier()`: [source,c++,NoName,linenums] ---- @@ -109,7 +120,7 @@ Queue.submit([&](cl::sycl::handler& cgh) { }); Queue.submit([&](cl::sycl::handler& cgh) { - cgh.barrier(); + cgh.ext_oneapi_barrier(); }); Queue.submit([&](cl::sycl::handler& cgh) { @@ -118,7 +129,7 @@ Queue.submit([&](cl::sycl::handler& cgh) { ... ---- -==== 2. Using `queue::submit_barrier()`: +==== 2. Using `queue::ext_oneapi_submit_barrier()`: [source,c++,NoName,linenums] ---- @@ -133,7 +144,7 @@ Queue.submit([&](cl::sycl::handler& cgh) { // CG3 }); -Queue.submit_barrier(); +Queue.ext_oneapi_submit_barrier(); Queue.submit([&](cl::sycl::handler& cgh) { // CG4 @@ -146,7 +157,7 @@ Queue.submit([&](cl::sycl::handler& cgh) { CG3 requires CG1 (in Queue1) and CG2 (in Queue2) to have completed before it (CG3) begins execution. -==== 1. Using `handler::barrier()`: +==== 1. Using `handler::ext_oneapi_barrier()`: [source,c++,NoName,linenums] ---- @@ -160,7 +171,7 @@ auto event_barrier2 = Queue2.submit([&](cl::sycl::handler& cgh) { }); Queue3.submit([&](cl::sycl::handler& cgh) { - cgh.barrier( vector_class{event_barrier1, event_barrier2} ); + cgh.ext_oneapi_barrier( vector_class{event_barrier1, event_barrier2} ); }); Queue3.submit([&](cl::sycl::handler& cgh) { @@ -169,7 +180,7 @@ Queue3.submit([&](cl::sycl::handler& cgh) { ... ---- -==== 2. Using `queue::submit_barrier()`: +==== 2. Using `queue::ext_oneapi_submit_barrier()`: [source,c++,NoName,linenums] ---- @@ -182,7 +193,7 @@ auto event_barrier2 = Queue2.submit([&](cl::sycl::handler& cgh) { // CG2 }); -Queue3.submit_barrier( vector_class{event_barrier1, event_barrier2} ); +Queue3.ext_oneapi_submit_barrier( vector_class{event_barrier1, event_barrier2} ); Queue3.submit([&](cl::sycl::handler& cgh) { // CG3 @@ -211,44 +222,45 @@ void wait(); template event submit(T cgf, const queue &secondaryQueue); -event submit_barrier(); +event ext_oneapi_submit_barrier(); -event submit_barrier( const vector_class &waitList ); +event ext_oneapi_submit_barrier( const vector_class &waitList ); void wait(); ... ---- -=== Add rows to Table 4.22 +=== Add rows to Table 28 [cols="70,300"] [grid="rows"] [options="header"] |======================================== |*Member functions*|*Description* -|`event submit_barrier()` | Same effect as submitting a `handler::barrier()` within a command group to this `queue`. The returned event enters the `info::event_command_status::complete` state when all events that the barrier is dependent on (implicitly from all previously submitted commands to the same queue) have entered the `info::event_command_status::complete` state. -|`event submit_barrier( const vector_class &waitList )` | Same effect as submitting a `handler:barrier( const vector_class &waitList )` within a command group to this `queue`. The returned event enters the `info::event_command_status::complete` state when all events that the barrier is dependent on (explicitly from `waitList`) have entered the `info::event_command_status::complete` state. +|`event ext_oneapi_submit_barrier()` | Same effect as submitting a `handler::ext_oneapi_barrier()` within a command group to this `queue`. The returned event enters the `info::event_command_status::complete` state when all events that the barrier is dependent on (implicitly from all previously submitted commands to the same queue) have entered the `info::event_command_status::complete` state. +|`event ext_oneapi_submit_barrier( const vector_class &waitList )` | Same effect as submitting a `handler:ext_oneapi_barrier( const vector_class &waitList )` within a command group to this `queue`. The returned event enters the `info::event_command_status::complete` state when all events that the barrier is dependent on (explicitly from `waitList`) have entered the `info::event_command_status::complete` state. |======================================== -=== Modify Section 4.8.2 +=== Modify Section 4.9.3 ==== Change first sentence from: -A command group scope in SYCL, as it is defined in Section 3.4.1, consists of a single kernel or explicit memory -operation (handler methods such as copy, update_host, fill), together with its requirements. +The member functions and objects defined in this scope will define the requirements for the kernel execution or +explicit memory operation, and will be used by the SYCL runtime to evaluate if the operation is ready for execution. ==== To: -A command group scope in SYCL, as it is defined in Section 3.4.1, consists of a single kernel, explicit memory -operation (handler methods such as copy, update_host, fill) or barrier, together with its requirements. +The member functions and objects defined in this scope will define the requirements for the kernel execution, +explicit memory operation or barrier, and will be used by the SYCL runtime to evaluate if the operation is ready for execution. + -=== Modify part of Section 4.8.3 +=== Modify part of Section 4.9.4 *Change from:* [source,c++,NoName,linenums] ---- ... -template -void fill(accessor dest, const T& src); +template +void fill(void *ptr, const T &pattern, size_t count); }; ... @@ -258,39 +270,36 @@ void fill(accessor dest, const T& src); [source,c++,NoName,linenums] ---- ... -template -void fill(accessor dest, const T& src); +template +void fill(void *ptr, const T &pattern, size_t count); -void barrier(); +void ext_oneapi_barrier(); -void barrier( const vector_class &waitList ); +void ext_oneapi_barrier( const vector_class &waitList ); }; ... ---- -=== Add a new section between Section 4.8.6 and 4.8.7 +=== Add a new section between Section 4.9.4 and 4.9.5 -4.8.X SYCL functions for enqueued synchronization barriers +4.9.X SYCL functions for enqueued synchronization barriers Barriers may be submitted to a queue, with the effect that they prevent later operations submitted to the same queue from executing until the barrier wait conditions have been satisfied. The wait conditions can be explicitly described by `waitList` or implicitly from all previously submitted commands to the same queue. There are no constraints on the context from which queues may participate in the `waitList`. Enqueued barriers do not block host program execution, but instead form additional dependence edges with the execution task graph. Barriers can be created by two members of the `handler` class that force synchronization on the SYCL command queue. The first variant of the `handler` barrier (`handler::barrier()`) takes no parameters, and waits for all previously submitted commands to the queue to enter the `info::event_command_status::complete` state before any command later submitted to the same queue is allowed to execute. The second variant of the `handler` barrier (`handler::barrier( const vector_class &waitList )`) accepts a list of events, with the behavior that no commands submitted to the same queue after barrier submission may execute until all events in the waitList have entered the `info::event_command_status::complete` state. -=== Add a new table in the new section between 4.8.6 and 4.8.7: Member functions of the handler class. +=== Add a new table in the new section between 4.9.4 and 4.9.5: Member functions of the handler class. [cols="70,300"] [grid="rows"] [options="header"] |======================================== |*Member functions*|*Description* -|`void barrier()` | Prevents any commands submitted afterward to this queue from executing until all commands previously submitted to this queue have entered the `info::event_command_status::complete` state. -|`void barrier( const vector_class &waitList` ) | Prevents any commands submitted afterward to this queue from executing until all events in `waitList` have entered the `info::event_command_status::complete` state. If `waitList` is empty, then the barrier has no effect. +|`void ext_oneapi_barrier()` | Prevents any commands submitted afterward to this queue from executing until all commands previously submitted to this queue have entered the `info::event_command_status::complete` state. +|`void ext_oneapi_barrier( const vector_class &waitList` ) | Prevents any commands submitted afterward to this queue from executing until all events in `waitList` have entered the `info::event_command_status::complete` state. If `waitList` is empty, then the barrier has no effect. |======================================== -== References -1. [[usmlink]]https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/USM/USM.adoc - == Issues None. @@ -303,6 +312,7 @@ None. |======================================== |Rev|Date|Author|Changes |1|2020-02-26|Ye Ting|*Initial public release* +|2|2021-08-30|Dmitry Vodopyanov|*Updated according to SYCL 2020 reqs for extensions* |======================================== //************************************************************************ diff --git a/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md b/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md index ce8b00d2ef51..55ccf83219f1 100644 --- a/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md +++ b/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md @@ -9,7 +9,6 @@ The currently supported targets are all Intel GPUs starting with Gen9. NOTE: This specification is a draft. While describing the currently implemented behaviors it is known to be not complete nor exhaustive. We shall continue to add more information, e.g. explain general mapping of SYCL programming model to Level-Zero API. - It will also be gradually changing to a SYCL-2020 conforming implementation. ## 2. Prerequisites @@ -23,7 +22,7 @@ The Level-Zero backend is added to the cl::sycl::backend enumeration: ``` C++ enum class backend { // ... - level_zero, + ext_oneapi_level_zero, // ... }; ``` @@ -55,7 +54,7 @@ and they must be included in the order shown: ``` C++ #include "level_zero/ze_api.h" - #include "sycl/backend/level_zero.hpp" + #include "sycl/ext/oneapi/backend/level_zero.hpp" ``` ### 4.1 Mapping of SYCL objects to Level-Zero handles @@ -71,7 +70,7 @@ These SYCL objects encapsulate the corresponding Level-Zero handles: ### 4.2 Obtaining of native Level-Zero handles from SYCL objects -The ```get_native()``` member function is how a raw native Level-Zero handle can be obtained +The ```get_native()``` member function is how a raw native Level-Zero handle can be obtained for a specific SYCL object. It is currently supported for SYCL ```platform```, ```device```, ```context```, ```queue```, ```event``` and ```program``` classes. There is also a free-function defined in ```cl::sycl``` namespace that can be used instead of the member function: ``` C++ @@ -81,7 +80,7 @@ auto get_native(const SyclObjectT &Obj) -> ``` ### 4.3 Construct a SYCL object from a Level-Zero handle -The following free functions defined in the ```cl::sycl::level_zero``` namespace allow an application to create +The following free functions defined in the ```cl::sycl::ext::oneapi::level_zero``` namespace allow an application to create a SYCL object that encapsulates a corresponding Level-Zero object: | Level-Zero interoperability function |Description| @@ -103,11 +102,15 @@ some interoperability API supports overriding this behavior and keep the ownersh Use this enumeration for explicit specification of the ownership: ``` C++ namespace sycl { +namespace ext { +namespace oneapi { namespace level_zero { enum class ownership { transfer, keep }; } // namespace level_zero +} // namespace oneapi +} // namespace ext } // namespace sycl ``` @@ -193,3 +196,4 @@ struct free_memory { |3|2021-04-13|James Brodman|Free Memory Query |4|2021-07-06|Rehana Begam|Introduced explicit ownership for queue |5|2021-07-25|Sergey Maslov|Introduced SYCL interop for events +|6|2021-08-30|Dmitry Vodopyanov|Updated according to SYCL 2020 reqs for extensions diff --git a/sycl/doc/extensions/MemChannel/SYCL_INTEL_mem_channel_property.asciidoc b/sycl/doc/extensions/MemChannel/MemChannel.asciidoc similarity index 74% rename from sycl/doc/extensions/MemChannel/SYCL_INTEL_mem_channel_property.asciidoc rename to sycl/doc/extensions/MemChannel/MemChannel.asciidoc index 15b309851043..50a0258a3297 100644 --- a/sycl/doc/extensions/MemChannel/SYCL_INTEL_mem_channel_property.asciidoc +++ b/sycl/doc/extensions/MemChannel/MemChannel.asciidoc @@ -1,4 +1,4 @@ -= SYCL_INTEL_mem_channel_property += SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY == Introduction 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. @@ -23,14 +23,30 @@ Because the interfaces defined by this specification are not final and are subje == Version Built On: {docdate} + -Revision: 1 +Revision: 2 == Dependencies -This extension is written against the SYCL 2020 provisional specification, Revision 1. +This extension is written against the SYCL 2020 specification, Revision 3. The use of this extension requires a target that supports cl_intel_mem_channel_property or equivalent if OpenCL is used as the underlying device runtime. +== 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_MEM_CHANNEL_PROPERTY` 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. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|=== + == Overview On some targets manual assignment of buffers to memory regions can improve memory bandwidth. This extension adds a buffer property to indicate in which memory channel a particular buffer should be allocated. This information is an optimization hint to the runtime and thus it is legal to ignore. @@ -59,7 +75,7 @@ Add a new constructor to Table 4.34: Constructors of the buffer property classes |=== -- -Add a new member function to Table 4.35: Member functions of the buffer property classes as follows: +Add a new member function to Table 42: Member functions of the buffer property classes as follows: -- [options="header"] @@ -87,7 +103,7 @@ enum class aspect { } // namespace sycl ``` -Add an entry for the new aspect to Table 4.20: Device aspects defined by the core SYCL specification: +Add an entry for the new aspect to Table 26: Device aspects defined by the core SYCL specification: -- [options="header"] @@ -107,4 +123,5 @@ Add an entry for the new aspect to Table 4.20: Device aspects defined by the cor |======================================== |Rev|Date|Author|Changes |1|2020-10-26|Joe Garvey|*Initial public draft* +|2|2021-08-30|Dmitry Vodopyanov|*Updated according to some SYCL 2020 reqs for extensions* |======================================== diff --git a/sycl/doc/extensions/SYCL_ONEAPI_dot_accumulate.asciidoc b/sycl/doc/extensions/SYCL_ONEAPI_dot_accumulate.asciidoc index 64471166407b..de7622f6121c 100755 --- a/sycl/doc/extensions/SYCL_ONEAPI_dot_accumulate.asciidoc +++ b/sycl/doc/extensions/SYCL_ONEAPI_dot_accumulate.asciidoc @@ -46,7 +46,7 @@ Because the interfaces defined by this specification are not final and are subje == Version Built On: {docdate} + -Revision: B +Revision: 3 == Contact @@ -135,9 +135,9 @@ None. [options="header"] |======================================== |Rev|Date|Author|Changes -|A|2019-12-13|Ben Ashbaugh|*Initial draft* -|B|2019-12-18|Ben Ashbaugh|Switched to standard C++ fixed width types. -|C|2020-10-26|Rajiv Deodhar|Added int32 types. +|1|2019-12-13|Ben Ashbaugh|*Initial draft* +|2|2019-12-18|Ben Ashbaugh|Switched to standard C++ fixed width types. +|3|2020-10-26|Rajiv Deodhar|Added int32 types. |======================================== //************************************************************************ diff --git a/sycl/doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc b/sycl/doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc index 73e9de475890..138f2874919a 100644 --- a/sycl/doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc +++ b/sycl/doc/extensions/USMAddressSpaces/usm_address_spaces.asciidoc @@ -1,4 +1,4 @@ -= SYCL_INTEL_usm_address_spaces += SYCL_EXT_INTEL_USM_ADDRESS_SPACES == Introduction This extension introduces two new address spaces and their corresponding multi_ptr specializations. @@ -11,9 +11,6 @@ NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are tradema 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. -== Name Strings -+SYCL_INTEL_usm_address_spaces+ - == Notice Copyright (c) 2020 Intel Corporation. All rights reserved. @@ -28,14 +25,30 @@ Because the interfaces defined by this specification are not final and are subje == Version Built On: {docdate} + -Revision: 1 +Revision: 2 == Dependencies -This extension is written against the SYCL 1.2.1 specification, Revision 7. It requires the Unified Shared Memory SYCL proposal. +This extension is written against the SYCL 2020 specification, Revision 3. If SPIR-V is used by the implementation, this extension also requires support for the SPV_INTEL_usm_storage_classes SPIR-V extension. +== 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. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|=== + == Overview This extension adds two new address spaces: device and host that are subsets of the global address space. @@ -45,9 +58,9 @@ The goal of this division of the global address space is to enable users to expl 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 the SYCL Specification, Version 1.2.1 revision 7 +== Modifications to the SYCL Specification, Version 2020 revision 3 -=== Section 3.5.2 SYCL Device Memory Model +=== Section 3.8.2 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. @@ -58,32 +71,33 @@ Add two new memory regions as follows: *Host memory* is a sub-region of global memory. USM pointers allocated with the host alloc type reside in this address space. -=== Section 3.5.2.1 Access to memory +=== Section 3.8.2.1 Access to memory -In the second last paragraph, add cl::sycl::device_ptr and cl::sycl::host_ptr to the list of explicit pointer classes. +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 In the overview of the multi_ptr class replace the address_space enum with the following: ```c++ enum class address_space : int { - global_space, - local_space, - constant_space, - private_space, - device_space, - host_space + global_space, + local_space, + constant_space, // Deprecated in SYCL 2020 + private_space, + generic_space, + ext_intel_global_device_space, + ext_intel_global_host_space }; ``` Add the following new conversion operator: ```c++ // Explicit conversion to global_space -// Only available if Space == address_space::device_space || Space == address_space::host_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; ``` -Add a new row to Table 4.54: Constructors of the SYCL multi_ptr class template, as follows: +Add a new row to Table 91: Constructors of the SYCL multi_ptr class template, as follows: -- [options="header"] @@ -93,12 +107,12 @@ a| ```c++ template +ext_intel_global_device_space> template multi_ptr( accessor) -``` | Constructs a multi_ptr from an accessor of access::target::global_buffer. +``` | Constructs a multi_ptr from an accessor of access::target::global_buffer. |=== -- @@ -107,10 +121,10 @@ device_space> Add device_ptr and host_ptr aliases to the list of multi_ptr aliases as follows: ```c++ template -using device_ptr = multi_ptr +using device_ptr = multi_ptr template -using host_ptr = multi_ptr +using host_ptr = multi_ptr ``` == Revision History @@ -120,5 +134,6 @@ using host_ptr = multi_ptr [options="header"] |======================================== |Rev|Date|Author|Changes -|A|2020-06-18|Joe Garvey|Initial public draft +|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/doc/extensions/accessor_properties/SYCL_INTEL_buffer_location.asciidoc b/sycl/doc/extensions/accessor_properties/SYCL_INTEL_buffer_location.asciidoc index d92a228aeba7..cc462dc7e34f 100644 --- a/sycl/doc/extensions/accessor_properties/SYCL_INTEL_buffer_location.asciidoc +++ b/sycl/doc/extensions/accessor_properties/SYCL_INTEL_buffer_location.asciidoc @@ -31,7 +31,7 @@ Because the interfaces defined by this specification are not final and are subje == Version Built On: {docdate} + -Revision: A +Revision: 1 == Contact Michael Kinsner, Intel (michael 'dot' kinsner 'at' intel 'dot' com) @@ -95,5 +95,5 @@ It also notifies the SYCL runtime to store the given accessor in that memory. | [options="header"] |======================================== |Rev|Date|Author|Changes -|A|2020-09-08|Joe Garvey|*Initial public draft* +|1|2020-09-08|Joe Garvey|*Initial public draft* |======================================== diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 8f4b9a92804e..cba9308111cc 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -46,8 +46,14 @@ enum class address_space : int { global_space = 1, constant_space = 2, local_space = 3, - global_device_space = 4, - global_host_space = 5 + ext_intel_global_device_space = 4, + ext_intel_host_device_space = 5, + global_device_space __SYCL2020_DEPRECATED( + "use 'ext_intel_global_device_space' instead") = + ext_intel_global_device_space, + global_host_space __SYCL2020_DEPRECATED( + "use 'ext_intel_host_device_space' instead") = + ext_intel_host_device_space, }; } // namespace access diff --git a/sycl/include/CL/sycl/backend/level_zero.hpp b/sycl/include/CL/sycl/backend/level_zero.hpp index 9661078915c6..8d808fec3e5e 100644 --- a/sycl/include/CL/sycl/backend/level_zero.hpp +++ b/sycl/include/CL/sycl/backend/level_zero.hpp @@ -8,186 +8,9 @@ #pragma once -#include -// This header should be included by users. -//#include +#include -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { +__SYCL_WARNING("CL/sycl/backend/level_zero.hpp usage is deprecated, include " + "sycl/ext/oneapi/backend/level_zero.hpp instead") -template <> struct interop { - using type = ze_driver_handle_t; -}; - -template <> struct interop { - using type = ze_device_handle_t; -}; - -template <> struct interop { - using type = ze_context_handle_t; -}; - -template <> struct interop { - using type = ze_command_queue_handle_t; -}; - -template <> struct interop { - using type = ze_event_handle_t; -}; - -template <> struct interop { - using type = ze_module_handle_t; -}; - -template -struct interop> { - using type = char *; -}; - -template -struct interop> { - using type = char *; -}; - -template -struct interop> { - using type = ze_image_handle_t; -}; - -namespace level_zero { -// Since Level-Zero is not doing any reference counting itself, we have to -// be explicit about the ownership of the native handles used in the -// interop functions below. -// -enum class ownership { transfer, keep }; -} // namespace level_zero - -namespace detail { - -template <> struct BackendInput { - using type = struct { - interop::type NativeHandle; - std::vector DeviceList; - level_zero::ownership Ownership; - }; -}; - -template <> struct BackendReturn { - using type = ze_kernel_handle_t; -}; - -template <> struct InteropFeatureSupportMap { - static constexpr bool MakePlatform = true; - static constexpr bool MakeDevice = true; - static constexpr bool MakeContext = true; - static constexpr bool MakeQueue = false; - static constexpr bool MakeEvent = true; - static constexpr bool MakeBuffer = false; - static constexpr bool MakeKernel = false; - static constexpr bool MakeKernelBundle = false; -}; -} // namespace detail - -namespace level_zero { -// Implementation of various "make" functions resides in libsycl.so and thus -// their interface needs to be backend agnostic. -// TODO: remove/merge with similar functions in sycl::detail -__SYCL_EXPORT platform make_platform(pi_native_handle NativeHandle); -__SYCL_EXPORT device make_device(const platform &Platform, - pi_native_handle NativeHandle); -__SYCL_EXPORT context make_context(const std::vector &DeviceList, - pi_native_handle NativeHandle, - bool keep_ownership = false); -__SYCL_EXPORT program make_program(const context &Context, - pi_native_handle NativeHandle); -__SYCL_EXPORT queue make_queue(const context &Context, - pi_native_handle InteropHandle, - bool keep_ownership = false); -__SYCL_EXPORT event make_event(const context &Context, - pi_native_handle InteropHandle, - bool keep_ownership = false); - -// Construction of SYCL platform. -template ::value> * = nullptr> -__SYCL_DEPRECATED("Use SYCL-2020 sycl::make_platform free function") -T make(typename interop::type Interop) { - return make_platform(reinterpret_cast(Interop)); -} - -// Construction of SYCL device. -template ::value> * = nullptr> -__SYCL_DEPRECATED("Use SYCL-2020 sycl::make_device free function") -T make(const platform &Platform, - typename interop::type Interop) { - return make_device(Platform, reinterpret_cast(Interop)); -} - -/// Construction of SYCL context. -/// \param DeviceList is a vector of devices which must be encapsulated by -/// created SYCL context. Provided devices and native context handle must -/// be associated with the same platform. -/// \param Interop is a Level Zero native context handle. -/// \param Ownership (optional) specifies who will assume ownership of the -/// native context handle. Default is that SYCL RT does, so it destroys -/// the native handle when the created SYCL object goes out of life. -/// -template ::value>::type * = nullptr> -__SYCL_DEPRECATED("Use SYCL-2020 sycl::make_context free function") -T make(const std::vector &DeviceList, - typename interop::type Interop, - ownership Ownership = ownership::transfer) { - return make_context(DeviceList, detail::pi::cast(Interop), - Ownership == ownership::keep); -} - -// Construction of SYCL program. -template ::value> * = nullptr> -T make(const context &Context, - typename interop::type Interop) { - return make_program(Context, reinterpret_cast(Interop)); -} - -// Construction of SYCL queue. -template ::value> * = nullptr> -T make(const context &Context, - typename interop::type Interop, - ownership Ownership = ownership::transfer) { - return make_queue(Context, reinterpret_cast(Interop), - Ownership == ownership::keep); -} - -// Construction of SYCL event. -template ::value> * = nullptr> -T make(const context &Context, - typename interop::type Interop, - ownership Ownership = ownership::transfer) { - return make_event(Context, reinterpret_cast(Interop), - Ownership == ownership::keep); -} -} // namespace level_zero - -// Specialization of sycl::make_context for Level-Zero backend. -template <> -context make_context( - const backend_input_t &BackendObject, - const async_handler &Handler) { - return level_zero::make_context( - BackendObject.DeviceList, - detail::pi::cast(BackendObject.NativeHandle), - BackendObject.Ownership == level_zero::ownership::keep); -} - -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) +#include diff --git a/sycl/include/CL/sycl/backend_types.hpp b/sycl/include/CL/sycl/backend_types.hpp index 0652adb14c2e..28c05a2ad215 100644 --- a/sycl/include/CL/sycl/backend_types.hpp +++ b/sycl/include/CL/sycl/backend_types.hpp @@ -21,7 +21,9 @@ namespace sycl { enum class backend : char { host = 0, opencl = 1, - level_zero = 2, + ext_oneapi_level_zero = 2, + level_zero __SYCL2020_DEPRECATED("use 'ext_oneapi_level_zero' instead") = + ext_oneapi_level_zero, cuda = 3, all = 4, esimd_cpu = 5, diff --git a/sycl/include/CL/sycl/bit_cast.hpp b/sycl/include/CL/sycl/bit_cast.hpp index 2a042c92811b..6fe0b85dfe63 100644 --- a/sycl/include/CL/sycl/bit_cast.hpp +++ b/sycl/include/CL/sycl/bit_cast.hpp @@ -22,7 +22,6 @@ namespace detail { inline void memcpy(void *Dst, const void *Src, std::size_t Size); } -// sycl::bit_cast ( no longer sycl::detail::bit_cast ) template #if __cpp_lib_bit_cast || __has_builtin(__builtin_bit_cast) constexpr @@ -54,11 +53,11 @@ constexpr namespace detail { template +__SYCL2020_DEPRECATED("use 'sycl::bit_cast' instead") #if __cpp_lib_bit_cast || __has_builtin(__builtin_bit_cast) constexpr #endif - To - bit_cast(const From &from) noexcept { + To bit_cast(const From &from) noexcept { return sycl::bit_cast(from); } } // namespace detail diff --git a/sycl/include/CL/sycl/feature_test.hpp b/sycl/include/CL/sycl/feature_test.hpp index 4625cfa06fed..08f51244adc9 100644 --- a/sycl/include/CL/sycl/feature_test.hpp +++ b/sycl/include/CL/sycl/feature_test.hpp @@ -24,6 +24,10 @@ namespace sycl { #define SYCL_EXT_ONEAPI_MATRIX 2 #endif #define SYCL_EXT_INTEL_BF16_CONVERSION 1 +#define SYCL_EXT_ONEAPI_ENQUEUE_BARRIER 1 +#define SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY 1 +#define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 1 +#define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 1 } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index ac8e50450da9..c0663563f1e9 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -2286,17 +2286,32 @@ class __SYCL_EXPORT handler { /// Prevents any commands submitted afterward to this queue from executing /// until all commands previously submitted to this queue have entered the /// complete state. - void barrier() { + void ext_oneapi_barrier() { throwIfActionIsCreated(); setType(detail::CG::Barrier); } + /// Prevents any commands submitted afterward to this queue from executing + /// until all commands previously submitted to this queue have entered the + /// complete state. + __SYCL2020_DEPRECATED("use 'ext_oneapi_barrier' instead") + void barrier() { ext_oneapi_barrier(); } + + /// Prevents any commands submitted afterward to this queue from executing + /// until all events in WaitList have entered the complete state. If WaitList + /// is empty, then the barrier has no effect. + /// + /// \param WaitList is a vector of valid SYCL events that need to complete + /// before barrier command can be executed. + void ext_oneapi_barrier(const std::vector &WaitList); + /// Prevents any commands submitted afterward to this queue from executing /// until all events in WaitList have entered the complete state. If WaitList /// is empty, then the barrier has no effect. /// /// \param WaitList is a vector of valid SYCL events that need to complete /// before barrier command can be executed. + __SYCL2020_DEPRECATED("use 'ext_oneapi_barrier' instead") void barrier(const std::vector &WaitList); /// Copies data from one memory region to another, both pointed by diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index d50840e8f345..01f6cbc99015 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -312,8 +312,38 @@ class __SYCL_EXPORT queue { /// \param CodeLoc is the code location of the submit call (default argument) /// \return a SYCL event object, which corresponds to the queue the command /// group is being enqueued on. + event ext_oneapi_submit_barrier(_CODELOCONLYPARAM(&CodeLoc)) { + return submit( + [=](handler &CGH) { CGH.ext_oneapi_barrier(); } _CODELOCFW(CodeLoc)); + } + + /// Prevents any commands submitted afterward to this queue from executing + /// until all commands previously submitted to this queue have entered the + /// complete state. + /// + /// \param CodeLoc is the code location of the submit call (default argument) + /// \return a SYCL event object, which corresponds to the queue the command + /// group is being enqueued on. + __SYCL2020_DEPRECATED("use 'ext_oneapi_submit_barrier' instead") event submit_barrier(_CODELOCONLYPARAM(&CodeLoc)) { - return submit([=](handler &CGH) { CGH.barrier(); } _CODELOCFW(CodeLoc)); + _CODELOCARG(&CodeLoc); + return ext_oneapi_submit_barrier(CodeLoc); + } + + /// Prevents any commands submitted afterward to this queue from executing + /// until all events in WaitList have entered the complete state. If WaitList + /// is empty, then ext_oneapi_submit_barrier has no effect. + /// + /// \param WaitList is a vector of valid SYCL events that need to complete + /// before barrier command can be executed. + /// \param CodeLoc is the code location of the submit call (default argument) + /// \return a SYCL event object, which corresponds to the queue the command + /// group is being enqueued on. + event ext_oneapi_submit_barrier( + const std::vector &WaitList _CODELOCPARAM(&CodeLoc)) { + return submit([=](handler &CGH) { + CGH.ext_oneapi_barrier(WaitList); + } _CODELOCFW(CodeLoc)); } /// Prevents any commands submitted afterward to this queue from executing @@ -325,10 +355,11 @@ class __SYCL_EXPORT queue { /// \param CodeLoc is the code location of the submit call (default argument) /// \return a SYCL event object, which corresponds to the queue the command /// group is being enqueued on. + __SYCL2020_DEPRECATED("use 'ext_oneapi_submit_barrier' instead") event submit_barrier(const std::vector &WaitList _CODELOCPARAM(&CodeLoc)) { - return submit( - [=](handler &CGH) { CGH.barrier(WaitList); } _CODELOCFW(CodeLoc)); + _CODELOCARG(&CodeLoc); + return ext_oneapi_submit_barrier(WaitList, CodeLoc); } /// Performs a blocking wait for the completion of all enqueued tasks in the diff --git a/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp b/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp new file mode 100644 index 000000000000..d4f436307cb8 --- /dev/null +++ b/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp @@ -0,0 +1,235 @@ +//===------- online_compiler.hpp - Online source compilation service ------===// +// +// 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 // for __SYCL_INLINE_NAMESPACE +#include // for __SYCL_EXPORT +#include + +#include +#include + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace ext { +namespace intel { +namespace experimental { + +using byte = unsigned char; + +enum class compiled_code_format { + spir_v = 0 // the only format supported for now +}; + +class device_arch { +public: + static constexpr int any = 0; + + device_arch(int Val) : Val(Val) {} + + enum gpu { + gpu_any = 1, + gpu_gen9 = 2, + gpu_skl = gpu_gen9, + gpu_gen9_5 = 3, + gpu_kbl = gpu_gen9_5, + gpu_cfl = gpu_gen9_5, + gpu_gen11 = 4, + gpu_icl = gpu_gen11, + gpu_gen12 = 5 + }; + + enum cpu { + cpu_any = 1, + }; + + enum fpga { + fpga_any = 1, + }; + + operator int() { return Val; } + +private: + int Val; +}; + +/// Represents an error happend during online compilation. +class online_compile_error : public sycl::exception { +public: + online_compile_error() = default; + online_compile_error(const std::string &Msg) : sycl::exception(Msg) {} +}; + +/// Designates a source language for the online compiler. +enum class source_language { opencl_c = 0, cm = 1 }; + +/// Represents an online compiler for the language given as template +/// parameter. +template class online_compiler { +public: + /// Constructs online compiler which can target any device and produces + /// given compiled code format. Produces 64-bit device code. + /// The created compiler is "optimistic" - it assumes all applicable SYCL + /// device capabilities are supported by the target device(s). + online_compiler(compiled_code_format fmt = compiled_code_format::spir_v) + : OutputFormat(fmt), OutputFormatVersion({0, 0}), + DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), + Is64Bit(true), DeviceStepping("") {} + + /// Constructs online compiler which targets given architecture and produces + /// given compiled code format. Produces 64-bit device code. + /// Throws online_compile_error if values of constructor arguments are + /// contradictory or not supported - e.g. if the source language is not + /// supported for given device type. + online_compiler(sycl::info::device_type dev_type, device_arch arch, + compiled_code_format fmt = compiled_code_format::spir_v) + : OutputFormat(fmt), OutputFormatVersion({0, 0}), DeviceType(dev_type), + DeviceArch(arch), Is64Bit(true), DeviceStepping("") {} + + /// Constructs online compiler for the target specified by given SYCL device. + // TODO: the initial version generates the generic code (SKL now), need + // to do additional device::info calls to determine the device by it's + // features. + online_compiler(const sycl::device &) + : OutputFormat(compiled_code_format::spir_v), OutputFormatVersion({0, 0}), + DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), + Is64Bit(true), DeviceStepping("") {} + + /// Compiles given in-memory \c Lang source to a binary blob. Blob format, + /// other parameters are set in the constructor by the compilation target + /// specification parameters. + /// Specialization for each language will provide exact signatures, which + /// can be different for different languages. + /// Throws online_compile_error if compilation is not successful. + template + std::vector compile(const std::string &src, const Tys &... args); + + /// Sets the compiled code format of the compilation target and returns *this. + online_compiler &setOutputFormat(compiled_code_format fmt) { + OutputFormat = fmt; + return *this; + } + + /// Sets the compiled code format version of the compilation target and + /// returns *this. + online_compiler &setOutputFormatVersion(int major, int minor) { + OutputFormatVersion = {major, minor}; + return *this; + } + + /// Sets the device type of the compilation target and returns *this. + online_compiler &setTargetDeviceType(sycl::info::device_type type) { + DeviceType = type; + return *this; + } + + /// Sets the device architecture of the compilation target and returns *this. + online_compiler &setTargetDeviceArch(device_arch arch) { + DeviceArch = arch; + return *this; + } + + /// Makes the compilation target 32-bit and returns *this. + online_compiler &set32bitTarget() { + Is64Bit = false; + return *this; + }; + + /// Makes the compilation target 64-bit and returns *this. + online_compiler &set64bitTarget() { + Is64Bit = true; + return *this; + }; + + /// Sets implementation-defined target device stepping of the compilation + /// target and returns *this. + online_compiler &setTargetDeviceStepping(const std::string &id) { + DeviceStepping = id; + return *this; + } + +private: + /// Compiled code format. + compiled_code_format OutputFormat; + + /// Compiled code format version - a pair of "major" and "minor" components + std::pair OutputFormatVersion; + + /// Target device type + sycl::info::device_type DeviceType; + + /// Target device architecture + device_arch DeviceArch; + + /// Whether the target device architecture is 64-bit + bool Is64Bit; + + /// Target device stepping (implementation defined) + std::string DeviceStepping; + + /// Handles to helper functions used by the implementation. + void *CompileToSPIRVHandle = nullptr; + void *FreeSPIRVOutputsHandle = nullptr; +}; + +// Specializations of the online_compiler class and 'compile' function for +// particular languages and parameter types. + +/// Compiles the given OpenCL source. May throw \c online_compile_error. +/// @param src - contents of the source. +/// @param options - compilation options (implementation defined); standard +/// OpenCL JIT compiler options must be supported. +template <> +template <> +__SYCL_EXPORT std::vector +online_compiler::compile( + const std::string &src, const std::vector &options); + +/// Compiles the given OpenCL source. May throw \c online_compile_error. +/// @param src - contents of the source. +template <> +template <> +std::vector +online_compiler::compile(const std::string &src) { + return compile(src, std::vector{}); +} + +/// Compiles the given CM source \p src. +/// @param src - contents of the source. +/// @param options - compilation options (implementation defined). +template <> +template <> +__SYCL_EXPORT std::vector online_compiler::compile( + const std::string &src, const std::vector &options); + +/// Compiles the given CM source \p src. +template <> +template <> +std::vector +online_compiler::compile(const std::string &src) { + return compile(src, std::vector{}); +} + +} // namespace experimental +} // namespace intel +} // namespace ext + +namespace ext { +namespace __SYCL2020_DEPRECATED( + "use 'ext::intel::experimental' instead") intel { + using namespace ext::intel::experimental; +} // namespace intel +} // namespace ext + +namespace __SYCL2020_DEPRECATED( + "use 'ext::intel::experimental' instead") INTEL { + using namespace ext::intel::experimental; +} // namespace INTEL +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/sycl/ext/intel/online_compiler.hpp b/sycl/include/sycl/ext/intel/online_compiler.hpp index ef104f3dd9e9..d7212bb1b26f 100644 --- a/sycl/include/sycl/ext/intel/online_compiler.hpp +++ b/sycl/include/sycl/ext/intel/online_compiler.hpp @@ -8,358 +8,10 @@ #pragma once -#include // for __SYCL_INLINE_NAMESPACE -#include // for __SYCL_EXPORT -#include +#include -#include -#include +__SYCL_WARNING( + "sycl/ext/intel/online_compiler.hpp usage is deprecated, include " + "sycl/ext/intel/experimental/online_compiler.hpp instead") -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -namespace ext { -namespace intel { - -using byte = unsigned char; - -enum class compiled_code_format { - spir_v = 0 // the only format supported for now -}; - -class device_arch { -public: - static constexpr int any = 0; - - device_arch(int Val) : Val(Val) {} - - enum gpu { - gpu_any = 1, - gpu_gen9 = 2, - gpu_skl = gpu_gen9, - gpu_gen9_5 = 3, - gpu_kbl = gpu_gen9_5, - gpu_cfl = gpu_gen9_5, - gpu_gen11 = 4, - gpu_icl = gpu_gen11, - gpu_gen12 = 5 - }; - - enum cpu { - cpu_any = 1, - }; - - enum fpga { - fpga_any = 1, - }; - - operator int() { return Val; } - -private: - int Val; -}; - -/// Represents an error happend during online compilation. -class online_compile_error : public sycl::exception { -public: - online_compile_error() = default; - online_compile_error(const std::string &Msg) : sycl::exception(Msg) {} -}; - -/// Designates a source language for the online compiler. -enum class source_language { opencl_c = 0, cm = 1 }; - -/// Represents an online compiler for the language given as template -/// parameter. -template class online_compiler { -public: - /// Constructs online compiler which can target any device and produces - /// given compiled code format. Produces 64-bit device code. - /// The created compiler is "optimistic" - it assumes all applicable SYCL - /// device capabilities are supported by the target device(s). - online_compiler(compiled_code_format fmt = compiled_code_format::spir_v) - : OutputFormat(fmt), OutputFormatVersion({0, 0}), - DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), - Is64Bit(true), DeviceStepping("") {} - - /// Constructs online compiler which targets given architecture and produces - /// given compiled code format. Produces 64-bit device code. - /// Throws online_compile_error if values of constructor arguments are - /// contradictory or not supported - e.g. if the source language is not - /// supported for given device type. - online_compiler(sycl::info::device_type dev_type, device_arch arch, - compiled_code_format fmt = compiled_code_format::spir_v) - : OutputFormat(fmt), OutputFormatVersion({0, 0}), DeviceType(dev_type), - DeviceArch(arch), Is64Bit(true), DeviceStepping("") {} - - /// Constructs online compiler for the target specified by given SYCL device. - // TODO: the initial version generates the generic code (SKL now), need - // to do additional device::info calls to determine the device by it's - // features. - online_compiler(const sycl::device &) - : OutputFormat(compiled_code_format::spir_v), OutputFormatVersion({0, 0}), - DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), - Is64Bit(true), DeviceStepping("") {} - - /// Compiles given in-memory \c Lang source to a binary blob. Blob format, - /// other parameters are set in the constructor by the compilation target - /// specification parameters. - /// Specialization for each language will provide exact signatures, which - /// can be different for different languages. - /// Throws online_compile_error if compilation is not successful. - template - std::vector compile(const std::string &src, const Tys &... args); - - /// Sets the compiled code format of the compilation target and returns *this. - online_compiler &setOutputFormat(compiled_code_format fmt) { - OutputFormat = fmt; - return *this; - } - - /// Sets the compiled code format version of the compilation target and - /// returns *this. - online_compiler &setOutputFormatVersion(int major, int minor) { - OutputFormatVersion = {major, minor}; - return *this; - } - - /// Sets the device type of the compilation target and returns *this. - online_compiler &setTargetDeviceType(sycl::info::device_type type) { - DeviceType = type; - return *this; - } - - /// Sets the device architecture of the compilation target and returns *this. - online_compiler &setTargetDeviceArch(device_arch arch) { - DeviceArch = arch; - return *this; - } - - /// Makes the compilation target 32-bit and returns *this. - online_compiler &set32bitTarget() { - Is64Bit = false; - return *this; - }; - - /// Makes the compilation target 64-bit and returns *this. - online_compiler &set64bitTarget() { - Is64Bit = true; - return *this; - }; - - /// Sets implementation-defined target device stepping of the compilation - /// target and returns *this. - online_compiler &setTargetDeviceStepping(const std::string &id) { - DeviceStepping = id; - return *this; - } - -private: - /// Compiled code format. - compiled_code_format OutputFormat; - - /// Compiled code format version - a pair of "major" and "minor" components - std::pair OutputFormatVersion; - - /// Target device type - sycl::info::device_type DeviceType; - - /// Target device architecture - device_arch DeviceArch; - - /// Whether the target device architecture is 64-bit - bool Is64Bit; - - /// Target device stepping (implementation defined) - std::string DeviceStepping; - - /// Handles to helper functions used by the implementation. - void *CompileToSPIRVHandle = nullptr; - void *FreeSPIRVOutputsHandle = nullptr; -}; - -// Specializations of the online_compiler class and 'compile' function for -// particular languages and parameter types. - -/// Compiles the given OpenCL source. May throw \c online_compile_error. -/// @param src - contents of the source. -/// @param options - compilation options (implementation defined); standard -/// OpenCL JIT compiler options must be supported. -template <> -template <> -__SYCL_EXPORT std::vector -online_compiler::compile( - const std::string &src, const std::vector &options); - -/// Compiles the given OpenCL source. May throw \c online_compile_error. -/// @param src - contents of the source. -template <> -template <> -std::vector -online_compiler::compile(const std::string &src) { - return compile(src, std::vector{}); -} - -/// Compiles the given CM source \p src. -/// @param src - contents of the source. -/// @param options - compilation options (implementation defined). -template <> -template <> -__SYCL_EXPORT std::vector online_compiler::compile( - const std::string &src, const std::vector &options); - -/// Compiles the given CM source \p src. -template <> -template <> -std::vector -online_compiler::compile(const std::string &src) { - return compile(src, std::vector{}); -} - -} // namespace intel -} // namespace ext - -namespace __SYCL2020_DEPRECATED("use 'ext::intel' instead") INTEL { - using namespace ext::intel; - - /// Designates a source language for the online compiler. - enum class source_language { opencl_c = 0, cm = 1 }; - - template class online_compiler { - public: - /// Constructs online compiler which can target any device and produces - /// given compiled code format. Produces 64-bit device code. - /// The created compiler is "optimistic" - it assumes all applicable SYCL - /// device capabilities are supported by the target device(s). - online_compiler(ext::intel::compiled_code_format fmt = - ext::intel::compiled_code_format::spir_v) { - MOnlineCompiler = - ext::intel::online_compiler<(ext::intel::source_language)Lang>(fmt); - } - - /// Constructs online compiler which targets given architecture and produces - /// given compiled code format. Produces 64-bit device code. - /// Throws online_compile_error if values of constructor arguments are - /// contradictory or not supported - e.g. if the source language is not - /// supported for given device type. - online_compiler(sycl::info::device_type dev_type, - ext::intel::device_arch arch, - ext::intel::compiled_code_format fmt = - ext::intel::compiled_code_format::spir_v) { - MOnlineCompiler = - ext::intel::online_compiler<(ext::intel::source_language)Lang>( - dev_type, arch, fmt); - } - - /// Constructs online compiler for the target specified by given SYCL - /// device. - // TODO: the initial version generates the generic code (SKL now), need - // to do additional device::info calls to determine the device by it's - // features. - online_compiler(const sycl::device &device) { - MOnlineCompiler = - ext::intel::online_compiler<(ext::intel::source_language)Lang>( - device); - } - - /// Compiles given in-memory \c Lang source to a binary blob. Blob format, - /// other parameters are set in the constructor by the compilation target - /// specification parameters. - /// Specialization for each language will provide exact signatures, which - /// can be different for different languages. - /// Throws online_compile_error if compilation is not successful. - template - std::vector compile(const std::string &src, const Tys &... args); - - /// Sets the compiled code format of the compilation target and returns - /// *this. - online_compiler & - setOutputFormat(ext::intel::compiled_code_format fmt) { - MOnlineCompiler.setOutputFormat(fmt); - return *this; - } - - /// Sets the compiled code format version of the compilation target and - /// returns *this. - online_compiler &setOutputFormatVersion(int major, int minor) { - MOnlineCompiler.setOutputFormatVersion(major, minor); - return *this; - } - - /// Sets the device type of the compilation target and returns *this. - online_compiler &setTargetDeviceType(sycl::info::device_type type) { - MOnlineCompiler.setTargetDeviceType(type); - return *this; - } - - /// Sets the device architecture of the compilation target and returns - /// *this. - online_compiler &setTargetDeviceArch(device_arch arch) { - MOnlineCompiler.setTargetDeviceArch(arch); - return *this; - } - - /// Makes the compilation target 32-bit and returns *this. - online_compiler &set32bitTarget() { - MOnlineCompiler.set32bitTarget(); - return *this; - }; - - /// Makes the compilation target 64-bit and returns *this. - online_compiler &set64bitTarget() { - MOnlineCompiler.set64bitTarget(); - return *this; - }; - - /// Sets implementation-defined target device stepping of the compilation - /// target and returns *this. - online_compiler &setTargetDeviceStepping(const std::string &id) { - MOnlineCompiler.setTargetDeviceStepping(id); - return *this; - } - - private: - ext::intel::online_compiler<(ext::intel::source_language)Lang> - MOnlineCompiler; - }; - - // Specializations of the online_compiler class and 'compile' function for - // particular languages and parameter types. - - /// Compiles the given OpenCL source. May throw \c online_compile_error. - /// @param src - contents of the source. - /// @param options - compilation options (implementation defined); standard - /// OpenCL JIT compiler options must be supported. - template <> - template <> - __SYCL_EXPORT std::vector - online_compiler::compile( - const std::string &src, const std::vector &options); - - /// Compiles the given OpenCL source. May throw \c online_compile_error. - /// @param src - contents of the source. - template <> - template <> - std::vector online_compiler::compile( - const std::string &src) { - return MOnlineCompiler.compile(src); - } - - /// Compiles the given CM source \p src. - /// @param src - contents of the source. - /// @param options - compilation options (implementation defined). - template <> - template <> - __SYCL_EXPORT std::vector online_compiler::compile( - const std::string &src, const std::vector &options); - - /// Compiles the given CM source \p src. - template <> - template <> - std::vector online_compiler::compile( - const std::string &src) { - return MOnlineCompiler.compile(src); - } - -} // namespace INTEL -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) +#include diff --git a/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp new file mode 100644 index 000000000000..cb41cc34447f --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp @@ -0,0 +1,205 @@ +//==--------- level_zero.hpp - SYCL Level-Zero backend ---------------------==// +// +// 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 +// This header should be included by users. +//#include + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { + +template <> struct interop { + using type = ze_driver_handle_t; +}; + +template <> struct interop { + using type = ze_device_handle_t; +}; + +template <> struct interop { + using type = ze_context_handle_t; +}; + +template <> struct interop { + using type = ze_command_queue_handle_t; +}; + +template <> struct interop { + using type = ze_event_handle_t; +}; + +template <> struct interop { + using type = ze_module_handle_t; +}; + +template +struct interop> { + using type = char *; +}; + +template +struct interop> { + using type = char *; +}; + +template +struct interop> { + using type = ze_image_handle_t; +}; + +namespace ext { +namespace oneapi { +namespace level_zero { +// Since Level-Zero is not doing any reference counting itself, we have to +// be explicit about the ownership of the native handles used in the +// interop functions below. +// +enum class ownership { transfer, keep }; +} // namespace level_zero +} // namespace oneapi +} // namespace ext + +namespace detail { + +template <> struct BackendInput { + using type = struct { + interop::type NativeHandle; + std::vector DeviceList; + ext::oneapi::level_zero::ownership Ownership; + }; +}; + +template <> struct BackendReturn { + using type = ze_kernel_handle_t; +}; + +template <> struct InteropFeatureSupportMap { + static constexpr bool MakePlatform = true; + static constexpr bool MakeDevice = true; + static constexpr bool MakeContext = true; + static constexpr bool MakeQueue = false; + static constexpr bool MakeEvent = true; + static constexpr bool MakeBuffer = false; + static constexpr bool MakeKernel = false; + static constexpr bool MakeKernelBundle = false; +}; +} // namespace detail + +namespace ext { +namespace oneapi { +namespace level_zero { +// Implementation of various "make" functions resides in libsycl.so and thus +// their interface needs to be backend agnostic. +// TODO: remove/merge with similar functions in sycl::detail +__SYCL_EXPORT platform make_platform(pi_native_handle NativeHandle); +__SYCL_EXPORT device make_device(const platform &Platform, + pi_native_handle NativeHandle); +__SYCL_EXPORT context make_context(const std::vector &DeviceList, + pi_native_handle NativeHandle, + bool keep_ownership = false); +__SYCL_EXPORT program make_program(const context &Context, + pi_native_handle NativeHandle); +__SYCL_EXPORT queue make_queue(const context &Context, + pi_native_handle InteropHandle, + bool keep_ownership = false); +__SYCL_EXPORT event make_event(const context &Context, + pi_native_handle InteropHandle, + bool keep_ownership = false); + +// Construction of SYCL platform. +template ::value> * = nullptr> +__SYCL_DEPRECATED("Use SYCL-2020 sycl::make_platform free function") +T make(typename interop::type Interop) { + return make_platform(reinterpret_cast(Interop)); +} + +// Construction of SYCL device. +template ::value> * = nullptr> +__SYCL_DEPRECATED("Use SYCL-2020 sycl::make_device free function") +T make(const platform &Platform, + typename interop::type Interop) { + return make_device(Platform, reinterpret_cast(Interop)); +} + +/// Construction of SYCL context. +/// \param DeviceList is a vector of devices which must be encapsulated by +/// created SYCL context. Provided devices and native context handle must +/// be associated with the same platform. +/// \param Interop is a Level Zero native context handle. +/// \param Ownership (optional) specifies who will assume ownership of the +/// native context handle. Default is that SYCL RT does, so it destroys +/// the native handle when the created SYCL object goes out of life. +/// +template ::value>::type * = nullptr> +__SYCL_DEPRECATED("Use SYCL-2020 sycl::make_context free function") +T make(const std::vector &DeviceList, + typename interop::type Interop, + ownership Ownership = ownership::transfer) { + return make_context(DeviceList, detail::pi::cast(Interop), + Ownership == ownership::keep); +} + +// Construction of SYCL program. +template ::value> * = nullptr> +T make(const context &Context, + typename interop::type Interop) { + return make_program(Context, reinterpret_cast(Interop)); +} + +// Construction of SYCL queue. +template ::value> * = nullptr> +T make(const context &Context, + typename interop::type Interop, + ownership Ownership = ownership::transfer) { + return make_queue(Context, reinterpret_cast(Interop), + Ownership == ownership::keep); +} + +// Construction of SYCL event. +template ::value> * = nullptr> +T make(const context &Context, + typename interop::type Interop, + ownership Ownership = ownership::transfer) { + return make_event(Context, reinterpret_cast(Interop), + Ownership == ownership::keep); +} +} // namespace level_zero +} // namespace oneapi +} // namespace ext + +// Specialization of sycl::make_context for Level-Zero backend. +template <> +context make_context( + const backend_input_t &BackendObject, + const async_handler &Handler) { + return ext::oneapi::level_zero::make_context( + BackendObject.DeviceList, + detail::pi::cast(BackendObject.NativeHandle), + BackendObject.Ownership == ext::oneapi::level_zero::ownership::keep); +} + +namespace __SYCL2020_DEPRECATED("use 'ext::oneapi::level_zero' instead") + level_zero { + using namespace ext::oneapi::level_zero; +} +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/sycl/ext/oneapi/group_local_memory.hpp b/sycl/include/sycl/ext/oneapi/group_local_memory.hpp index 1f37c29c5bec..89aba9bf4a43 100644 --- a/sycl/include/sycl/ext/oneapi/group_local_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/group_local_memory.hpp @@ -35,7 +35,7 @@ std::enable_if_t::value && return reinterpret_cast<__attribute__((opencl_local)) T *>(AllocatedMem); #else throw feature_not_supported( - "SYCL_INTEL_local_memory extension is not supported on host device", + "SYCL_EXT_ONEAPI_LOCAL_MEMORY extension is not supported on host device", PI_INVALID_OPERATION); #endif } @@ -60,7 +60,7 @@ std::enable_if_t::value && // Silence unused variable warning [&args...] {}(); throw feature_not_supported( - "SYCL_INTEL_local_memory extension is not supported on host device", + "SYCL_EXT_ONEAPI_LOCAL_MEMORY extension is not supported on host device", PI_INVALID_OPERATION); #endif } diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index f3eecbd2de11..f8744de9e17a 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -52,8 +52,12 @@ function(add_sycl_rt_library LIB_NAME) target_compile_options(${LIB_OBJ_NAME} PUBLIC -fvisibility=hidden -fvisibility-inlines-hidden) set(linker_script "${CMAKE_CURRENT_SOURCE_DIR}/ld-version-script.txt") + set(abi_linker_script "${CMAKE_CURRENT_SOURCE_DIR}/abi_replacements_linux.txt") target_link_libraries( - ${LIB_NAME} PRIVATE "-Wl,--version-script=${linker_script}") + ${LIB_NAME} PRIVATE "-Wl,${abi_linker_script}") + set_target_properties(${LIB_NAME} PROPERTIES LINK_DEPENDS ${abi_linker_script}) + target_link_libraries( + ${LIB_NAME} PRIVATE "-Wl,--version-script=${linker_script}") set_target_properties(${LIB_NAME} PROPERTIES LINK_DEPENDS ${linker_script}) if (SYCL_ENABLE_XPTI_TRACING) target_link_libraries(${LIB_NAME} PRIVATE dl) diff --git a/sycl/source/abi_replacements_linux.txt b/sycl/source/abi_replacements_linux.txt new file mode 100644 index 000000000000..edf56b2ebe73 --- /dev/null +++ b/sycl/source/abi_replacements_linux.txt @@ -0,0 +1,12 @@ +_ZN2cl4sycl10level_zero10make_queueERKNS0_7contextEmb = _ZN2cl4sycl3ext6oneapi10level_zero10make_queueERKNS0_7contextEmb; +_ZN2cl4sycl10level_zero10make_queueERKNS0_7contextEm = _ZN2cl4sycl3ext6oneapi10level_zero10make_queueERKNS0_7contextEm; +_ZN2cl4sycl10level_zero12make_programERKNS0_7contextEm = _ZN2cl4sycl3ext6oneapi10level_zero12make_programERKNS0_7contextEm; +_ZN2cl4sycl10level_zero12make_contextERKSt6vectorINS0_6deviceESaIS3_EEmb = _ZN2cl4sycl3ext6oneapi10level_zero12make_contextERKSt6vectorINS0_6deviceESaIS5_EEmb; +_ZN2cl4sycl10level_zero12make_contextERKSt6vectorINS0_6deviceESaIS3_EEm = _ZN2cl4sycl3ext6oneapi10level_zero12make_contextERKSt6vectorINS0_6deviceESaIS5_EEm; +_ZN2cl4sycl10level_zero11make_deviceERKNS0_8platformEm = _ZN2cl4sycl3ext6oneapi10level_zero11make_deviceERKNS0_8platformEm; +_ZN2cl4sycl10level_zero10make_eventERKNS0_7contextEmb = _ZN2cl4sycl3ext6oneapi10level_zero10make_eventERKNS0_7contextEmb; +_ZN2cl4sycl10level_zero13make_platformEm = _ZN2cl4sycl3ext6oneapi10level_zero13make_platformEm; +_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISC_EEEEES6_IhSaIhEERKSC_DpRKT_ = _ZN2cl4sycl3ext5intel12experimental15online_compilerILNS3_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_; +_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISC_EEEEES6_IhSaIhEERKSC_DpRKT_ = _ZN2cl4sycl3ext5intel12experimental15online_compilerILNS3_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_; +_ZN2cl4sycl3ext5intel15online_compilerILNS2_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISD_EEEEES7_IhSaIhEERKSD_DpRKT_ = _ZN2cl4sycl3ext5intel12experimental15online_compilerILNS3_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_; +_ZN2cl4sycl3ext5intel15online_compilerILNS2_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISD_EEEEES7_IhSaIhEERKSD_DpRKT_ = _ZN2cl4sycl3ext5intel12experimental15online_compilerILNS3_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_; \ No newline at end of file diff --git a/sycl/source/abi_replacements_windows.cpp b/sycl/source/abi_replacements_windows.cpp index 35f782a38064..609954225e91 100644 --- a/sycl/source/abi_replacements_windows.cpp +++ b/sycl/source/abi_replacements_windows.cpp @@ -7,3 +7,45 @@ #pragma comment( \ linker, \ "/export:?accessGlobalOffset@stream_impl@detail@sycl@cl@@QEAA?AV?$accessor@I$00$0EAF@$0HNO@$0A@V?$accessor_property_list@$$V@ONEAPI@sycl@cl@@@34@AEAVhandler@34@@Z=?accessGlobalOffset@stream_impl@detail@sycl@cl@@QEAA?AV?$accessor@I$00$0EAF@$0HNO@$0A@V?$accessor_property_list@$$V@oneapi@ext@sycl@cl@@@34@AEAVhandler@34@@Z") +#pragma comment( \ + linker, \ + "/export:?make_context@level_zero@sycl@cl@@YA?AVcontext@23@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@_K_N@Z=?make_context@level_zero@oneapi@ext@sycl@cl@@YA?AVcontext@45@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@_K_N@Z") +#pragma comment( \ + linker, \ + "/export:?make_event@level_zero@sycl@cl@@YA?AVevent@23@AEBVcontext@23@_K_N@Z=?make_event@level_zero@oneapi@ext@sycl@cl@@YA?AVevent@45@AEBVcontext@45@_K_N@Z") +#pragma comment( \ + linker, \ + "/export:?make_queue@level_zero@sycl@cl@@YA?AVqueue@23@AEBVcontext@23@_K@Z=?make_queue@level_zero@oneapi@ext@sycl@cl@@YA?AVqueue@45@AEBVcontext@45@_K@Z") +#pragma comment( \ + linker, \ + "/export:??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$00@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@6@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@6@@Z=??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$00@experimental@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@7@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@7@@Z") +#pragma comment( \ + linker, \ + "/export:?make_device@level_zero@sycl@cl@@YA?AVdevice@23@AEBVplatform@23@_K@Z=?make_device@level_zero@oneapi@ext@sycl@cl@@YA?AVdevice@45@AEBVplatform@45@_K@Z") +#pragma comment( \ + linker, \ + "/export:?make_platform@level_zero@sycl@cl@@YA?AVplatform@23@_K@Z=?make_platform@level_zero@oneapi@ext@sycl@cl@@YA?AVplatform@45@_K@Z") +#pragma comment( \ + linker, \ + "/export:??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$0A@@INTEL@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@5@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@5@@Z=??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$0A@@experimental@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@7@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@7@@Z") +#pragma comment( \ + linker, \ + "/export:?submit_barrier@queue@sycl@cl@@QEAA?AVevent@23@AEBV?$vector@Vevent@sycl@cl@@V?$allocator@Vevent@sycl@cl@@@std@@@std@@AEBUcode_location@detail@23@@Z=?ext_oneapi_submit_barrier@queue@sycl@cl@@QEAA?AVevent@23@AEBV?$vector@Vevent@sycl@cl@@V?$allocator@Vevent@sycl@cl@@@std@@@std@@AEBUcode_location@detail@23@@Z") +#pragma comment( \ + linker, \ + "/export:?make_queue@level_zero@sycl@cl@@YA?AVqueue@23@AEBVcontext@23@_K_N@Z=?make_queue@level_zero@oneapi@ext@sycl@cl@@YA?AVqueue@45@AEBVcontext@45@_K_N@Z") +#pragma comment( \ + linker, \ + "/export:?make_program@level_zero@sycl@cl@@YA?AVprogram@23@AEBVcontext@23@_K@Z=?make_program@level_zero@oneapi@ext@sycl@cl@@YA?AVprogram@45@AEBVcontext@45@_K@Z") +#pragma comment( \ + linker, \ + "/export:??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$0A@@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@6@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@6@@Z=??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$0A@@experimental@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@7@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@7@@Z") +#pragma comment( \ + linker, \ + "/export:?make_context@level_zero@sycl@cl@@YA?AVcontext@23@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@_K@Z=?make_context@level_zero@oneapi@ext@sycl@cl@@YA?AVcontext@45@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@_K@Z") +#pragma comment( \ + linker, \ + "/export:??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$00@INTEL@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@5@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@5@@Z=??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$00@experimental@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@7@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@7@@Z") +#pragma comment( \ + linker, \ + "/export:?submit_barrier@queue@sycl@cl@@QEAA?AVevent@23@AEBUcode_location@detail@23@@Z=?ext_oneapi_submit_barrier@queue@sycl@cl@@QEAA?AVevent@23@AEBUcode_location@detail@23@@Z") diff --git a/sycl/source/backend/level_zero.cpp b/sycl/source/backend/level_zero.cpp index 6463cc6e20bf..fe10236415ec 100644 --- a/sycl/source/backend/level_zero.cpp +++ b/sycl/source/backend/level_zero.cpp @@ -15,6 +15,8 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { +namespace ext { +namespace oneapi { namespace level_zero { using namespace detail; @@ -103,5 +105,7 @@ __SYCL_EXPORT event make_event(const context &Context, } } // namespace level_zero +} // namespace oneapi +} // namespace ext } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/online_compiler/online_compiler.cpp b/sycl/source/detail/online_compiler/online_compiler.cpp index 2c304f602af8..96a1641750bc 100644 --- a/sycl/source/detail/online_compiler/online_compiler.cpp +++ b/sycl/source/detail/online_compiler/online_compiler.cpp @@ -8,7 +8,7 @@ #include #include -#include +#include #include @@ -18,6 +18,7 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace ext { namespace intel { +namespace experimental { namespace detail { static std::vector @@ -229,27 +230,20 @@ __SYCL_EXPORT std::vector online_compiler::compile( DeviceStepping, CompileToSPIRVHandle, FreeSPIRVOutputsHandle, CMUserArgs); } - +} // namespace experimental } // namespace intel } // namespace ext -namespace __SYCL2020_DEPRECATED("use 'ext::intel' instead") INTEL { - using namespace ext::intel; - - template <> - template <> - __SYCL_EXPORT std::vector - online_compiler::compile( - const std::string &Source, const std::vector &UserArgs) { - return MOnlineCompiler.compile(Source, UserArgs); - } +namespace ext { +namespace __SYCL2020_DEPRECATED( + "use 'ext::intel::experimental' instead") intel { + using namespace ext::intel::experimental; +} // namespace intel +} // namespace ext - template <> - template <> - __SYCL_EXPORT std::vector online_compiler::compile( - const std::string &Source, const std::vector &UserArgs) { - return MOnlineCompiler.compile(Source, UserArgs); - } +namespace __SYCL2020_DEPRECATED( + "use 'ext::intel::experimental' instead") INTEL { + using namespace ext::intel::experimental; } // namespace INTEL } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 202c4ec7ea8d..645800a787af 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -490,7 +490,7 @@ std::string handler::getKernelName() { return MKernel->get_info(); } -void handler::barrier(const std::vector &WaitList) { +void handler::ext_oneapi_barrier(const std::vector &WaitList) { throwIfActionIsCreated(); MCGType = detail::CG::BarrierWaitlist; MEventsWaitWithBarrier.resize(WaitList.size()); @@ -499,6 +499,11 @@ void handler::barrier(const std::vector &WaitList) { [](const event &Event) { return detail::getSyclObjImpl(Event); }); } +__SYCL2020_DEPRECATED("use 'ext_oneapi_barrier' instead") +void handler::barrier(const std::vector &WaitList) { + handler::ext_oneapi_barrier(WaitList); +} + using namespace sycl::detail; bool handler::DisableRangeRounding() { return SYCLConfig::get(); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 48df1d6dd9ff..d42ded40e098 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3634,8 +3634,18 @@ _ZN2cl4sycl20aligned_alloc_sharedEmmRKNS0_5queueE _ZN2cl4sycl20aligned_alloc_sharedEmmRKNS0_5queueERKNS0_13property_listE _ZN2cl4sycl20aligned_alloc_sharedEmmRKNS0_6deviceERKNS0_7contextE _ZN2cl4sycl20aligned_alloc_sharedEmmRKNS0_6deviceERKNS0_7contextERKNS0_13property_listE +_ZN2cl4sycl3ext5intel12experimental15online_compilerILNS3_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_ +_ZN2cl4sycl3ext5intel12experimental15online_compilerILNS3_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_ _ZN2cl4sycl3ext5intel15online_compilerILNS2_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISD_EEEEES7_IhSaIhEERKSD_DpRKT_ _ZN2cl4sycl3ext5intel15online_compilerILNS2_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISD_EEEEES7_IhSaIhEERKSD_DpRKT_ +_ZN2cl4sycl3ext6oneapi10level_zero10make_eventERKNS0_7contextEmb +_ZN2cl4sycl3ext6oneapi10level_zero10make_queueERKNS0_7contextEm +_ZN2cl4sycl3ext6oneapi10level_zero10make_queueERKNS0_7contextEmb +_ZN2cl4sycl3ext6oneapi10level_zero11make_deviceERKNS0_8platformEm +_ZN2cl4sycl3ext6oneapi10level_zero12make_contextERKSt6vectorINS0_6deviceESaIS5_EEm +_ZN2cl4sycl3ext6oneapi10level_zero12make_contextERKSt6vectorINS0_6deviceESaIS5_EEmb +_ZN2cl4sycl3ext6oneapi10level_zero12make_programERKNS0_7contextEm +_ZN2cl4sycl3ext6oneapi10level_zero13make_platformEm _ZN2cl4sycl3ext6oneapi15filter_selectorC1ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZN2cl4sycl3ext6oneapi15filter_selectorC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZN2cl4sycl3ext6oneapi6detail16reduGetMaxWGSizeESt10shared_ptrINS0_6detail10queue_implEEm @@ -3904,6 +3914,7 @@ _ZN2cl4sycl7handler10mem_adviseEPKvmi _ZN2cl4sycl7handler10processArgEPvRKNS0_6detail19kernel_param_kind_tEimRmb _ZN2cl4sycl7handler10processArgEPvRKNS0_6detail19kernel_param_kind_tEimRmbb _ZN2cl4sycl7handler13getKernelNameB5cxx11Ev +_ZN2cl4sycl7handler18ext_oneapi_barrierERKSt6vectorINS0_5eventESaIS3_EE _ZN2cl4sycl7handler18RangeRoundingTraceEv _ZN2cl4sycl7handler18extractArgsAndReqsEv _ZN2cl4sycl7handler20DisableRangeRoundingEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index dc854bf1395e..a091826e8dfe 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -8,8 +8,10 @@ # UNSUPPORTED: libcxx ??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$00@INTEL@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@5@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@5@@Z +??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$00@experimental@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@7@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@7@@Z ??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$00@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@6@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@6@@Z ??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$0A@@INTEL@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@5@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@5@@Z +??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$0A@@experimental@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@7@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@7@@Z ??$compile@V?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@@?$online_compiler@$0A@@intel@ext@sycl@cl@@QEAA?AV?$vector@EV?$allocator@E@std@@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@6@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@6@@Z ??$create_sub_devices@$0BAIG@@device@sycl@cl@@QEBA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@_K@Z ??$create_sub_devices@$0BAIH@@device@sycl@cl@@QEBA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@AEBV?$vector@_KV?$allocator@_K@std@@@4@@Z @@ -1738,6 +1740,10 @@ ?expm1@__host_std@cl@@YA?AVhalf@half_impl@detail@sycl@2@V34562@@Z ?expm1@__host_std@cl@@YAMM@Z ?expm1@__host_std@cl@@YANN@Z +?ext_oneapi_barrier@handler@sycl@cl@@QEAAXAEBV?$vector@Vevent@sycl@cl@@V?$allocator@Vevent@sycl@cl@@@std@@@std@@@Z +?ext_oneapi_barrier@handler@sycl@cl@@QEAAXXZ +?ext_oneapi_submit_barrier@queue@sycl@cl@@QEAA?AVevent@23@AEBUcode_location@detail@23@@Z +?ext_oneapi_submit_barrier@queue@sycl@cl@@QEAA?AVevent@23@AEBV?$vector@Vevent@sycl@cl@@V?$allocator@Vevent@sycl@cl@@@std@@@std@@AEBUcode_location@detail@23@@Z ?extractArgsAndReqs@handler@sycl@cl@@AEAAXXZ ?extractArgsAndReqsFromLambda@handler@sycl@cl@@AEAAXPEAD_KPEBUkernel_param_desc_t@detail@23@@Z ?extractArgsAndReqsFromLambda@handler@sycl@cl@@AEAAXPEAD_KPEBUkernel_param_desc_t@detail@23@_N@Z @@ -2545,25 +2551,33 @@ ?mad@__host_std@cl@@YANNNN@Z ?makeDir@OSUtil@detail@sycl@cl@@SAHPEBD@Z ?make_context@detail@sycl@cl@@YA?AVcontext@23@_KAEBV?$function@$$A6AXVexception_list@sycl@cl@@@Z@std@@W4backend@23@@Z +?make_context@level_zero@oneapi@ext@sycl@cl@@YA?AVcontext@45@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@_K@Z +?make_context@level_zero@oneapi@ext@sycl@cl@@YA?AVcontext@45@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@_K_N@Z ?make_context@level_zero@sycl@cl@@YA?AVcontext@23@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@_K@Z ?make_context@level_zero@sycl@cl@@YA?AVcontext@23@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@_K_N@Z ?make_context@opencl@sycl@cl@@YA?AVcontext@23@_K@Z ?make_device@detail@sycl@cl@@YA?AVdevice@23@_KW4backend@23@@Z +?make_device@level_zero@oneapi@ext@sycl@cl@@YA?AVdevice@45@AEBVplatform@45@_K@Z ?make_device@level_zero@sycl@cl@@YA?AVdevice@23@AEBVplatform@23@_K@Z ?make_device@opencl@sycl@cl@@YA?AVdevice@23@_K@Z ?make_error_code@sycl@cl@@YA?AVerror_code@std@@W4errc@12@@Z ?make_event@detail@sycl@cl@@YA?AVevent@23@_KAEBVcontext@23@W4backend@23@@Z ?make_event@detail@sycl@cl@@YA?AVevent@23@_KAEBVcontext@23@_NW4backend@23@@Z +?make_event@level_zero@oneapi@ext@sycl@cl@@YA?AVevent@45@AEBVcontext@45@_K_N@Z ?make_event@level_zero@sycl@cl@@YA?AVevent@23@AEBVcontext@23@_K_N@Z ?make_kernel@detail@sycl@cl@@YA?AVkernel@23@_KAEBVcontext@23@W4backend@23@@Z ?make_kernel_bundle@detail@sycl@cl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@_KAEBVcontext@23@W4bundle_state@23@W4backend@23@@Z ?make_platform@detail@sycl@cl@@YA?AVplatform@23@_KW4backend@23@@Z +?make_platform@level_zero@oneapi@ext@sycl@cl@@YA?AVplatform@45@_K@Z ?make_platform@level_zero@sycl@cl@@YA?AVplatform@23@_K@Z ?make_platform@opencl@sycl@cl@@YA?AVplatform@23@_K@Z +?make_program@level_zero@oneapi@ext@sycl@cl@@YA?AVprogram@45@AEBVcontext@45@_K@Z ?make_program@level_zero@sycl@cl@@YA?AVprogram@23@AEBVcontext@23@_K@Z ?make_program@opencl@sycl@cl@@YA?AVprogram@23@AEBVcontext@23@_K@Z ?make_queue@detail@sycl@cl@@YA?AVqueue@23@_KAEBVcontext@23@AEBV?$function@$$A6AXVexception_list@sycl@cl@@@Z@std@@W4backend@23@@Z ?make_queue@detail@sycl@cl@@YA?AVqueue@23@_KAEBVcontext@23@_NAEBV?$function@$$A6AXVexception_list@sycl@cl@@@Z@std@@W4backend@23@@Z +?make_queue@level_zero@oneapi@ext@sycl@cl@@YA?AVqueue@45@AEBVcontext@45@_K@Z +?make_queue@level_zero@oneapi@ext@sycl@cl@@YA?AVqueue@45@AEBVcontext@45@_K_N@Z ?make_queue@level_zero@sycl@cl@@YA?AVqueue@23@AEBVcontext@23@_K@Z ?make_queue@level_zero@sycl@cl@@YA?AVqueue@23@AEBVcontext@23@_K_N@Z ?make_queue@opencl@sycl@cl@@YA?AVqueue@23@AEBVcontext@23@_K@Z diff --git a/sycl/test/basic_tests/interop-level-zero-2020.cpp b/sycl/test/basic_tests/interop-level-zero-2020.cpp index eace5afa7d8c..4e40829cffca 100644 --- a/sycl/test/basic_tests/interop-level-zero-2020.cpp +++ b/sycl/test/basic_tests/interop-level-zero-2020.cpp @@ -41,9 +41,11 @@ int main() { // return_type is used when retrieving the backend specific native object from // a SYCL object. See the relevant backend specification for details. - backend_traits::return_type ZeDriver; - backend_traits::return_type ZeDevice; - backend_traits::return_type ZeContext; + backend_traits::return_type + ZeDriver; + backend_traits::return_type ZeDevice; + backend_traits::return_type + ZeContext; // 4.5.1.2 For each SYCL runtime class T which supports SYCL application // interoperability, a specialization of get_native must be defined, which @@ -52,20 +54,20 @@ int main() { // application interoperability. The lifetime of the object returned are // backend-defined and specified in the backend specification. - ZeDriver = get_native(Platform); - ZeDevice = get_native(Device); - ZeContext = get_native(Context); + ZeDriver = get_native(Platform); + ZeDevice = get_native(Device); + ZeContext = get_native(Context); // Check deprecated // expected-warning@+2 {{'get_native' is deprecated: Use SYCL-2020 sycl::get_native free function}} - // expected-warning@+1 {{'get_native' is deprecated: Use SYCL-2020 sycl::get_native free function}} - ZeDriver = Platform.get_native(); + // expected-warning@+1 {{'get_native' is deprecated: Use SYCL-2020 sycl::get_native free function}} + ZeDriver = Platform.get_native(); // expected-warning@+2 {{'get_native' is deprecated: Use SYCL-2020 sycl::get_native free function}} - // expected-warning@+1 {{'get_native' is deprecated: Use SYCL-2020 sycl::get_native free function}} - ZeDevice = Device.get_native(); + // expected-warning@+1 {{'get_native' is deprecated: Use SYCL-2020 sycl::get_native free function}} + ZeDevice = Device.get_native(); // expected-warning@+2 {{'get_native' is deprecated: Use SYCL-2020 sycl::get_native free function}} - // expected-warning@+1 {{'get_native' is deprecated: Use SYCL-2020 sycl::get_native free function}} - ZeContext = Context.get_native(); + // expected-warning@+1 {{'get_native' is deprecated: Use SYCL-2020 sycl::get_native free function}} + ZeContext = Context.get_native(); // 4.5.1.1 For each SYCL runtime class T which supports SYCL application // interoperability with the SYCL backend, a specialization of input_type must @@ -83,27 +85,31 @@ int main() { // behavior of these template functions is defined by the SYCL backend // specification document. - backend_input_t InteropPlatformInput{ZeDriver}; + backend_input_t + InteropPlatformInput{ZeDriver}; platform InteropPlatform = - make_platform(InteropPlatformInput); + make_platform(InteropPlatformInput); - backend_input_t InteropDeviceInput{ZeDevice}; - device InteropDevice = make_device(InteropDeviceInput); + backend_input_t InteropDeviceInput{ + ZeDevice}; + device InteropDevice = + make_device(InteropDeviceInput); - backend_input_t InteropContextInput{ + backend_input_t InteropContextInput{ ZeContext, std::vector(1, InteropDevice), - level_zero::ownership::keep}; + ext::oneapi::level_zero::ownership::keep}; context InteropContext = - make_context(InteropContextInput); + make_context(InteropContextInput); // Check deprecated // expected-warning@+1 {{'make' is deprecated: Use SYCL-2020 sycl::make_platform free function}} - auto P = level_zero::make(ZeDriver); + auto P = ext::oneapi::level_zero::make(ZeDriver); // expected-warning@+1 {{'make' is deprecated: Use SYCL-2020 sycl::make_device free function}} - auto D = level_zero::make(P, ZeDevice); + auto D = ext::oneapi::level_zero::make(P, ZeDevice); // expected-warning@+1 {{'make' is deprecated: Use SYCL-2020 sycl::make_context free function}} - auto C = level_zero::make(std::vector(1, D), ZeContext, - level_zero::ownership::keep); + auto C = ext::oneapi::level_zero::make( + std::vector(1, D), ZeContext, + ext::oneapi::level_zero::ownership::keep); return 0; } diff --git a/sycl/test/warnings/sycl_2020_deprecations.cpp b/sycl/test/warnings/sycl_2020_deprecations.cpp index e45f5289b3c6..30fd07e480ad 100644 --- a/sycl/test/warnings/sycl_2020_deprecations.cpp +++ b/sycl/test/warnings/sycl_2020_deprecations.cpp @@ -144,10 +144,29 @@ int main() { sycl::ONEAPI::atomic_fence(sycl::ONEAPI::memory_order::relaxed, sycl::ONEAPI::memory_scope::work_group); - // expected-warning@+1{{'INTEL' is deprecated: use 'ext::intel' instead}} + // expected-warning@+1{{'INTEL' is deprecated: use 'ext::intel::experimental' instead}} auto SL = sycl::INTEL::source_language::opencl_c; (void)SL; + // expected-warning@+1{{'intel' is deprecated: use 'ext::intel::experimental' instead}} + auto SLExtIntel = sycl::ext::intel::source_language::opencl_c; + (void)SLExtIntel; + + // expected-warning@+1{{'level_zero' is deprecated: use 'ext_oneapi_level_zero' instead}} + auto LevelZeroBackend = sycl::backend::level_zero; + (void)LevelZeroBackend; + + sycl::half Val = 1.0f; + // expected-warning@+1{{'bit_cast' is deprecated: use 'sycl::bit_cast' instead}} + auto BitCastRes = sycl::detail::bit_cast(Val); + (void)BitCastRes; + + // expected-warning@+1{{'submit_barrier' is deprecated: use 'ext_oneapi_submit_barrier' instead}} + Queue.submit_barrier(); + + // expected-warning@+1{{'barrier' is deprecated: use 'ext_oneapi_barrier' instead}} + Queue.submit([&](sycl::handler &CGH) { CGH.barrier(); }); + // expected-warning@+1{{'half' is deprecated: use 'sycl::half' instead}} half H; (void)H;