Skip to content

[SYCL][Doc] Fix KernelParameterPassing.md markup #2077

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 1 commit into from
Jul 10, 2020
Merged

[SYCL][Doc] Fix KernelParameterPassing.md markup #2077

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
51 changes: 26 additions & 25 deletions sycl/doc/KernelParameterPassing.md
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
<h2>SYCL Kernel Parameter Handling and Array Support</h2>
# SYCL Kernel Parameter Handling and Array Support

<h3>Introduction</h3>
## Introduction

This document describes how parameters of SYCL kernels are passed
from host to device. Support for arrays as kernel parameters was added
Expand Down Expand Up @@ -28,7 +28,8 @@ The first few sections describe the overall design.
The last three sections provide additional details of array support.
The implementation of this design is confined to four classes in the
file `SemaSYCL.cpp`.
<h3>A SYCL Kernel</h3>

## A SYCL Kernel

The SYCL constructs `single_task`, `parallel_for`, and
`parallel_for_work_group` each take a function object or a lambda function
Expand All @@ -37,7 +38,7 @@ lambda function is executed on the device.
To enable execution of the kernel on OpenCL devices, the lambda/function object
is converted into the format of an OpenCL kernel.

<h3>SYCL Kernel Code Generation</h3>
## SYCL Kernel Code Generation

Consider a source code example that captures an int, a struct and an accessor
by value:
Expand Down Expand Up @@ -112,7 +113,7 @@ spir_kernel void caller(
// Reassemble capture object from parts
local.i = i;
local.s = s;
// Call accessor�s init function
// Call accessor's init function
sycl::accessor::init(&local.outAcc, AccData, AccR1, AccR2, I);

// Call the kernel body
Expand Down Expand Up @@ -140,7 +141,7 @@ host to device separately. The values received on the device
are passed to the `init` functions executed on the device,
which results in the reassembly of the SYCL object in a form usable on the device.

There is one other aspect of code generation. An �integration header�
There is one other aspect of code generation. An "integration header"
is generated for use during host compilation.
This header file contains entries for each kernel.
Among the items it defines is a table of sizes and offsets of the
Expand Down Expand Up @@ -169,7 +170,7 @@ object which contains three values:
The previous sections described how kernel arguments are handled today.
The next three sections describe support for arrays.

<h3>Fix 1: Kernel Arguments that are Standard-Layout Arrays</h3>
## Fix 1: Kernel Arguments that are Standard-Layout Arrays

As described earlier, each variable captured by a lambda that comprises a
SYCL kernel becomes a parameter of the kernel caller function.
Expand All @@ -180,7 +181,7 @@ the purposes of passing to the device. Each array element is passed as a
separate parameter. The array elements received on the device
are copied into the array within the local capture object.

<h4>Source code fragment:</h4>
**Source code fragment:**

```C++
constexpr int num_items = 2;
Expand All @@ -197,7 +198,7 @@ are copied into the array within the local capture object.
});
```

<h4>Integration header produced:</h4>
**Integration header produced:**

```C++
static constexpr
Expand All @@ -211,7 +212,7 @@ const kernel_param_desc_t kernel_signatures[] = {

```

<h4>The changes to device code made to support this extension, in pseudo-code:</h4>
**The changes to device code made to support this extension, in pseudo-code:**

```C++
struct Capture {
Expand All @@ -238,22 +239,22 @@ spir_kernel void caller(
// Initialize array using existing clang Initialization mechanisms
local.array[0] = p_array_0;
local.array[1] = p_array_1;
// Call accessor�s init function
// Call accessor's init function
sycl::accessor::init(&local.outAcc, AccData, AccR1, AccR2, I);

callee(&local, id<1> wi);
}
```

<h3>Fix 2: Kernel Arguments that are Arrays of Accessors</h3>
## Fix 2: Kernel Arguments that are Arrays of Accessors

Arrays of accessors are supported in a manner similar to that of a plain
accessor. For each accessor array element, the four values required to
call its init function are passed as separate arguments to the kernel.
Reassembly within the kernel caller is done by calling the `init` functions
of each accessor array element in ascending index value.

<h4>Source code fragment:</h4>
**Source code fragment:**

```C++
myQueue.submit([&](handler &cgh) {
Expand All @@ -269,7 +270,7 @@ of each accessor array element in ascending index value.
});
```

<h4>Integration header:</h4>
**Integration header:**

```C++
static constexpr
Expand All @@ -281,7 +282,7 @@ const kernel_param_desc_t kernel_signatures[] = {
};
```

<h4>Device code generated in pseudo-code form:</h4>
**Device code generated in pseudo-code form:**

```C++
struct Capture {
Expand Down Expand Up @@ -311,20 +312,20 @@ spir_kernel void caller(
struct Capture local;

// Reassemble capture object from parts
// Call outAcc accessor�s init function
// Call outAcc accessor's init function
sycl::accessor::init(&local.outAcc, outAccData, outAccR1, outAccR2, outI);

// Call inAcc[0] accessor�s init function
// Call inAcc[0] accessor's init function
sycl::accessor::init(&local.inAcc[0], inAccData_0, inAccR1_0, inAccR2_0, inI_0);

// Call inAcc[1] accessor�s init function
// Call inAcc[1] accessor's init function
sycl::accessor::init(&local.inAcc[1], inAccData_1, inAccR1_1, inAccR2_1, inI_1);

callee(&local, id<1> wi);
}
```

<h3>Fix 3: Accessor Arrays within Structs</h3>
## Fix 3: Accessor Arrays within Structs

Kernel parameters that are structs are traversed member
by member, recursively, to enumerate member structs that are one of
Expand All @@ -340,7 +341,7 @@ Within the kernel caller function, the lambda object is reassembled
in a manner similar to other instances of accessor arrays.


<h4>Source code fragment:</h4>
**Source code fragment:**

```C++
myQueue.submit([&](handler &cgh) {
Expand All @@ -361,7 +362,7 @@ in a manner similar to other instances of accessor arrays.
});
```

<h4>Integration header:</h4>
**Integration header:**

```C++
static constexpr
Expand All @@ -375,7 +376,7 @@ const kernel_param_desc_t kernel_signatures[] = {
};
```

<h4>Device code generated in pseudo-code form:</h4>
**Device code generated in pseudo-code form:**

```C++
struct Capture {
Expand Down Expand Up @@ -411,15 +412,15 @@ spir_kernel void caller(
local.s = s;

// 2. Initialize accessors by calling init functions
// 2a. Call outAcc accessor�s init function
// 2a. Call outAcc accessor's init function
sycl::accessor::init(
&local.outAcc, outAccData, outAccR1, outAccR2, outI);

// 2b. Call s.inAcc[0] accessor�s init function
// 2b. Call s.inAcc[0] accessor's init function
sycl::accessor::init(
&local.s.inAcc[0], inAccData_0, inAccR1_0, inAccR2_0, inI_0);

// 2c. Call s.inAcc[1] accessor�s init function
// 2c. Call s.inAcc[1] accessor's init function
sycl::accessor::init(
&local.s.inAcc[1], inAccData_1, inAccR1_1, inAccR2_1, inI_1);

Expand Down
1 change: 1 addition & 0 deletions sycl/doc/index.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,7 @@ Developing oneAPI DPC++ Compiler

API Reference <https://intel.github.io/llvm-docs/doxygen>
CompilerAndRuntimeDesign
KernelParameterPassing
EnvironmentVariables
PluginInterface
ABIPolicyGuide