Skip to content

[SYCL][Doc] Fix spec for passing SPIRV struct args #12338

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 2 commits into from
Jan 19, 2024
Merged

[SYCL][Doc] Fix spec for passing SPIRV struct args #12338

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
6c0e8f0
[SYCL][Doc] Fix spec for passing SPIRV struct args
gmlueck Jan 9, 2024
851a60f
Merge branch 'sycl' into gmlueck/spirv-struct-args
gmlueck Jan 19, 2024
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
31 changes: 18 additions & 13 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -172,7 +172,7 @@ The SPIR-V *OpEntryPoint* that defines a kernel has an associated *OpFunction*.
This *OpFunction* is followed by a list of *OpFunctionParameter* instructions,
one for each kernel argument.
The following table defines the set of argument types that are supported by
this extension, and explains how to pass each type of argument from SYCL.
this extension and explains how to pass each type of argument from SYCL.
However, the set of supported argument types may be further limited by the
backend's SPIR-V client API specification.

Expand All @@ -198,21 +198,26 @@ of 32, `double` when the *OpTypeFloat* has a width of 64, and `sycl::half` when
the *OpTypeFloat* has a width of 16.
_{endnote}_]

|*OpTypePointer* with storage class *CrossWorkgroup*
|Either a pointer to USM memory or an `accessor` whose target is
`target::device`.
|*OpTypePointer* with _Storage Class_ *CrossWorkgroup*
|Either a {cpp} pointer (typically a pointer to USM memory) or an `accessor`
whose target is `target::device`.

|*OpTypePointer* with storage class *Workgroup*
|*OpTypePointer* with _Storage Class_ *Workgroup*
|A `local_accessor`.

|*OpTypeStruct* whose member types are limited to *OpTypeInt*, *OpTypeFloat*,
and *OpTypePointer* (with storage class *CrossWorkgroup*).
The *OpTypeStruct* may contain members which are also *OpTypeStruct* so long
as its members are limited to the same types.
|A C++ struct or class that is device copyable.
Each member variable must have the corresponding type as defined above,
except that an *OpTypePointer* member must correspond to a USM pointer.
It is not valid to pass an `accessor` for these members.
|*OpTypePointer* with _Storage Class_ *Function* and _Type_ *OpTypeStruct*
(i.e. the pointed-at type is *OpTypeStruct*).
|A C++ struct or class that is device copyable, which has the same size and
data representation as the *OpTypeStruct*.

[_Note:_ The SYCL argument is a structure even though the SPIR-V argument type
is a pointer because structures are passed by reference.

The SYCL argument must not contain any `accessor` or `local_accessor` members
because these types are not device copyable.
If the *OpTypeStruct* contains an *OpTypePointer* member, the corresponding SYCL
structure member is typically a USM pointer.
_{endnote}_]
|===

When data allocated on the host is accessed by the kernel via a pointer, the
Expand Down