Skip to content

[SYCL][Docs] Align kernel_args_restrict extension with SYCL 2020 to inherit critical fixes #5793

New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
May 13, 2022
Merged

[SYCL][Docs] Align kernel_args_restrict extension with SYCL 2020 to inherit critical fixes #5793

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
8c43200
Align kernel_args_restrict extension with fixes that we factored into…
Mar 11, 2022
4c7f2b3
Update extension to use new template
Mar 15, 2022
fcee412
Apply wording clarifications from PR discussion
Apr 11, 2022
37238e6
Remove wording about SPIR-V/OpenCL equivalence
Apr 11, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
173 changes: 81 additions & 92 deletions sycl/doc/extensions/supported/sycl_ext_intel_kernel_args_restrict.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,5 @@
= SYCL_INTEL_kernel_restrict_all
= sycl_intel_kernel_restrict_all

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

Expand All @@ -8,138 +9,126 @@
:toc: left
:encoding: utf-8
:lang: en

:blank: pass:[ +]
:dpcpp: pass:[DPC++]

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

== Introduction
IMPORTANT: This specification is a draft.

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.

NOTE: This document is better viewed when rendered as html with asciidoctor. GitHub does not render image icons.

This document describes an extension that adds a function type attribute which has the same effect as adding the C99 `restrict` attribute to all pointer arguments when applied to a kernel function.


== Name Strings

+SYCL_INTEL_kernel_restrict_all+

== Notice

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

== Status
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.

Working Draft

This is a preview extension specification, intended to provide early access to a feature for review and community feedback. When the feature matures, this specification may be released as a formal extension.
== Contact

Because the interfaces defined by this specification are not final and are subject to change they are not intended to be used by shipping software products.
To report problems with this extension, please open a new issue at:

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

Built On: {docdate} +
Revision: 1

== Contact
Michael Kinsner, Intel (michael 'dot' kinsner 'at' intel 'dot' com)

== Dependencies

This extension is written against the SYCL 1.2.1 specification, Revision v1.2.1-5.

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

This extension adds a function type attribute that has the same effect as adding the C99 (or equivalently OpenCL C kernel language) `restrict` attribute to all pointers and the pointer member of any accessors, that are function arguments, lambda captures, or functor members, of the callable to which the attribute was applied. The attribute can be applied to kernel lambdas, function call operators of a functor, or arbitrary functions (the effect on arbitrary functions, if any, is implementaton defined), to provide the compiler with additional information for optimization.

A new attribute is added by this extension because there is no clear location on which to manually annotate C99 `restrict` on lambda captures, functor members, or accessors.
== Status

== Add new paragraphs to end of section 6.7 (Attributes)
This extension is implemented and fully supported by {dpcpp}.

The attribute `intel::kernel_args_restrict` is legal on device functions, and can be ignored on non-device functions. When applied to a lambda or function call operator (of a functor) that defines a kernel, the attribute is a hint to the compiler equivalent to specifying the C99 `restrict` attribute on all pointer arguments or the pointer member of any accessors, which are a function argument, lambda capture, or functor member, of the callable to which the attribute was applied. This effect is equivalent to annotating `restrict` on *all* kernel pointer arguments in an OpenCL or SPIR-V kernel, if the callable is a kernel. If `intel::kernel_args_restrict` is applied to a function called from a device kernel, the effect is implementation defined. The attribute forms an unchecked assertion, in that implementations do not need to check/confirm the pre-condition in any way. If a user applies `intel::kernel_args_restrict` to a kernel, but there is in fact aliasing between kernel pointer arguments at runtime, the behavior is undefined.
== Overview

The attribute-token `intel::kernel_args_restrict` shall appear at most once in each _attribute-list_ and no _attribute-argument-clause_ shall be present. The attribute may be applied to the _function-type_ in a function declaration. The first declaration of a function shall specify the `intel::kernel_args_restrict` attribute if any declaration of that function specifies the `intel::kernel_args_restrict` attribute. If a function is declared with the `intel::kernel_args_restrict` attribute in one translation unit and the same function is declared without the `intel::kernel_args_restrict` attribute in another translation unit, the program is ill-formed and no diagnostic is required.
This extension adds a kernel function attribute that has similar effect as
the C99 `restrict` type qualifier. When a kernel is decorated with this
attribute, all pointers and accessors (treated as if each accessor was a
pointer) that are captured as kernel arguments
are assumed to point to disjoint objects in memory. This provides the
compiler with additional information for optimization.

The `intel::kernel_args_restrict` attribute has an effect when applied to a function, and no effect otherwise.
A new attribute is added by this extension because there is no clear location
on which to manually annotate C99 `restrict` on lambda captures, function
object members, or accessors.

== Example uses
=== Example uses

[source,c++,Restrict on lambda,linenums]
----
Q.submit([&](handler &cgh) {
auto acc1 = out_buf_1.get_access<access::mode::write>(cgh);
auto acc2 = out_buf_2.get_access<access::mode::write>(cgh);
cgh.parallel_for<class lambda_foo>(
range<1>(N), [=](id<1> wiid) [[intel::kernel_args_restrict]] {
int id = wiid[0];
acc1[id]=id;
acc2[id]=id*2;
});
});
Q.submit([&](handler &cgh) {
auto acc1 = out_buf_1.get_access<access::mode::write>(cgh);
auto acc2 = out_buf_2.get_access<access::mode::write>(cgh);
cgh.parallel_for<class lambda_foo>(
range<1>(N), [=](id<1> wiid) [[intel::kernel_args_restrict]] {
int id = wiid[0];
acc1[id]=id;
acc2[id]=id*2;
});
});
----

[source,c++,Restrict on functor,linenums]
[source,c++,Restrict on function object,linenums]
----
class functor_foo {
...
void operator()(item<1> item) [[intel::kernel_args_restrict]]
{
int id = item[0];
buf1_m[id]=id;
buf2_m[id]=id*2;
}
...
[[intel::kernel_args_restrict]]
void operator()(item<1> item)
{
int id = item[0];
buf1_m[id]=id;
buf2_m[id]=id*2;
}
};
----



== Issues

None.
== Specification

//. Title
//+
//--
//*RESOLUTION*: Description
//--

== Feature test macro
=== Feature test macro

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

[%header,cols="1,5"]
|===
|Value |Description
|1 |Initial extension version. Base features are supported.
|Value
|Description

|1
|Initial version of this extension.
|===

== Revision History

[cols="5,15,15,70"]
[grid="rows"]
[options="header"]
|========================================
|Rev|Date|Author|Changes
|1|2019-11-11|Michael Kinsner|*Initial public working draft*
|========================================

//************************************************************************
//Other formatting suggestions:
//
//* Use *bold* text for host APIs, or [source] syntax highlighting.
//* Use +mono+ text for device APIs, or [source] syntax highlighting.
//* Use +mono+ text for extension names, types, or enum values.
//* Use _italics_ for parameters.
//************************************************************************

=== Add new entry to Table 180 in section 5.8.1 (Kernel attributes)

[width="100%",options="header",separator="@",cols="65%,35%"]
|====
@ SYCL attribute @ Description
a@
[source]
----
intel::kernel_args_restrict
----
a@ Hint to the compiler equivalent to specifying the C99 `restrict`
attribute on all pointers and accessors (treated as if each accessor was a
pointer) that are captured as kernel arguments.

The attribute forms an unchecked assertion, in that implementations do not need
to check/confirm the pre-condition in any way. If a user applies `intel::kernel_args_restrict`
to a kernel, but there is in fact aliasing between accessors and/or pointers at runtime,
the behavior is undefined.

|====