Skip to content

[ESIMD] Fix leakage of detail namespace #3673

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
May 2, 2021
Merged

[ESIMD] Fix leakage of detail namespace #3673

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
62 changes: 31 additions & 31 deletions sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,8 +18,6 @@ namespace sycl {
namespace INTEL {
namespace gpu {

using namespace sycl::INTEL::gpu::detail;

/// The simd vector class.
///
/// This is a wrapper class for llvm vector values. Additionally this class
Expand Down Expand Up @@ -51,13 +49,13 @@ template <typename Ty, int N> class simd {
if constexpr (std::is_same<SrcTy, Ty>::value)
set(other.data());
else
set(__builtin_convertvector(other.data(), vector_type_t<Ty, N>));
set(__builtin_convertvector(other.data(), detail::vector_type_t<Ty, N>));
}
template <typename SrcTy> constexpr simd(simd<SrcTy, N> &&other) {
if constexpr (std::is_same<SrcTy, Ty>::value)
set(other.data());
else
set(__builtin_convertvector(other.data(), vector_type_t<Ty, N>));
set(__builtin_convertvector(other.data(), detail::vector_type_t<Ty, N>));
}
constexpr simd(const vector_type &Val) { set(Val); }

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

/// View this simd object in a different element type.
template <typename EltTy> auto format() & {
using TopRegionTy = compute_format_type_t<simd, EltTy>;
using TopRegionTy = detail::compute_format_type_t<simd, EltTy>;
using RetTy = simd_view<simd, TopRegionTy>;
TopRegionTy R(0);
return RetTy{*this, R};
Expand All @@ -144,7 +142,8 @@ template <typename Ty, int N> class simd {
//
/// View as a 2-dimensional simd_view.
template <typename EltTy, int Height, int Width> auto format() & {
using TopRegionTy = compute_format_type_2d_t<simd, EltTy, Height, Width>;
using TopRegionTy =
detail::compute_format_type_2d_t<simd, EltTy, Height, Width>;
using RetTy = simd_view<simd, TopRegionTy>;
TopRegionTy R(0, 0);
return RetTy{*this, R};
Expand Down Expand Up @@ -190,7 +189,7 @@ template <typename Ty, int N> class simd {
/// Read multiple elements by their indices in vector
template <int Size>
simd<Ty, Size> iselect(const simd<uint16_t, Size> &Indices) {
vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(Ty);
detail::vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(Ty);
return __esimd_rdindirect<Ty, N, Size>(data(), Offsets);
}
// TODO ESIMD_EXPERIMENTAL
Expand All @@ -205,7 +204,7 @@ template <typename Ty, int N> class simd {
template <int Size>
void iupdate(const simd<uint16_t, Size> &Indices, const simd<Ty, Size> &Val,
mask_type_t<Size> Mask) {
vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(Ty);
detail::vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(Ty);
set(__esimd_wrindirect<Ty, N, Size>(data(), Val.data(), Offsets, Mask));
}

Expand All @@ -217,18 +216,18 @@ template <typename Ty, int N> class simd {
// * if not different, then auto should not be used
#define DEF_BINOP(BINOP, OPASSIGN) \
ESIMD_INLINE friend auto operator BINOP(const simd &X, const simd &Y) { \
using ComputeTy = compute_type_t<simd>; \
auto V0 = convert<typename ComputeTy::vector_type>(X.data()); \
auto V1 = convert<typename ComputeTy::vector_type>(Y.data()); \
using ComputeTy = detail::compute_type_t<simd>; \
auto V0 = detail::convert<typename ComputeTy::vector_type>(X.data()); \
auto V1 = detail::convert<typename ComputeTy::vector_type>(Y.data()); \
auto V2 = V0 BINOP V1; \
return ComputeTy(V2); \
} \
ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const simd &RHS) { \
using ComputeTy = compute_type_t<simd>; \
auto V0 = convert<typename ComputeTy::vector_type>(LHS.data()); \
auto V1 = convert<typename ComputeTy::vector_type>(RHS.data()); \
using ComputeTy = detail::compute_type_t<simd>; \
auto V0 = detail::convert<typename ComputeTy::vector_type>(LHS.data()); \
auto V1 = detail::convert<typename ComputeTy::vector_type>(RHS.data()); \
auto V2 = V0 BINOP V1; \
LHS.write(convert<vector_type>(V2)); \
LHS.write(detail::convert<vector_type>(V2)); \
return LHS; \
} \
ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const Ty &RHS) { \
Expand All @@ -255,7 +254,7 @@ template <typename Ty, int N> class simd {
const simd &Y) { \
auto R = X.data() RELOP Y.data(); \
mask_type_t<N> M(1); \
return M & convert<mask_type_t<N>>(R); \
return M & detail::convert<mask_type_t<N>>(R); \
}

DEF_RELOP(>)
Expand All @@ -276,7 +275,7 @@ template <typename Ty, int N> class simd {
ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const simd &RHS) { \
static_assert(std::is_integral<Ty>(), "not integeral type"); \
auto V2 = LHS.data() BITWISE_OP RHS.data(); \
LHS.write(convert<vector_type>(V2)); \
LHS.write(detail::convert<vector_type>(V2)); \
return LHS; \
} \
ESIMD_INLINE friend simd &operator OPASSIGN(simd &LHS, const Ty &RHS) { \
Expand Down Expand Up @@ -401,17 +400,18 @@ template <typename Ty, int N> class simd {

/// Write a simd-vector into a basic region of a simd object.
template <typename RTy>
ESIMD_INLINE void writeRegion(
RTy Region,
const vector_type_t<typename RTy::element_type, RTy::length> &Val) {
ESIMD_INLINE void
writeRegion(RTy Region,
const detail::vector_type_t<typename RTy::element_type,
RTy::length> &Val) {
using ElemTy = typename RTy::element_type;
if constexpr (N * sizeof(Ty) == RTy::length * sizeof(ElemTy))
// update the entire vector
set(bitcast<Ty, ElemTy, RTy::length>(Val));
set(detail::bitcast<Ty, ElemTy, RTy::length>(Val));
else {
static_assert(!RTy::Is_2D);
// If element type differs, do bitcast conversion first.
auto Base = bitcast<ElemTy, Ty, N>(data());
auto Base = detail::bitcast<ElemTy, Ty, N>(data());
constexpr int BN = (N * sizeof(Ty)) / sizeof(ElemTy);
// Access the region information.
constexpr int M = RTy::Size_x;
Expand All @@ -422,28 +422,28 @@ template <typename Ty, int N> class simd {
auto Merged = __esimd_wrregion<ElemTy, BN, M,
/*VS*/ 0, M, Stride>(Base, Val, Offset);
// Convert back to the original element type, if needed.
set(bitcast<Ty, ElemTy, BN>(Merged));
set(detail::bitcast<Ty, ElemTy, BN>(Merged));
}
}

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

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

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

template <typename U, typename T, int n>
ESIMD_INLINE simd<U, n> convert(simd<T, n> val) {
return __builtin_convertvector(val.data(), vector_type_t<U, n>);
return __builtin_convertvector(val.data(), detail::vector_type_t<U, n>);
}

} // namespace gpu
Expand Down
43 changes: 24 additions & 19 deletions sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -113,7 +113,8 @@ ESIMD_NODEBUG ESIMD_INLINE
std::is_integral<U>::value,
simd<T0, SZ>>
esimd_shl(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
typedef
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;

Expand Down Expand Up @@ -151,7 +152,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
std::is_integral<T1>::value && std::is_integral<T2>::value,
typename sycl::detail::remove_const_t<T0>>
esimd_shl(T1 src0, T2 src1, int flag = GENX_NOSAT) {
typedef typename computation_type<T1, T2>::type ComputationTy;
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
simd<T0, 1> Result = esimd_shl<T0>(Src0, Src1, flag);
Expand All @@ -166,7 +167,8 @@ ESIMD_NODEBUG ESIMD_INLINE
std::is_integral<U>::value,
simd<T0, SZ>>
esimd_shr(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
typedef
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
typename detail::simd_type<ComputationTy>::type Result =
Expand All @@ -185,7 +187,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
std::is_integral<T1>::value && std::is_integral<T2>::value,
typename sycl::detail::remove_const_t<T0>>
esimd_shr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
typedef typename computation_type<T1, T2>::type ComputationTy;
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
simd<T0, 1> Result = esimd_shr<T0>(Src0, Src1, flag);
Expand All @@ -207,7 +209,8 @@ ESIMD_NODEBUG ESIMD_INLINE
std::is_integral<U>::value,
simd<T0, SZ>>
esimd_rol(simd<T1, SZ> src0, U src1) {
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
typedef
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
return __esimd_rol<T0>(Src0.data(), Src1.data());
Expand All @@ -220,7 +223,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
std::is_integral<T1>::value && std::is_integral<T2>::value,
typename sycl::detail::remove_const_t<T0>>
esimd_rol(T1 src0, T2 src1) {
typedef typename computation_type<T1, T2>::type ComputationTy;
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
simd<T0, 1> Result = esimd_rol<T0>(Src0, Src1);
Expand All @@ -242,7 +245,8 @@ ESIMD_NODEBUG ESIMD_INLINE
std::is_integral<U>::value,
simd<T0, SZ>>
esimd_ror(simd<T1, SZ> src0, U src1) {
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
typedef
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
return __esimd_ror<T0>(Src0.data(), Src1.data());
Expand All @@ -255,7 +259,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
std::is_integral<T1>::value && std::is_integral<T2>::value,
typename sycl::detail::remove_const_t<T0>>
esimd_ror(T1 src0, T2 src1) {
typedef typename computation_type<T1, T2>::type ComputationTy;
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
simd<T0, 1> Result = esimd_ror<T0>(Src0, Src1);
Expand All @@ -270,7 +274,7 @@ ESIMD_NODEBUG ESIMD_INLINE
std::is_integral<U>::value,
simd<T0, SZ>>
esimd_lsr(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
typedef typename computation_type<T1, T1>::type IntermedTy;
typedef typename detail::computation_type<T1, T1>::type IntermedTy;
typedef typename std::make_unsigned<IntermedTy>::type ComputationTy;
simd<ComputationTy, SZ> Src0 = src0;
simd<ComputationTy, SZ> Result = Src0.data() >> src1.data();
Expand All @@ -288,7 +292,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
std::is_integral<T1>::value && std::is_integral<T2>::value,
typename sycl::detail::remove_const_t<T0>>
esimd_lsr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
typedef typename computation_type<T1, T2>::type ComputationTy;
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
simd<T0, 1> Result = esimd_lsr<T0>(Src0, Src1, flag);
Expand All @@ -313,7 +317,7 @@ ESIMD_NODEBUG ESIMD_INLINE
std::is_integral<U>::value,
simd<T0, SZ>>
esimd_asr(simd<T1, SZ> src0, U src1, int flag = GENX_NOSAT) {
typedef typename computation_type<T1, T1>::type IntermedTy;
typedef typename detail::computation_type<T1, T1>::type IntermedTy;
typedef typename std::make_signed<IntermedTy>::type ComputationTy;
simd<ComputationTy, SZ> Src0 = src0;
simd<ComputationTy, SZ> Result = Src0 >> src1;
Expand All @@ -331,7 +335,7 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
std::is_integral<T1>::value && std::is_integral<T2>::value,
typename sycl::detail::remove_const_t<T0>>
esimd_asr(T1 src0, T2 src1, int flag = GENX_NOSAT) {
typedef typename computation_type<T1, T2>::type ComputationTy;
typedef typename detail::computation_type<T1, T2>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
simd<T0, 1> Result = esimd_asr<T0>(Src0, Src1, flag);
Expand All @@ -358,7 +362,8 @@ ESIMD_NODEBUG ESIMD_INLINE
detail::is_dword_type<U>::value,
simd<T0, SZ>>
esimd_imul(simd<T0, SZ> &rmd, simd<T1, SZ> src0, U src1) {
typedef typename computation_type<decltype(src0), U>::type ComputationTy;
typedef
typename detail::computation_type<decltype(src0), U>::type ComputationTy;
typename detail::simd_type<ComputationTy>::type Src0 = src0;
typename detail::simd_type<ComputationTy>::type Src1 = src1;
rmd = Src0 * Src1;
Expand All @@ -378,8 +383,8 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
detail::is_dword_type<U>::value && SZ == 1,
simd<T0, SZ>>
esimd_imul(simd<T0, SZ> &rmd, simd<T1, SZ> src0, U src1) {
typedef
typename computation_type<decltype(rmd), long long>::type ComputationTy;
typedef typename detail::computation_type<decltype(rmd), long long>::type
ComputationTy;
ComputationTy Product = convert<long long>(src0);
Product *= src1;
rmd = Product.format<T0>().select<1, 1>[0];
Expand All @@ -392,8 +397,8 @@ ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
detail::is_dword_type<U>::value && SZ != 1,
simd<T0, SZ>>
esimd_imul(simd<T0, SZ> &rmd, simd<T1, SZ> src0, U src1) {
typedef
typename computation_type<decltype(rmd), long long>::type ComputationTy;
typedef typename detail::computation_type<decltype(rmd), long long>::type
ComputationTy;
ComputationTy Product = convert<long long>(src0);
Product *= src1;
rmd = Product.format<T0>().select<SZ, 2>(0);
Expand Down Expand Up @@ -1964,11 +1969,11 @@ ESIMD_INLINE ESIMD_NODEBUG T0 hmin(simd<T1, SZ> v) {

template <typename T0, typename T1, int SZ, typename BinaryOperation>
ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd<T1, SZ> v, BinaryOperation op) {
if constexpr (std::is_same<remove_cvref_t<BinaryOperation>,
if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
std::plus<>>::value) {
T0 retv = detail::esimd_sum<T0>(v);
return retv;
} else if constexpr (std::is_same<remove_cvref_t<BinaryOperation>,
} else if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
std::multiplies<>>::value) {
T0 retv = detail::esimd_prod<T0>(v);
return retv;
Expand Down
Loading