Skip to content

Commit fb8ca1f

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (8 commits)
2 parents 62b8a00 + 0f39e01 commit fb8ca1f

30 files changed

+570
-152
lines changed

‎devops/actions/run-tests/benchmark/action.yml‎

Lines changed: 8 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -37,9 +37,6 @@ inputs:
3737
type: string
3838
required: False
3939

40-
env:
41-
BENCHMARK_RESULTS_BRANCH: "sycl-benchmark-ci-results"
42-
4340
runs:
4441
# composite actions don't make use of 'name', so copy-paste names as a comment in the first line of each step
4542
using: "composite"
@@ -164,6 +161,14 @@ runs:
164161
sudo cmake --install build
165162
166163
cd -
164+
- name: Set env var for results branch
165+
shell: bash
166+
run: |
167+
# Set env var for results branch
168+
# Set BENCHMARK_RESULTS_BRANCH globally for all subsequent steps.
169+
# This has to be done this way because of limits of composite actions.
170+
BENCHMARK_RESULTS_BRANCH="sycl-benchmark-ci-results"
171+
echo "BENCHMARK_RESULTS_BRANCH=$BENCHMARK_RESULTS_BRANCH" >> $GITHUB_ENV
167172
- name: Checkout results repo
168173
uses: actions/checkout@v5
169174
with:

‎devops/dependencies.json‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -19,9 +19,9 @@
1919
"root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
2020
},
2121
"level_zero": {
22-
"github_tag": "v1.24.3",
23-
"version": "v1.24.3",
24-
"url": "https://github.com/oneapi-src/level-zero/releases/tag/v1.24.3",
22+
"github_tag": "v1.25.2",
23+
"version": "v1.25.2",
24+
"url": "https://github.com/oneapi-src/level-zero/releases/tag/v1.25.2",
2525
"root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
2626
},
2727
"tbb": {

‎sycl/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,7 @@ option(SYCL_ADD_DEV_VERSION_POSTFIX "Adds -V postfix to version string" ON)
1111
option(SYCL_ENABLE_COVERAGE "Enables code coverage for runtime and unit tests" OFF)
1212
option(SYCL_ENABLE_STACK_PRINTING "Enables stack printing on crashes of SYCL applications" OFF)
1313
option(SYCL_LIB_WITH_DEBUG_SYMBOLS "Builds SYCL runtime libraries with debug symbols" OFF)
14+
option(SYCL_ENABLE_IPO "Builds SYCL runtime libraries with interprocedural optimization" OFF)
1415

1516
if (NOT SYCL_COVERAGE_PATH)
1617
set(SYCL_COVERAGE_PATH "${CMAKE_CURRENT_BINARY_DIR}/profiles")

‎sycl/doc/design/spirv-extensions/SPV_INTEL_bfloat16_arithmetic.asciidoc‎

Lines changed: 145 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,145 @@
1+
:extension_name: SPV_INTEL_bfloat16_arithmetic
2+
:capability_name: BFloat16ArithmeticINTEL
3+
:capability_token: 6226
4+
:dep_ext_name: SPV_KHR_bfloat16
5+
:dep_ext_rev: 1
6+
:dep_ext_capability: BFloat16TypeKHR
7+
8+
= {extension_name}
9+
10+
== Name Strings
11+
12+
{extension_name}
13+
14+
== Contact
15+
16+
To report problems with this extension, please open a new issue at:
17+
18+
https://github.com/intel/llvm
19+
20+
== Contributors
21+
22+
* Ben Ashbaugh, Intel
23+
* Victor Mustya, Intel
24+
25+
== Notice
26+
27+
Copyright (c) 2025 Intel Corporation. All rights reserved.
28+
29+
== Status
30+
31+
* Working Draft
32+
33+
This is a preview extension specification, intended to provide early access to
34+
a feature for review and community feedback. When the feature matures, this
35+
specification may be released as a formal extension.
36+
37+
Because the interfaces defined by this specification are not final and are
38+
subject to change they are not intended to be used by shipping software
39+
products. If you are interested in using this feature in your software product,
40+
please let us know!
41+
42+
== Version
43+
44+
[width="40%",cols="25,25"]
45+
|========================================
46+
| Last Modified Date | {docdate}
47+
| Revision | 1
48+
|========================================
49+
50+
== Dependencies
51+
52+
This extension is written against the SPIR-V Specification, Version 1.6 Revision
53+
5, and `{dep_ext_name}`, Revision {dep_ext_rev}.
54+
55+
This extension requires SPIR-V 1.0 and `{dep_ext_name}`, Revision {dep_ext_rev}.
56+
57+
== Overview
58+
59+
This extension allows to use *OpTypeFloat* instruction to define *BFloat16KHR*
60+
data type and use it for all the floating-point insructions.
61+
62+
== Extension Name
63+
64+
To use this extension within a SPIR-V module, the following *OpExtension* must
65+
be present in the module:
66+
67+
[subs="attributes"]
68+
----
69+
OpExtension "{extension_name}"
70+
----
71+
72+
== Modifications to the SPIR-V Specification, Version 1.6
73+
74+
=== Validation Rules
75+
76+
Add validation rules to section 2.16.1 Universal Validation Rules from:
77+
78+
* Variables with a type that is or includes a floating-point type with the
79+
*BFloat16KHR* encoding must only be used with the following instructions,
80+
if *{capability_name}* is declared:
81+
82+
** https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_mode_setting_instructions[Arithmetic Instructions]:
83+
84+
*** *OpFNegate*
85+
*** *OpFAdd*
86+
*** *OpFSub*
87+
*** *OpFMul*
88+
*** *OpFDiv*
89+
*** *OpFRem*
90+
*** *OpFMod*
91+
*** *OpVectorTimesScalar*
92+
93+
** https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_relational_and_logical_instructions[Relational and Logical Instructions]
94+
95+
* Variables with a type that is or includes a floating-point type with the
96+
*BFloat16KHR* encoding can be used with *OpExtInst* instruction with the
97+
opcodes from the https://registry.khronos.org/SPIR-V/specs/unified1/OpenCL.ExtendedInstructionSet.100.html[OpenCL extended instruction set].
98+
99+
=== Capabilities
100+
101+
Modify Section 3.31, "Capability", adding these rows to the Capability table:
102+
103+
--
104+
[cols="^.^2,16,15",options="header"]
105+
|====
106+
2+^.^| Capability | Implicitly Declares
107+
| {capability_token} | *{capability_name}* +
108+
Allows *OpTypeFloat* with the *BFloat16KHR* <<Floating_Point_Encoding, floating point encoding>>
109+
to be used with arithmetic instructions.
110+
| *{dep_ext_capability}*
111+
|====
112+
--
113+
114+
== Interactions with Other Extensions
115+
116+
If `SPV_INTEL_masked_gather_scatter` is supported then *OpMaskedGatherINTEL*
117+
and *OpMaskedGatherINTEL* also accept vector of *BFloat16KHR* and vector of
118+
pointers to *BFloat16KHR*.
119+
120+
== Issues
121+
122+
. Should the extension be INTEL or EXT/KHR?
123+
+
124+
--
125+
*UNRESOLVED*: The functionality is generic, so it probably may be useful for
126+
non-Intel implementations.
127+
--
128+
129+
. Should we define *BFloat16* atomics and interaction with
130+
`SPV_EXT_shader_atomic_float_min_max` and `SPV_EXT_shader_atomic_float_add`
131+
by this extension?
132+
+
133+
--
134+
*RESOLVED*: The atomic operations should be defined by a separate extension.
135+
--
136+
137+
== Revision History
138+
139+
[cols="5,15,15,70"]
140+
[grid="rows"]
141+
[options="header"]
142+
|========================================
143+
|Rev|Date|Author|Changes
144+
|1|2025-05-07|Victor Mustya|Initial public revision
145+
|========================================

‎sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc‎

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -994,6 +994,45 @@ int main() {
994994
}
995995
----
996996

997+
=== Using "scratch" work-group local memory
998+
999+
Free function kernels can use work-group local memory via `local_accessor` or
1000+
via other extensions.
1001+
This example demonstrates how to use work-group local memory via the
1002+
link:../experimental/sycl_ext_oneapi_work_group_scratch_memory.asciidoc[
1003+
sycl_ext_oneapi_work_group_scratch_memory] extension because it also illustrates
1004+
how to pass a kernel launch parameter.
1005+
1006+
[source,c++]
1007+
----
1008+
#include <sycl/sycl.hpp>
1009+
namespace syclexp = sycl::ext::oneapi::experimental;
1010+
1011+
static constexpr size_t NUM = 1024;
1012+
static constexpr size_t WGSIZE = 16;
1013+
1014+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
1015+
void mykernel() {
1016+
// Gets a pointer to WGSIZE int's
1017+
int *ptr = static_cast<int *>(syclexp::get_work_group_scratch_memory());
1018+
}
1019+
1020+
int main() {
1021+
sycl::queue q;
1022+
1023+
syclexp::launch_config cfg{
1024+
sycl::nd_range{{NUM}, {WGSIZE}},
1025+
syclexp::properties{
1026+
syclexp::work_group_scratch_size{WGSIZE * sizeof(int)}
1027+
}
1028+
};
1029+
syclexp::nd_launch(q, cfg, syclexp::kernel_function<mykernel>);
1030+
1031+
q.wait();
1032+
}
1033+
----
1034+
1035+
9971036

9981037
[[level-zero-and-opencl-compatibility]]
9991038
== {dpcpp} guaranteed compatibility with Level Zero and OpenCL backends

‎sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -157,7 +157,7 @@ void single_task(queue Q, const KernelType &KernelObj,
157157
if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
158158
void>::value)) {
159159
detail::submit_kernel_direct_single_task<KernelName>(
160-
std::move(Q), KernelObj, empty_properties_t{}, CodeLoc);
160+
std::move(Q), KernelObj, {}, empty_properties_t{}, CodeLoc);
161161
} else {
162162
submit(
163163
std::move(Q),
@@ -312,7 +312,7 @@ void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
312312
LaunchConfigAccess(Config);
313313

314314
detail::submit_kernel_direct_parallel_for<KernelName>(
315-
std::move(Q), LaunchConfigAccess.getRange(), KernelObj,
315+
std::move(Q), LaunchConfigAccess.getRange(), KernelObj, {},
316316
LaunchConfigAccess.getProperties());
317317
} else {
318318
submit(std::move(Q), [&](handler &CGH) {

‎sycl/include/sycl/khr/free_function_commands.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -317,7 +317,7 @@ void launch_task(const sycl::queue &q, KernelType &&k,
317317
if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
318318
void>::value)) {
319319
detail::submit_kernel_direct_single_task(
320-
q, std::forward<KernelType>(k),
320+
q, std::forward<KernelType>(k), {},
321321
ext::oneapi::experimental::empty_properties_t{}, codeLoc);
322322
} else {
323323
submit(q, [&](handler &h) { launch_task<KernelType>(h, k); }, codeLoc);

0 commit comments

Comments
 (0)