Skip to content

Commit e1dd434

Browse files
wenju-hewenju-he
authored
Merge branch 'sycl' into libspirv-bitfield_extract-insert
2 parents e8560eb + 066e266 commit e1dd434

File tree

20 files changed

+100
-42
lines changed

20 files changed

+100
-42
lines changed

‎.github/workflows/sycl-linux-run-tests.yml‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -330,7 +330,8 @@ jobs:
330330
- run: which clang++ sycl-ls
331331
- run: sycl-ls --verbose
332332
- run: SYCL_UR_TRACE=1 sycl-ls
333-
- run: |
333+
- name: Print IGC version
334+
run: |
334335
if [ -f /usr/local/lib/igc/IGCTAG.txt ]; then
335336
cat /usr/local/lib/igc/IGCTAG.txt
336337
fi

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 14 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -7302,22 +7302,30 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
73027302
}
73037303

73047304
if (FreeFunctionCount > 0) {
7305+
// GlobalMapUpdater has to be in an anonymous namespace.
7306+
// Otherwise, if multiple translation units include the same integration
7307+
// header, there will be multiple varying definitions of GlobalMapUpdater
7308+
// with the same name across translation units, violating the C++'s One
7309+
// Definition Rule. Putting it in an anonymous namespace gives each
7310+
// translation unit its own unique definition.
7311+
73057312
O << "\n#include <sycl/kernel_bundle.hpp>\n";
73067313
O << "#include <sycl/detail/kernel_global_info.hpp>\n";
7307-
O << "namespace sycl {\n";
7308-
O << "inline namespace _V1 {\n";
7309-
O << "namespace detail {\n";
7314+
O << "namespace {\n";
73107315
O << "struct GlobalMapUpdater {\n";
73117316
O << " GlobalMapUpdater() {\n";
73127317
O << " sycl::detail::free_function_info_map::add("
73137318
<< "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
73147319
<< KernelDescs.size() << ");\n";
73157320
O << " }\n";
7321+
O << " ~GlobalMapUpdater() {\n";
7322+
O << " sycl::detail::free_function_info_map::remove("
7323+
<< "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
7324+
<< KernelDescs.size() << ");\n";
7325+
O << " }\n";
73167326
O << "};\n";
73177327
O << "static GlobalMapUpdater updater;\n";
7318-
O << "} // namespace detail\n";
7319-
O << "} // namespace _V1\n";
7320-
O << "} // namespace sycl\n";
7328+
O << "} // namespace\n";
73217329
}
73227330
}
73237331

‎clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp‎

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1087,15 +1087,14 @@ namespace Testing::Tests {
10871087

10881088
// CHECK: #include <sycl/kernel_bundle.hpp>
10891089
// CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
1090-
// CHECK-NEXT: namespace sycl {
1091-
// CHECK-NEXT: inline namespace _V1 {
1092-
// CHECK-NEXT: namespace detail {
1090+
// CHECK-NEXT: namespace {
10931091
// CHECK-NEXT: struct GlobalMapUpdater {
10941092
// CHECK-NEXT: GlobalMapUpdater() {
10951093
// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 28);
10961094
// CHECK-NEXT: }
1095+
// CHECK-NEXT: ~GlobalMapUpdater() {
1096+
// CHECK-NEXT: sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 28);
1097+
// CHECK-NEXT: }
10971098
// CHECK-NEXT: };
10981099
// CHECK-NEXT: static GlobalMapUpdater updater;
1099-
// CHECK-NEXT: } // namespace detail
1100-
// CHECK-NEXT: } // namespace _V1
1101-
// CHECK-NEXT: } // namespace sycl
1100+
// CHECK-NEXT: }

‎clang/test/CodeGenSYCL/free_function_int_header.cpp‎

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1567,15 +1567,14 @@ void ff_24(int arg) {
15671567

15681568
// CHECK: #include <sycl/kernel_bundle.hpp>
15691569
// CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
1570-
// CHECK-NEXT: namespace sycl {
1571-
// CHECK-NEXT: inline namespace _V1 {
1572-
// CHECK-NEXT: namespace detail {
1570+
// CHECK-NEXT: namespace {
15731571
// CHECK-NEXT: struct GlobalMapUpdater {
15741572
// CHECK-NEXT: GlobalMapUpdater() {
15751573
// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 33);
15761574
// CHECK-NEXT: }
1575+
// CHECK-NEXT: ~GlobalMapUpdater() {
1576+
// CHECK-NEXT: sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 33);
1577+
// CHECK-NEXT: }
15771578
// CHECK-NEXT: };
15781579
// CHECK-NEXT: static GlobalMapUpdater updater;
1579-
// CHECK-NEXT: } // namespace detail
1580-
// CHECK-NEXT: } // namespace _V1
1581-
// CHECK-NEXT: } // namespace sycl
1580+
// CHECK-NEXT: }

‎clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp‎

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -75,15 +75,14 @@ int main(){
7575

7676
// CHECK-NORTC: #include <sycl/kernel_bundle.hpp>
7777
// CHECK-NORTC-NEXT: #include <sycl/detail/kernel_global_info.hpp>
78-
// CHECK-NORTC-NEXT: namespace sycl {
79-
// CHECK-NORTC-NEXT: inline namespace _V1 {
80-
// CHECK-NORTC-NEXT: namespace detail {
78+
// CHECK-NORTC-NEXT: namespace {
8179
// CHECK-NORTC-NEXT: struct GlobalMapUpdater {
8280
// CHECK-NORTC-NEXT: GlobalMapUpdater() {
8381
// CHECK-NORTC-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
8482
// CHECK-NORTC-NEXT: }
83+
// CHECK-NORTC-NEXT: ~GlobalMapUpdater() {
84+
// CHECK-NORTC-NEXT: sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
85+
// CHECK-NORTC-NEXT: }
8586
// CHECK-NORTC-NEXT: };
8687
// CHECK-NORTC-NEXT: static GlobalMapUpdater updater;
87-
// CHECK-NORTC-NEXT: } // namespace detail
88-
// CHECK-NORTC-NEXT: } // namespace _V1
89-
// CHECK-NORTC-NEXT: } // namespace sycl
88+
// CHECK-NORTC-NEXT: }

‎clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp‎

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -10,14 +10,11 @@
1010

1111
#include <sycl/detail/export.hpp>
1212

13-
namespace sycl {
14-
inline namespace _V1 {
15-
namespace detail {
13+
namespace {
1614
namespace free_function_info_map {
1715

1816
__SYCL_EXPORT void add(const void *DeviceGlobalPtr, const char *UniqueId);
17+
__SYCL_EXPORT void remove(const void *DeviceGlobalPtr, const char *UniqueId);
1918

2019
} // namespace free_function_info_map
21-
} // namespace detail
22-
} // namespace _V1
23-
} // namespace sycl
20+
}

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

Lines changed: 9 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,7 @@ inputs:
3838
required: False
3939

4040
runs:
41+
# composite actions don't make use of 'name', so copy-paste names as a comment in the first line of each step
4142
using: "composite"
4243
steps:
4344
- name: Check specified runner type / target backend
@@ -46,6 +47,7 @@ runs:
4647
TARGET_DEVICE: ${{ inputs.target_devices }}
4748
PRESET: ${{ inputs.preset }}
4849
run: |
50+
# Check specified runner type / target backend
4951
case "$RUNNER_TAG" in
5052
'["PVC_PERF"]' ) ;;
5153
'["BMG_PERF"]' ) ;;
@@ -114,6 +116,7 @@ runs:
114116
- name: Build Unified Runtime
115117
shell: bash
116118
run: |
119+
# Build Unified Runtime
117120
# Sparse-checkout UR at build ref:
118121
git clone --depth 1 --no-checkout https://github.com/intel/llvm ur
119122
cd ur
@@ -140,14 +143,16 @@ runs:
140143
- name: Checkout results repo
141144
shell: bash
142145
run: |
146+
# Checkout results repo
143147
git clone -b unify-ci https://github.com/intel/llvm-ci-perf-results
144-
- name: Run compute-benchmarks
148+
- name: Build and run benchmarks
145149
env:
146150
# Need to append "_<device>_<backend>" to save name in order to follow
147151
# conventions:
148152
SAVE_PREFIX: ${{ inputs.save_name }}
149153
shell: bash
150154
run: |
155+
# Build and run benchmarks
151156
# TODO generate summary + display helpful message here
152157
export CMPLR_ROOT=./toolchain
153158
echo "-----"
@@ -231,7 +236,8 @@ runs:
231236
- name: Cache changes and upload github summary
232237
if: always()
233238
shell: bash
234-
run: |
239+
run: |
240+
# Cache changes and upload github summary
235241
[ -f "github_summary.md" ] && cat github_summary.md >> $GITHUB_STEP_SUMMARY
236242
237243
cd "./llvm-ci-perf-results"
@@ -244,6 +250,7 @@ runs:
244250
if: inputs.upload_results == 'true' && always()
245251
shell: bash
246252
run: |
253+
# Push benchmarks results
247254
cd "./llvm-ci-perf-results"
248255
git config user.name "SYCL Benchmarking Bot"
249256
git config user.email "[email protected]"

‎libclc/libspirv/lib/generic/SOURCES‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,7 @@ integer/add_sat.cl
4242
integer/bitfield_extract_signed.cl
4343
integer/bitfield_extract_unsigned.cl
4444
integer/bitfield_insert.cl
45+
integer/bit_reverse.cl
4546
integer/clamp.cl
4647
integer/clz.cl
4748
integer/ctz.cl

‎libclc/libspirv/lib/generic/integer/BitCount.cl‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,10 +6,11 @@
66
//
77
//===----------------------------------------------------------------------===//
88

9+
#include <clc/integer/clc_popcount.h>
910
#include <libspirv/spirv.h>
1011

1112
#define __CLC_FUNCTION __spirv_BitCount
12-
#define __CLC_IMPL_FUNCTION(x) __spirv_ocl_popcount
13+
#define __CLC_IMPL_FUNCTION(x) __clc_popcount
1314
#define __CLC_BODY <clc/shared/unary_def.inc>
1415

1516
#include <clc/integer/gentype.inc>

‎libclc/libspirv/lib/generic/integer/abs.cl‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,6 @@
66
//
77
//===----------------------------------------------------------------------===//
88

9-
#include <clc/internal/clc.h>
109
#include <clc/integer/clc_abs.h>
1110
#include <libspirv/spirv.h>
1211

0 commit comments

Comments
 (0)