Skip to content

Commit 8662149

Browse files
DenisBakhvalovDenisBakhvalov
authored
[ESIMD] Fix leakage of detail namespace (#3673)
This patch removes `using namespace sycl::INTEL::gpu::detail;` which leaks into the global namespace.
1 parent 041ca27 commit 8662149

File tree

4 files changed

+77
-69
lines changed

4 files changed

+77
-69
lines changed

‎sycl/include/CL/sycl/INTEL/esimd/esimd.hpp

Lines changed: 31 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,6 @@ namespace sycl {
1818
namespace INTEL {
1919
namespace gpu {
2020

21-
using namespace sycl::INTEL::gpu::detail;
22-
2321
/// The simd vector class.
2422
///
2523
/// This is a wrapper class for llvm vector values. Additionally this class
@@ -51,13 +49,13 @@ template <typename Ty, int N> class simd {
5149
if constexpr (std::is_same<SrcTy, Ty>::value)
5250
set(other.data());
5351
else
54-
set(__builtin_convertvector(other.data(), vector_type_t<Ty, N>));
52+
set(__builtin_convertvector(other.data(), detail::vector_type_t<Ty, N>));
5553
}
5654
template <typename SrcTy> constexpr simd(simd<SrcTy, N> &&other) {
5755
if constexpr (std::is_same<SrcTy, Ty>::value)
5856
set(other.data());
5957
else
60-
set(__builtin_convertvector(other.data(), vector_type_t<Ty, N>));
58+
set(__builtin_convertvector(other.data(), detail::vector_type_t<Ty, N>));
6159
}
6260
constexpr simd(const vector_type &Val) { set(Val); }
6361

@@ -134,7 +132,7 @@ template <typename Ty, int N> class simd {
134132

135133
/// View this simd object in a different element type.
136134
template <typename EltTy> auto format() & {
137-
using TopRegionTy = compute_format_type_t<simd, EltTy>;
135+
using TopRegionTy = detail::compute_format_type_t<simd, EltTy>;
138136
using RetTy = simd_view<simd, TopRegionTy>;
139137
TopRegionTy R(0);
140138
return RetTy{*this, R};
@@ -144,7 +142,8 @@ template <typename Ty, int N> class simd {
144142
//
145143
/// View as a 2-dimensional simd_view.
146144
template <typename EltTy, int Height, int Width> auto format() & {
147-
using TopRegionTy = compute_format_type_2d_t<simd, EltTy, Height, Width>;
145+
using TopRegionTy =
146+
detail::compute_format_type_2d_t<simd, EltTy, Height, Width>;
148147
using RetTy = simd_view<simd, TopRegionTy>;
149148
TopRegionTy R(0, 0);
150149
return RetTy{*this, R};
@@ -190,7 +189,7 @@ template <typename Ty, int N> class simd {
190189
/// Read multiple elements by their indices in vector
191190
template <int Size>
192191
simd<Ty, Size> iselect(const simd<uint16_t, Size> &Indices) {
193-
vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(Ty);
192+
detail::vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(Ty);
194193
return __esimd_rdindirect<Ty, N, Size>(data(), Offsets);
195194
}
196195
// TODO ESIMD_EXPERIMENTAL
@@ -205,7 +204,7 @@ template <typename Ty, int N> class simd {
205204
template <int Size>
206205
void iupdate(const simd<uint16_t, Size> &Indices, const simd<Ty, Size> &Val,
207206
mask_type_t<Size> Mask) {
208-
vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(Ty);
207+
detail::vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(Ty);
209208
set(__esimd_wrindirect<Ty, N, Size>(data(), Val.data(), Offsets, Mask));
210209
}
211210

@@ -217,18 +216,18 @@ template <typename Ty, int N> class simd {
217216
// * if not different, then auto should not be used
218217
#define DEF_BINOP(BINOP, OPASSIGN) \
219218
ESIMD_INLINE friend auto operator BINOP(const simd &X, const simd &Y) { \
220-
using ComputeTy = compute_type_t<simd>; \
221-
auto V0 = convert<typename ComputeTy::vector_type>(X.data()); \
222-
auto V1 = convert<typename ComputeTy::vector_type>(Y.data()); \
219+
using ComputeTy = detail::compute_type_t<simd>; \
220+
auto V0 = detail::convert<typename ComputeTy::vector_type>(X.data()); \
221+
auto V1 = detail::convert<typename ComputeTy::vector_type>(Y.data()); \
223222
auto V2 = V0 BINOP V1; \
224223
return ComputeTy(V2); \
225224
} \
226225
ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const simd &RHS) { \
227-
using ComputeTy = compute_type_t<simd>; \
228-
auto V0 = convert<typename ComputeTy::vector_type>(LHS.data()); \
229-
auto V1 = convert<typename ComputeTy::vector_type>(RHS.data()); \
226+
using ComputeTy = detail::compute_type_t<simd>; \
227+
auto V0 = detail::convert<typename ComputeTy::vector_type>(LHS.data()); \
228+
auto V1 = detail::convert<typename ComputeTy::vector_type>(RHS.data()); \
230229
auto V2 = V0 BINOP V1; \
231-
LHS.write(convert<vector_type>(V2)); \
230+
LHS.write(detail::convert<vector_type>(V2)); \
232231
return LHS; \
233232
} \
234233
ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const Ty &RHS) { \
@@ -255,7 +254,7 @@ template <typename Ty, int N> class simd {
255254
const simd &Y) { \
256255
auto R = X.data() RELOP Y.data(); \
257256
mask_type_t<N> M(1); \
258-
return M & convert<mask_type_t<N>>(R); \
257+
return M & detail::convert<mask_type_t<N>>(R); \
259258
}
260259

261260
DEF_RELOP(>)
@@ -276,7 +275,7 @@ template <typename Ty, int N> class simd {
276275
ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const simd &RHS) { \
277276
static_assert(std::is_integral<Ty>(), "not integeral type"); \
278277
auto V2 = LHS.data() BITWISE_OP RHS.data(); \
279-
LHS.write(convert<vector_type>(V2)); \
278+
LHS.write(detail::convert<vector_type>(V2)); \
280279
return LHS; \
281280
} \
282281
ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const Ty &RHS) { \
@@ -401,17 +400,18 @@ template <typename Ty, int N> class simd {
401400

402401
/// Write a simd-vector into a basic region of a simd object.
403402
template <typename RTy>
404-
ESIMD_INLINE void writeRegion(
405-
RTy Region,
406-
const vector_type_t<typename RTy::element_type, RTy::length> &Val) {
403+
ESIMD_INLINE void
404+
writeRegion(RTy Region,
405+
const detail::vector_type_t<typename RTy::element_type,
406+
RTy::length> &Val) {
407407
using ElemTy = typename RTy::element_type;
408408
if constexpr (N * sizeof(Ty) == RTy::length * sizeof(ElemTy))
409409
// update the entire vector
410-
set(bitcast<Ty, ElemTy, RTy::length>(Val));
410+
set(detail::bitcast<Ty, ElemTy, RTy::length>(Val));
411411
else {
412412
static_assert(!RTy::Is_2D);
413413
// If element type differs, do bitcast conversion first.
414-
auto Base = bitcast<ElemTy, Ty, N>(data());
414+
auto Base = detail::bitcast<ElemTy, Ty, N>(data());
415415
constexpr int BN = (N * sizeof(Ty)) / sizeof(ElemTy);
416416
// Access the region information.
417417
constexpr int M = RTy::Size_x;
@@ -422,28 +422,28 @@ template <typename Ty, int N> class simd {
422422
auto Merged = __esimd_wrregion<ElemTy, BN, M,
423423
/*VS*/ 0, M, Stride>(Base, Val, Offset);
424424
// Convert back to the original element type, if needed.
425-
set(bitcast<Ty, ElemTy, BN>(Merged));
425+
set(detail::bitcast<Ty, ElemTy, BN>(Merged));
426426
}
427427
}
428428

429429
/// Write a simd-vector into a nested region of a simd object.
430430
template <typename TR, typename UR>
431-
ESIMD_INLINE void
432-
writeRegion(std::pair<TR, UR> Region,
433-
const vector_type_t<typename TR::element_type, TR::length> &Val) {
431+
ESIMD_INLINE void writeRegion(
432+
std::pair<TR, UR> Region,
433+
const detail::vector_type_t<typename TR::element_type, TR::length> &Val) {
434434
// parent-region type
435435
using PaTy = typename shape_type<UR>::type;
436436
using ElemTy = typename TR::element_type;
437437
using BT = typename PaTy::element_type;
438438
constexpr int BN = PaTy::length;
439439

440440
if constexpr (PaTy::Size_in_bytes == TR::Size_in_bytes) {
441-
writeRegion(Region.second, bitcast<BT, ElemTy, TR::length>(Val));
441+
writeRegion(Region.second, detail::bitcast<BT, ElemTy, TR::length>(Val));
442442
} else {
443443
// Recursively read the base
444-
auto Base = readRegion<Ty, N>(data(), Region.second);
444+
auto Base = detail::readRegion<Ty, N>(data(), Region.second);
445445
// If element type differs, do bitcast conversion first.
446-
auto Base1 = bitcast<ElemTy, BT, BN>(Base);
446+
auto Base1 = detail::bitcast<ElemTy, BT, BN>(Base);
447447
constexpr int BN1 = PaTy::Size_in_bytes / sizeof(ElemTy);
448448

449449
if constexpr (!TR::Is_2D) {
@@ -474,7 +474,7 @@ template <typename Ty, int N> class simd {
474474
Base1, Val, Offset);
475475
}
476476
// Convert back to the original element type, if needed.
477-
auto Merged1 = bitcast<BT, ElemTy, BN1>(Base1);
477+
auto Merged1 = detail::bitcast<BT, ElemTy, BN1>(Base1);
478478
// recursively write it back to the base
479479
writeRegion(Region.second, Merged1);
480480
}
@@ -495,7 +495,7 @@ template <typename Ty, int N> class simd {
495495

496496
template <typename U, typename T, int n>
497497
ESIMD_INLINE simd<U, n> convert(simd<T, n> val) {
498-
return __builtin_convertvector(val.data(), vector_type_t<U, n>);
498+
return __builtin_convertvector(val.data(), detail::vector_type_t<U, n>);
499499
}
500500

501501
} // namespace gpu

‎sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp

Lines changed: 24 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -113,7 +113,8 @@ ESIMD_NODEBUG ESIMD_INLINE
113113
std::is_integral<U>::value,
114114
simd<T0, SZ>>
115115
esimd_shl(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
116-
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
116+
typedef
117+
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
117118
typename detail::simd_type<ComputationTy>::type Src0 = src0;
118119
typename detail::simd_type<ComputationTy>::type Src1 = src1;
119120

@@ -151,7 +152,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
151152
std::is_integral<T1>::value && std::is_integral<T2>::value,
152153
typename sycl::detail::remove_const_t<T0>>
153154
esimd_shl(T1 src0, T2 src1, int flag = GENX_NOSAT) {
154-
typedef typename computation_type<T1, T2>::type ComputationTy;
155+
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
155156
typename detail::simd_type<ComputationTy>::type Src0 = src0;
156157
typename detail::simd_type<ComputationTy>::type Src1 = src1;
157158
simd<T0, 1> Result = esimd_shl<T0>(Src0, Src1, flag);
@@ -166,7 +167,8 @@ ESIMD_NODEBUG ESIMD_INLINE
166167
std::is_integral<U>::value,
167168
simd<T0, SZ>>
168169
esimd_shr(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
169-
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
170+
typedef
171+
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
170172
typename detail::simd_type<ComputationTy>::type Src0 = src0;
171173
typename detail::simd_type<ComputationTy>::type Src1 = src1;
172174
typename detail::simd_type<ComputationTy>::type Result =
@@ -185,7 +187,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
185187
std::is_integral<T1>::value && std::is_integral<T2>::value,
186188
typename sycl::detail::remove_const_t<T0>>
187189
esimd_shr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
188-
typedef typename computation_type<T1, T2>::type ComputationTy;
190+
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
189191
typename detail::simd_type<ComputationTy>::type Src0 = src0;
190192
typename detail::simd_type<ComputationTy>::type Src1 = src1;
191193
simd<T0, 1> Result = esimd_shr<T0>(Src0, Src1, flag);
@@ -207,7 +209,8 @@ ESIMD_NODEBUG ESIMD_INLINE
207209
std::is_integral<U>::value,
208210
simd<T0, SZ>>
209211
esimd_rol(simd<T1, SZ> src0, U src1) {
210-
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
212+
typedef
213+
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
211214
typename detail::simd_type<ComputationTy>::type Src0 = src0;
212215
typename detail::simd_type<ComputationTy>::type Src1 = src1;
213216
return __esimd_rol<T0>(Src0.data(), Src1.data());
@@ -220,7 +223,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
220223
std::is_integral<T1>::value && std::is_integral<T2>::value,
221224
typename sycl::detail::remove_const_t<T0>>
222225
esimd_rol(T1 src0, T2 src1) {
223-
typedef typename computation_type<T1, T2>::type ComputationTy;
226+
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
224227
typename detail::simd_type<ComputationTy>::type Src0 = src0;
225228
typename detail::simd_type<ComputationTy>::type Src1 = src1;
226229
simd<T0, 1> Result = esimd_rol<T0>(Src0, Src1);
@@ -242,7 +245,8 @@ ESIMD_NODEBUG ESIMD_INLINE
242245
std::is_integral<U>::value,
243246
simd<T0, SZ>>
244247
esimd_ror(simd<T1, SZ> src0, U src1) {
245-
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
248+
typedef
249+
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
246250
typename detail::simd_type<ComputationTy>::type Src0 = src0;
247251
typename detail::simd_type<ComputationTy>::type Src1 = src1;
248252
return __esimd_ror<T0>(Src0.data(), Src1.data());
@@ -255,7 +259,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
255259
std::is_integral<T1>::value && std::is_integral<T2>::value,
256260
typename sycl::detail::remove_const_t<T0>>
257261
esimd_ror(T1 src0, T2 src1) {
258-
typedef typename computation_type<T1, T2>::type ComputationTy;
262+
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
259263
typename detail::simd_type<ComputationTy>::type Src0 = src0;
260264
typename detail::simd_type<ComputationTy>::type Src1 = src1;
261265
simd<T0, 1> Result = esimd_ror<T0>(Src0, Src1);
@@ -270,7 +274,7 @@ ESIMD_NODEBUG ESIMD_INLINE
270274
std::is_integral<U>::value,
271275
simd<T0, SZ>>
272276
esimd_lsr(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
273-
typedef typename computation_type<T1, T1>::type IntermedTy;
277+
typedef typename detail::computation_type<T1, T1>::type IntermedTy;
274278
typedef typename std::make_unsigned<IntermedTy>::type ComputationTy;
275279
simd<ComputationTy, SZ> Src0 = src0;
276280
simd<ComputationTy, SZ> Result = Src0.data() >> src1.data();
@@ -288,7 +292,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
288292
std::is_integral<T1>::value && std::is_integral<T2>::value,
289293
typename sycl::detail::remove_const_t<T0>>
290294
esimd_lsr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
291-
typedef typename computation_type<T1, T2>::type ComputationTy;
295+
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
292296
typename detail::simd_type<ComputationTy>::type Src0 = src0;
293297
typename detail::simd_type<ComputationTy>::type Src1 = src1;
294298
simd<T0, 1> Result = esimd_lsr<T0>(Src0, Src1, flag);
@@ -313,7 +317,7 @@ ESIMD_NODEBUG ESIMD_INLINE
313317
std::is_integral<U>::value,
314318
simd<T0, SZ>>
315319
esimd_asr(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
316-
typedef typename computation_type<T1, T1>::type IntermedTy;
320+
typedef typename detail::computation_type<T1, T1>::type IntermedTy;
317321
typedef typename std::make_signed<IntermedTy>::type ComputationTy;
318322
simd<ComputationTy, SZ> Src0 = src0;
319323
simd<ComputationTy, SZ> Result = Src0 >> src1;
@@ -331,7 +335,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
331335
std::is_integral<T1>::value && std::is_integral<T2>::value,
332336
typename sycl::detail::remove_const_t<T0>>
333337
esimd_asr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
334-
typedef typename computation_type<T1, T2>::type ComputationTy;
338+
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
335339
typename detail::simd_type<ComputationTy>::type Src0 = src0;
336340
typename detail::simd_type<ComputationTy>::type Src1 = src1;
337341
simd<T0, 1> Result = esimd_asr<T0>(Src0, Src1, flag);
@@ -358,7 +362,8 @@ ESIMD_NODEBUG ESIMD_INLINE
358362
detail::is_dword_type<U>::value,
359363
simd<T0, SZ>>
360364
esimd_imul(simd<T0, SZ> &rmd, simd<T1, SZ> src0, U src1) {
361-
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
365+
typedef
366+
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
362367
typename detail::simd_type<ComputationTy>::type Src0 = src0;
363368
typename detail::simd_type<ComputationTy>::type Src1 = src1;
364369
rmd = Src0 * Src1;
@@ -378,8 +383,8 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
378383
detail::is_dword_type<U>::value && SZ == 1,
379384
simd<T0, SZ>>
380385
esimd_imul(simd<T0, SZ> &rmd, simd<T1, SZ> src0, U src1) {
381-
typedef
382-
typename computation_type<decltype(rmd), long long>::type ComputationTy;
386+
typedef typename detail::computation_type<decltype(rmd), long long>::type
387+
ComputationTy;
383388
ComputationTy Product = convert<long long>(src0);
384389
Product *= src1;
385390
rmd = Product.format<T0>().select<1, 1>[0];
@@ -392,8 +397,8 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
392397
detail::is_dword_type<U>::value && SZ != 1,
393398
simd<T0, SZ>>
394399
esimd_imul(simd<T0, SZ> &rmd, simd<T1, SZ> src0, U src1) {
395-
typedef
396-
typename computation_type<decltype(rmd), long long>::type ComputationTy;
400+
typedef typename detail::computation_type<decltype(rmd), long long>::type
401+
ComputationTy;
397402
ComputationTy Product = convert<long long>(src0);
398403
Product *= src1;
399404
rmd = Product.format<T0>().select<SZ, 2>(0);
@@ -1964,11 +1969,11 @@ ESIMD_INLINE ESIMD_NODEBUG T0 hmin(simd<T1, SZ> v) {
19641969

19651970
template <typename T0, typename T1, int SZ, typename BinaryOperation>
19661971
ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd<T1, SZ> v, BinaryOperation op) {
1967-
if constexpr (std::is_same<remove_cvref_t<BinaryOperation>,
1972+
if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
19681973
std::plus<>>::value) {
19691974
T0 retv = detail::esimd_sum<T0>(v);
19701975
return retv;
1971-
} else if constexpr (std::is_same<remove_cvref_t<BinaryOperation>,
1976+
} else if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
19721977
std::multiplies<>>::value) {
19731978
T0 retv = detail::esimd_prod<T0>(v);
19741979
return retv;

0 commit comments

Comments
 (0)