Skip to content

Commit 90094f5

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (9 commits)
2 parents 2fe0875 + 36ce10e commit 90094f5

File tree

93 files changed

+256
-244
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

93 files changed

+256
-244
lines changed

‎.github/workflows/sycl-post-commit.yml‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -52,6 +52,11 @@ jobs:
5252
runner: '["Linux", "arc"]'
5353
extra_lit_opts: --param matrix-xmx8=True
5454
reset_intel_gpu: true
55+
- name: Intel Battlemage Graphics with Level Zero
56+
runner: '["Linux", "bmg"]'
57+
target_devices: level_zero:gpu
58+
# The new Xe kernel driver used by BMG doesn't support resetting.
59+
reset_intel_gpu: false
5560
- name: AMD/HIP
5661
runner: '["Linux", "amdgpu"]'
5762
image_options: -u 1001 --device=/dev/dri --device=/dev/kfd

‎clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp‎

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -2362,14 +2362,6 @@ Expected<SmallVector<StringRef>> linkAndWrapDeviceFiles(
23622362
HasNonSYCLOffloadKinds = true;
23632363
}
23642364

2365-
// Write any remaining device inputs to an output file.
2366-
SmallVector<StringRef> InputFiles;
2367-
for (const OffloadFile &File : Input) {
2368-
auto FileNameOrErr = writeOffloadFile(File);
2369-
if (!FileNameOrErr)
2370-
return FileNameOrErr.takeError();
2371-
InputFiles.emplace_back(*FileNameOrErr);
2372-
}
23732365
if (HasSYCLOffloadKind) {
23742366
SmallVector<StringRef> InputFiles;
23752367
// Write device inputs to an output file for the linker.

‎devops/dependencies-igc-dev.json‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,10 @@
11
{
22
"linux": {
33
"igc_dev": {
4-
"github_tag": "igc-dev-61b96b3",
5-
"version": "61b96b3",
6-
"updated_at": "2025-01-15T17:43:30Z",
7-
"url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2435370337/zip",
4+
"github_tag": "igc-dev-4cc8dff",
5+
"version": "4cc8dff",
6+
"updated_at": "2025-02-10T10:27:30Z",
7+
"url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2564401848/zip",
88
"root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
99
}
1010
}

‎devops/scripts/install_build_tools.sh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,6 @@ apt update && apt install -yqq \
1010
python3-psutil \
1111
python-is-python3 \
1212
python3-pip \
13-
zstd \
1413
ocl-icd-opencl-dev \
1514
vim \
1615
libffi-dev \
@@ -21,6 +20,7 @@ apt update && apt install -yqq \
2120
zstd \
2221
zip \
2322
unzip \
23+
pigz \
2424
jq \
2525
curl \
2626
libhwloc-dev \

‎devops/scripts/update_drivers.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@ def uplift_linux_igfx_driver(config, platform_tag, igc_dev_only):
3838
config[platform_tag]["igc_dev"]["version"] = igcdevver
3939
config[platform_tag]["igc_dev"]["updated_at"] = igc_dev["updated_at"]
4040
config[platform_tag]["igc_dev"]["url"] = get_artifacts_download_url(
41-
"intel/intel-graphics-compiler", "IGC_Ubuntu22.04_llvm14_clang-" + igcdevver
41+
"intel/intel-graphics-compiler", "IGC_Ubuntu24.04_llvm14_clang-" + igcdevver
4242
)
4343
return config
4444

‎sycl/include/sycl/detail/vector_arith.hpp‎

Lines changed: 71 additions & 81 deletions
Original file line numberDiff line numberDiff line change
@@ -59,11 +59,39 @@ struct UnaryPlus {
5959
}
6060
};
6161

62-
struct VecOperators {
62+
// Tag to map/templatize the mixin for prefix/postfix inc/dec operators.
63+
struct IncDec {};
64+
65+
template <typename SelfOperandTy> struct IncDecImpl {
66+
using element_type = typename from_incomplete<SelfOperandTy>::element_type;
67+
using vec_t = simplify_if_swizzle_t<std::remove_const_t<SelfOperandTy>>;
68+
69+
public:
70+
friend SelfOperandTy &operator++(SelfOperandTy &x) {
71+
x += element_type{1};
72+
return x;
73+
}
74+
friend SelfOperandTy &operator--(SelfOperandTy &x) {
75+
x -= element_type{1};
76+
return x;
77+
}
78+
friend auto operator++(SelfOperandTy &x, int) {
79+
vec_t tmp{x};
80+
x += element_type{1};
81+
return tmp;
82+
}
83+
friend auto operator--(SelfOperandTy &x, int) {
84+
vec_t tmp{x};
85+
x -= element_type{1};
86+
return tmp;
87+
}
88+
};
89+
90+
template <typename Self> struct VecOperators {
91+
static_assert(is_vec_v<Self>);
92+
6393
template <typename OpTy, typename... ArgTys>
6494
static constexpr auto apply(const ArgTys &...Args) {
65-
using Self = nth_type_t<0, ArgTys...>;
66-
static_assert(is_vec_v<Self>);
6795
static_assert(((std::is_same_v<Self, ArgTys> && ...)));
6896

6997
using element_type = typename Self::element_type;
@@ -163,6 +191,41 @@ struct VecOperators {
163191
res[i] = Op(Args[i]...);
164192
return res;
165193
}
194+
195+
// Uglier than possible due to
196+
// https://gcc.gnu.org/bugzilla/show_bug.cgi?id=85282.
197+
template <typename Op, typename = void> struct OpMixin;
198+
199+
template <typename Op>
200+
struct OpMixin<Op, std::enable_if_t<std::is_same_v<Op, IncDec>>>
201+
: public IncDecImpl<Self> {};
202+
203+
#define __SYCL_VEC_UOP_MIXIN(OP, OPERATOR) \
204+
template <typename Op> \
205+
struct OpMixin<Op, std::enable_if_t<std::is_same_v<Op, OP>>> { \
206+
friend auto operator OPERATOR(const Self &v) { return apply<OP>(v); } \
207+
};
208+
209+
__SYCL_VEC_UOP_MIXIN(std::negate<void>, -)
210+
__SYCL_VEC_UOP_MIXIN(std::logical_not<void>, !)
211+
__SYCL_VEC_UOP_MIXIN(UnaryPlus, +)
212+
213+
template <typename Op>
214+
struct OpMixin<Op, std::enable_if_t<std::is_same_v<Op, std::bit_not<void>>>> {
215+
template <typename T = typename from_incomplete<Self>::element_type>
216+
friend std::enable_if_t<!is_vgenfloat_v<T>, Self> operator~(const Self &v) {
217+
return apply<std::bit_not<void>>(v);
218+
}
219+
};
220+
221+
#undef __SYCL_VEC_UOP_MIXIN
222+
223+
template <typename... Op>
224+
struct __SYCL_EBO CombineImpl : public OpMixin<Op>... {};
225+
226+
struct Combined
227+
: public CombineImpl<std::negate<void>, std::logical_not<void>,
228+
std::bit_not<void>, UnaryPlus, IncDec> {};
166229
};
167230

168231
// Macros to populate binary operation on sycl::vec.
@@ -174,7 +237,7 @@ struct VecOperators {
174237
template <typename T = DataT> \
175238
friend std::enable_if_t<(COND), vec_t> operator BINOP(const vec_t & Lhs, \
176239
const vec_t & Rhs) { \
177-
return VecOperators::apply<FUNCTOR>(Lhs, Rhs); \
240+
return VecOperators<vec_t>::template apply<FUNCTOR>(Lhs, Rhs); \
178241
} \
179242
\
180243
template <typename T = DataT> \
@@ -200,65 +263,11 @@ struct VecOperators {
200263
return Lhs; \
201264
}
202265

203-
/****************************************************************
204-
* vec_arith_common
205-
* / | \
206-
* / | \
207-
* vec_arith<int> vec_arith<float> ... vec_arith<byte>
208-
* \ | /
209-
* \ | /
210-
* sycl::vec<T>
211-
*
212-
* vec_arith_common is the base class for vec_arith. It contains
213-
* the common math operators of sycl::vec for all types.
214-
* vec_arith is the derived class that contains the math operators
215-
* specialized for certain types. sycl::vec inherits from vec_arith.
216-
* *************************************************************/
217-
template <typename DataT, int NumElements> class vec_arith_common;
218-
template <typename DataT> struct vec_helper;
219-
220266
template <typename DataT, int NumElements>
221-
class vec_arith : public vec_arith_common<DataT, NumElements> {
267+
class vec_arith : public VecOperators<vec<DataT, NumElements>>::Combined {
222268
protected:
223269
using vec_t = vec<DataT, NumElements>;
224270
using ocl_t = detail::fixed_width_signed<sizeof(DataT)>;
225-
template <typename T> using vec_data = vec_helper<T>;
226-
227-
// operator!.
228-
friend vec<ocl_t, NumElements> operator!(const vec_t &Rhs) {
229-
return VecOperators::apply<std::logical_not<void>>(Rhs);
230-
}
231-
232-
// operator +.
233-
friend vec_t operator+(const vec_t &Lhs) {
234-
return VecOperators::apply<UnaryPlus>(Lhs);
235-
}
236-
237-
// operator -.
238-
friend vec_t operator-(const vec_t &Lhs) {
239-
return VecOperators::apply<std::negate<void>>(Lhs);
240-
}
241-
242-
// Unary operations on sycl::vec
243-
// FIXME: Don't allow Unary operators on vec<bool> after
244-
// https://github.com/KhronosGroup/SYCL-CTS/issues/896 gets fixed.
245-
#ifdef __SYCL_UOP
246-
#error "Undefine __SYCL_UOP macro"
247-
#endif
248-
#define __SYCL_UOP(UOP, OPASSIGN) \
249-
friend vec_t &operator UOP(vec_t & Rhs) { \
250-
Rhs OPASSIGN DataT{1}; \
251-
return Rhs; \
252-
} \
253-
friend vec_t operator UOP(vec_t &Lhs, int) { \
254-
vec_t Ret(Lhs); \
255-
Lhs OPASSIGN DataT{1}; \
256-
return Ret; \
257-
}
258-
259-
__SYCL_UOP(++, +=)
260-
__SYCL_UOP(--, -=)
261-
#undef __SYCL_UOP
262271

263272
// The logical operations on scalar types results in 0/1, while for vec<>,
264273
// logical operations should result in 0 and -1 (similar to OpenCL vectors).
@@ -272,7 +281,7 @@ class vec_arith : public vec_arith_common<DataT, NumElements> {
272281
template <typename T = DataT> \
273282
friend std::enable_if_t<(COND), vec<ocl_t, NumElements>> operator RELLOGOP( \
274283
const vec_t & Lhs, const vec_t & Rhs) { \
275-
return VecOperators::apply<FUNCTOR>(Lhs, Rhs); \
284+
return VecOperators<vec_t>::template apply<FUNCTOR>(Lhs, Rhs); \
276285
} \
277286
\
278287
template <typename T = DataT> \
@@ -325,13 +334,13 @@ class vec_arith : public vec_arith_common<DataT, NumElements> {
325334
#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
326335
template <int NumElements>
327336
class vec_arith<std::byte, NumElements>
328-
: public vec_arith_common<std::byte, NumElements> {
337+
: public VecOperators<vec<std::byte, NumElements>>::template OpMixin<
338+
std::bit_not<void>> {
329339
protected:
330340
// NumElements can never be zero. Still using the redundant check to avoid
331341
// incomplete type errors.
332342
using DataT = typename std::conditional_t<NumElements == 0, int, std::byte>;
333343
using vec_t = vec<DataT, NumElements>;
334-
template <typename T> using vec_data = vec_helper<T>;
335344

336345
// Special <<, >> operators for std::byte.
337346
// std::byte is not an arithmetic type and it only supports the following
@@ -376,25 +385,6 @@ class vec_arith<std::byte, NumElements>
376385
};
377386
#endif // (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
378387

379-
template <typename DataT, int NumElements> class vec_arith_common {
380-
protected:
381-
using vec_t = vec<DataT, NumElements>;
382-
383-
static constexpr bool IsBfloat16 =
384-
std::is_same_v<DataT, sycl::ext::oneapi::bfloat16>;
385-
386-
// operator~() available only when: dataT != float && dataT != double
387-
// && dataT != half
388-
template <typename T = DataT>
389-
friend std::enable_if_t<!detail::is_vgenfloat_v<T>, vec_t>
390-
operator~(const vec_t &Rhs) {
391-
return VecOperators::apply<std::bit_not<void>>(Rhs);
392-
}
393-
394-
// friends
395-
template <typename T1, int T2> friend class __SYCL_EBO vec;
396-
};
397-
398388
#undef __SYCL_BINOP
399389

400390
} // namespace detail

‎sycl/include/sycl/vector.hpp‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -435,7 +435,6 @@ class __SYCL_EBO vec
435435
template <typename T1, int T2> friend class __SYCL_EBO vec;
436436
// To allow arithmetic operators access private members of vec.
437437
template <typename T1, int T2> friend class detail::vec_arith;
438-
template <typename T1, int T2> friend class detail::vec_arith_common;
439438
};
440439
///////////////////////// class sycl::vec /////////////////////////
441440

‎sycl/test-e2e/AddressSanitizer/lit.local.cfg‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,3 +19,6 @@ config.unsupported_features += ['gpu-intel-gen12']
1919

2020
# CMPLRLLVM-64052
2121
config.unsupported_features += ['spirv-backend']
22+
23+
# https://github.com/intel/llvm/issues/16920
24+
config.unsupported_features += ['arch-intel_gpu_bmg_g21']

‎sycl/test-e2e/AtomicRef/atomic_memory_order_acq_rel.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %{build} -O3 -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %}
1+
// RUN: %{build} -O3 -o %t.out %if target-nvidia %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %}
22
// RUN: %{run} %t.out
33

44
// NOTE: Tests fetch_add for acquire and release memory ordering.

‎sycl/test-e2e/AtomicRef/atomic_memory_order_seq_cst.cpp‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,7 @@
1-
// RUN: %{build} -O3 -o %t.out %if any-device-is-cuda %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %}
1+
// RUN: %{build} -O3 -o %t.out %if target-nvidia %{ -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70 %}
22
// RUN: %{run} %t.out
3+
// UNSUPPORTED: arch-intel_gpu_bmg_g21
4+
// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/16924
35

46
#include "atomic_memory_order.h"
57
#include <cmath>

0 commit comments

Comments
 (0)