Skip to content

[SYCL][Doc] Fix Sphinx docs index #2599

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.

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
Oct 6, 2020
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
26 changes: 13 additions & 13 deletions sycl/doc/SpecializationConstants.md
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@ ordinal number. This complicates the design, as the compiler

Simple source code example:

```cpp
```
class MyInt32Const;
...
sycl::program p(q.get_context());
Expand Down Expand Up @@ -46,7 +46,7 @@ primitive numeric types. POD types support is described further in the document.

Key `spec_constant::get()` function implementation for the device code:

```cpp
```
template <typename T, typename ID = T> class spec_constant {
...
public:
Expand Down Expand Up @@ -87,7 +87,7 @@ After this pass the sycl-post-link tool will output the
attaching this info to the device binary image via the offload wrapper tool as
a property set:

```cpp
```
struct pi_device_binary_struct {
...
// Array of preperty sets; e.g. specialization constants symbol-int ID map is
Expand All @@ -112,7 +112,7 @@ the value of a spec constant.

Given the `__spirv_SpecConstant` intrinsic calls produced by the
`SpecConstants` pass:
```cpp
```
; Function Attrs: alwaysinline
define dso_local spir_func i32 @get() local_unnamed_addr #0 {
; args are "ID" and "default value":
Expand All @@ -124,7 +124,7 @@ define dso_local spir_func i32 @get() local_unnamed_addr #0 {
the translator will generate `OpSpecConstant` SPIR-V instructions with proper
`SpecId` decorations:

```cpp
```
OpDecorate %i32 SpecId 42 ; ID
%i32 = OpSpecConstant %int 0 ; Default value
%1 = OpTypeFunction %int
Expand Down Expand Up @@ -152,7 +152,7 @@ unaware of the clang-specific built-ins.
Before JIT-ing a program, the runtime "flushes" the spec constants: it iterates
through the value map and invokes the

```cpp
```
pi_result piextProgramSetSpecializationConstant(pi_program prog,
pi_uint32 spec_id,
size_t spec_size,
Expand All @@ -167,7 +167,7 @@ Plugin Interface function for each entry, taking the `spec_id` from the ID map.

Say, the POD type is

```cpp
```
struct A {
int x;
float y;
Expand All @@ -181,7 +181,7 @@ struct POD {

and the user says

```cpp
```
POD gold{
{
{ goldi, goldf },
Expand All @@ -199,7 +199,7 @@ and the user says

- The SpecConstants pass in the post-link will have the following IR as input (`sret` conversion is omitted for clarity):

```cpp
```
%spec_const = call %struct.POD __sycl_getCompositeSpecConstantValue<POD type mangling> ("MyConst_mangled")
```

Expand All @@ -214,7 +214,7 @@ where `__sycl_getCompositeSpecConstantValue` is a new "intrinsic"
specialization constant's type (`%struct.POD`), the pass will traverse its leaf
fields and generate 5 "primitive" spec constants using already existing SPIR-V intrinsic:

```cpp
```
%gold_POD_a0x = call i32 __spirv_SpecConstant(i32 10, i32 0)
%gold_POD_a0y = call float __spirv_SpecConstant(i32 11, float 0)
%gold_POD_a1x = call i32 __spirv_SpecConstant(i32 12, i32 0)
Expand All @@ -224,7 +224,7 @@ where `__sycl_getCompositeSpecConstantValue` is a new "intrinsic"

And 1 "composite"

```cpp
```
%gold_POD = call %struct.POD __spirvCompositeSpecConstant<POD type mangling>(i32 10, i32 11, i32 12, i32 13, i32 14)
```

Expand All @@ -244,15 +244,15 @@ passed to the runtime. Also, for a composite specialization constant there is
no ID map entry within the meta information, and the composite constant is
referenced by its symbolic ID. For example:

```cpp
```
MyConst_mangled [10,int,0,4],[11,float,4,4],[12,int,8,4],[13,float,12,4],[14,int,16,4]
```

#### LLVMIR-\>SPIR-V translator

The translator aims to create the following code (pseudo-code)

```cpp
```
%gold_POD_a0x = OpSpecConstant(0) [SpecId = 10]
%gold_POD_a0y = OpSpecConstant(0.0f) [SpecId = 11]
%gold_POD_a1x = OpSpecConstant(0) [SpecId = 12]
Expand Down
1 change: 1 addition & 0 deletions sycl/doc/index.rst
Original file line number Diff line number Diff line change
Expand Up @@ -27,3 +27,4 @@ Developing oneAPI DPC++ Compiler
EnvironmentVariables
PluginInterface
ABIPolicyGuide
SpecializationConstants