@@ -17,7 +17,7 @@ ordinal number. This complicates the design, as the compiler
1717
1818Simple source code example:
1919
20- ``` cpp
20+ ```
2121class MyInt32Const;
2222...
2323 sycl::program p(q.get_context());
@@ -46,7 +46,7 @@ primitive numeric types. POD types support is described further in the document.
4646
4747Key ` spec_constant::get() ` function implementation for the device code:
4848
49- ```cpp
49+ ```
5050template <typename T, typename ID = T> class spec_constant {
5151...
5252public:
@@ -87,7 +87,7 @@ After this pass the sycl-post-link tool will output the
8787attaching this info to the device binary image via the offload wrapper tool as
8888a property set:
8989
90- ``` cpp
90+ ```
9191struct pi_device_binary_struct {
9292...
9393 // Array of preperty sets; e.g. specialization constants symbol-int ID map is
@@ -112,7 +112,7 @@ the value of a spec constant.
112112
113113Given the ` __spirv_SpecConstant ` intrinsic calls produced by the
114114` SpecConstants ` pass:
115- ```cpp
115+ ```
116116; Function Attrs: alwaysinline
117117define dso_local spir_func i32 @get() local_unnamed_addr #0 {
118118 ; args are "ID" and "default value":
@@ -124,7 +124,7 @@ define dso_local spir_func i32 @get() local_unnamed_addr #0 {
124124the translator will generate ` OpSpecConstant ` SPIR-V instructions with proper
125125` SpecId ` decorations:
126126
127- ``` cpp
127+ ```
128128 OpDecorate %i32 SpecId 42 ; ID
129129 %i32 = OpSpecConstant %int 0 ; Default value
130130 %1 = OpTypeFunction %int
@@ -152,7 +152,7 @@ unaware of the clang-specific built-ins.
152152Before JIT-ing a program, the runtime "flushes" the spec constants: it iterates
153153through the value map and invokes the
154154
155- ```cpp
155+ ```
156156pi_result piextProgramSetSpecializationConstant(pi_program prog,
157157 pi_uint32 spec_id,
158158 size_t spec_size,
@@ -167,7 +167,7 @@ Plugin Interface function for each entry, taking the `spec_id` from the ID map.
167167
168168Say, the POD type is
169169
170- ```cpp
170+ ```
171171struct A {
172172 int x;
173173 float y;
@@ -181,7 +181,7 @@ struct POD {
181181
182182and the user says
183183
184- ``` cpp
184+ ```
185185 POD gold{
186186 {
187187 { goldi, goldf },
@@ -199,7 +199,7 @@ and the user says
199199
200200 - The SpecConstants pass in the post-link will have the following IR as input (` sret ` conversion is omitted for clarity):
201201
202- ```cpp
202+ ```
203203 %spec_const = call %struct.POD __sycl_getCompositeSpecConstantValue<POD type mangling> ("MyConst_mangled")
204204```
205205
@@ -214,7 +214,7 @@ where `__sycl_getCompositeSpecConstantValue` is a new "intrinsic"
214214 specialization constant's type (` %struct.POD ` ), the pass will traverse its leaf
215215 fields and generate 5 "primitive" spec constants using already existing SPIR-V intrinsic:
216216
217- ``` cpp
217+ ```
218218%gold_POD_a0x = call i32 __spirv_SpecConstant(i32 10, i32 0)
219219%gold_POD_a0y = call float __spirv_SpecConstant(i32 11, float 0)
220220%gold_POD_a1x = call i32 __spirv_SpecConstant(i32 12, i32 0)
@@ -224,7 +224,7 @@ where `__sycl_getCompositeSpecConstantValue` is a new "intrinsic"
224224
225225And 1 "composite"
226226
227- ```cpp
227+ ```
228228 %gold_POD = call %struct.POD __spirvCompositeSpecConstant<POD type mangling>(i32 10, i32 11, i32 12, i32 13, i32 14)
229229```
230230
@@ -244,15 +244,15 @@ passed to the runtime. Also, for a composite specialization constant there is
244244no ID map entry within the meta information, and the composite constant is
245245referenced by its symbolic ID. For example:
246246
247- ``` cpp
247+ ```
248248MyConst_mangled [10,int,0,4],[11,float,4,4],[12,int,8,4],[13,float,12,4],[14,int,16,4]
249249```
250250
251251#### LLVMIR-\> SPIR-V translator
252252
253253The translator aims to create the following code (pseudo-code)
254254
255- ``` cpp
255+ ```
256256%gold_POD_a0x = OpSpecConstant(0) [SpecId = 10]
257257%gold_POD_a0y = OpSpecConstant(0.0f) [SpecId = 11]
258258%gold_POD_a1x = OpSpecConstant(0) [SpecId = 12]
0 commit comments