Skip to content

Commit b185269

Browse files
gmlueckgmlueck
authored
[SYCL][DOC] Deprecate old spec constant extension (#5676)
This old extension for specialization constants is marked deprecated in the headers, so make the specification deprecated too.
1 parent 4bd6d20 commit b185269

File tree

2 files changed

+77
-73
lines changed

2 files changed

+77
-73
lines changed

‎sycl/doc/extensions/deprecated/sycl_ext_oneapi_spec_constants.md

+74
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,74 @@
1+
# Specialization constants.
2+
3+
Specialization constant is basically a variable in a SYCL program set by host
4+
code and used in device code which appears to be constant for the online (JIT)
5+
compiler of the device code. Things like optimal tile size in a tiled matrix
6+
multiplication kernel may depend on the hardware and can be expressed via a
7+
specialization constant for better code generation.
8+
9+
This version of oneAPI provides experimental implementation of specialization
10+
constants based on the
11+
[proposal](https://github.com/codeplaysoftware/standards-proposals/blob/master/spec-constant/index.md)
12+
from Codeplay.
13+
14+
**NOTE:** This extension is now deprecated. Use the core SYCL specialization
15+
constant APIs defined in the
16+
[SYCL 2020 specification](https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html)
17+
instead.
18+
19+
A specialization constant is identified by a C++ type name, similarly to a
20+
kernel, its value is set via `program::set_spec_constant` class API and is
21+
"frozen" once the program is built. The following example shows how
22+
different values of a specialization constant can be used within the same
23+
kernel:
24+
25+
```cpp
26+
for (int i = 0; i < n_sc_sets; i++) {
27+
cl::sycl::program program(q.get_context());
28+
const int *sc_set = &sc_vals[i][0];
29+
cl::sycl::ext::oneapi::experimental::spec_constant<int32_t, SC0> sc0 =
30+
program.set_spec_constant<SC0>(sc_set[0]);
31+
cl::sycl::ext::oneapi::experimental::spec_constant<int32_t, SC1> sc1 =
32+
program.set_spec_constant<SC1>(sc_set[1]);
33+
34+
program.build_with_kernel_type<KernelAAA>();
35+
36+
try {
37+
cl::sycl::buffer<int, 1> buf(vec.data(), vec.size());
38+
39+
q.submit([&](cl::sycl::handler &cgh) {
40+
auto acc = buf.get_access<cl::sycl::access::mode::write>(cgh);
41+
cgh.single_task<KernelAAA>(
42+
program.get_kernel<KernelAAA>(),
43+
[=]() {
44+
acc[i] = sc0.get() + sc1.get();
45+
});
46+
});
47+
} catch (cl::sycl::exception &e) {
48+
std::cout << "*** Exception caught: " << e.what() << "\n";
49+
return 1;
50+
}
51+
...
52+
}
53+
```
54+
Here the values of specialization constants `SC0` and `SC1` are changed on
55+
every loop iteration. All what's needed is re-creating a `program` class
56+
instance, setting new values and rebuilding it via
57+
`program::build_with_kernel_type`. JIT compiler will effectively replace
58+
`sc0.get()` and `sc1.get()` within thhe device code with the corresponding
59+
constant values (`sc_vals[i][0]` and `sc_vals[i][1]`). Full runnable example
60+
can be found on
61+
[github](https://github.com/intel/llvm-test-suite/blob/intel/SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp).
62+
63+
Specialization constants can be used in programs compiled Ahead-Of-Time, in this
64+
case a specialization constant takes default value for its type (as specified by
65+
[C++ standard](https://en.cppreference.com/w/cpp/language/value_initialization)).
66+
67+
#### Limitations
68+
- The implementation does not support the `template <unsigned NID> struct spec_constant_id`
69+
API design for interoperability with OpenCL - to set specializataion constants
70+
in SYCL programs originating from external SPIRV modules and wrapped by OpenCL
71+
program objects. In SPIRV/OpenCL specialization constants are identified by an
72+
integer number, and the `spec_constant_id` class models that.
73+
- Only primitive numeric types are supported.
74+

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_spec_constants.md

+3-73
Original file line numberDiff line numberDiff line change
@@ -1,74 +1,4 @@
1-
# Specialization constants.
2-
3-
Specialization constant is basically a variable in a SYCL program set by host
4-
code and used in device code which appears to be constant for the online (JIT)
5-
compiler of the device code. Things like optimal tile size in a tiled matrix
6-
multiplication kernel may depend on the hardware and can be expressed via a
7-
specialization constant for better code generation.
8-
9-
This version of oneAPI provides experimental implementation of specialization
10-
constants based on the
11-
[proposal](https://github.com/codeplaysoftware/standards-proposals/blob/master/spec-constant/index.md)
12-
from Codeplay.
13-
14-
**NOTE:** This extension is now deprecated. Use the core SYCL specialization
15-
constant APIs defined in the
16-
[SYCL 2020 specification](https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html)
17-
instead.
18-
19-
A specialization constant is identified by a C++ type name, similarly to a
20-
kernel, its value is set via `program::set_spec_constant` class API and is
21-
"frozen" once the program is built. The following example shows how
22-
different values of a specialization constant can be used within the same
23-
kernel:
24-
25-
```cpp
26-
for (int i = 0; i < n_sc_sets; i++) {
27-
cl::sycl::program program(q.get_context());
28-
const int *sc_set = &sc_vals[i][0];
29-
cl::sycl::ext::oneapi::experimental::spec_constant<int32_t, SC0> sc0 =
30-
program.set_spec_constant<SC0>(sc_set[0]);
31-
cl::sycl::ext::oneapi::experimental::spec_constant<int32_t, SC1> sc1 =
32-
program.set_spec_constant<SC1>(sc_set[1]);
33-
34-
program.build_with_kernel_type<KernelAAA>();
35-
36-
try {
37-
cl::sycl::buffer<int, 1> buf(vec.data(), vec.size());
38-
39-
q.submit([&](cl::sycl::handler &cgh) {
40-
auto acc = buf.get_access<cl::sycl::access::mode::write>(cgh);
41-
cgh.single_task<KernelAAA>(
42-
program.get_kernel<KernelAAA>(),
43-
[=]() {
44-
acc[i] = sc0.get() + sc1.get();
45-
});
46-
});
47-
} catch (cl::sycl::exception &e) {
48-
std::cout << "*** Exception caught: " << e.what() << "\n";
49-
return 1;
50-
}
51-
...
52-
}
53-
```
54-
Here the values of specialization constants `SC0` and `SC1` are changed on
55-
every loop iteration. All what's needed is re-creating a `program` class
56-
instance, setting new values and rebuilding it via
57-
`program::build_with_kernel_type`. JIT compiler will effectively replace
58-
`sc0.get()` and `sc1.get()` within thhe device code with the corresponding
59-
constant values (`sc_vals[i][0]` and `sc_vals[i][1]`). Full runnable example
60-
can be found on
61-
[github](https://github.com/intel/llvm-test-suite/blob/intel/SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp).
62-
63-
Specialization constants can be used in programs compiled Ahead-Of-Time, in this
64-
case a specialization constant takes default value for its type (as specified by
65-
[C++ standard](https://en.cppreference.com/w/cpp/language/value_initialization)).
66-
67-
#### Limitations
68-
- The implementation does not support the `template <unsigned NID> struct spec_constant_id`
69-
API design for interoperability with OpenCL - to set specializataion constants
70-
in SYCL programs originating from external SPIRV modules and wrapped by OpenCL
71-
program objects. In SPIRV/OpenCL specialization constants are identified by an
72-
integer number, and the `spec_constant_id` class models that.
73-
- Only primitive numeric types are supported.
1+
This extension has been deprecated, but the specification is still available
2+
[here][1].
743

4+
[1]: <../deprecated/sycl_ext_oneapi_spec_constants.md>

0 commit comments

Comments
 (0)