Skip to content

Commit bef8a98

Browse files
sarnexsarnex
authored
[SYCL][ESIMD] Revert some SPIR-V intrinsic changes (#17301)
These require a new driver driver than the minimum supported version by customers, so we need to revert. Signed-off-by: Sarnie, Nick <[email protected]>
1 parent 607dff4 commit bef8a98

File tree

10 files changed

+79
-118
lines changed

10 files changed

+79
-118
lines changed

‎llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp‎

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1245,6 +1245,21 @@ static Instruction *addCastInstIfNeeded(Instruction *OldI, Instruction *NewI,
12451245
return NewI;
12461246
}
12471247

1248+
// Translates the following intrinsics:
1249+
// %res = call float @llvm.fmuladd.f32(float %a, float %b, float %c)
1250+
// %res = call double @llvm.fmuladd.f64(double %a, double %b, double %c)
1251+
// To
1252+
// %mul = fmul <type> %a, <type> %b
1253+
// %res = fadd <type> %mul, <type> %c
1254+
// TODO: Remove when newer GPU driver is used in CI.
1255+
void translateFmuladd(CallInst *CI) {
1256+
assert(CI->getIntrinsicID() == Intrinsic::fmuladd);
1257+
IRBuilder<> Bld(CI);
1258+
auto *Mul = Bld.CreateFMul(CI->getOperand(0), CI->getOperand(1));
1259+
auto *Res = Bld.CreateFAdd(Mul, CI->getOperand(2));
1260+
CI->replaceAllUsesWith(Res);
1261+
}
1262+
12481263
// Translates an LLVM intrinsic to a form, digestable by the BE.
12491264
bool translateLLVMIntrinsic(CallInst *CI) {
12501265
Function *F = CI->getCalledFunction();
@@ -1256,6 +1271,9 @@ bool translateLLVMIntrinsic(CallInst *CI) {
12561271
// no translation - it will be simply removed.
12571272
// TODO: make use of 'assume' info in the BE
12581273
break;
1274+
case Intrinsic::fmuladd:
1275+
translateFmuladd(CI);
1276+
break;
12591277
default:
12601278
return false; // "intrinsic wasn't translated, keep the original call"
12611279
}

‎llvm/test/SYCLLowerIR/ESIMD/lower_llvm_intrin.ll‎

Lines changed: 8 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
; RUN: opt -passes=LowerESIMD -S < %s | FileCheck %s
22

3-
; This test checks that LowerESIMD pass does not lower some llvm intrinsics
4-
; which can now be handled by the VC BE.
3+
; This test checks that LowerESIMD pass correctly lowers some llvm intrinsics
4+
; which can't be handled by the VC BE.
55
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
66
target triple = "spir64-unknown-unknown"
77

@@ -10,15 +10,17 @@ declare double @llvm.fmuladd.f64(double %x, double %y, double %z)
1010

1111
define spir_func float @test_fmuladd_f32(float %x, float %y, float %z) {
1212
%1 = call float @llvm.fmuladd.f32(float %x, float %y, float %z)
13-
; CHECK: %[[A:[0-9a-zA-Z\._]+]] = call float @llvm.fmuladd.f32(float %x, float %y, float %z)
13+
; CHECK: %[[A:[0-9a-zA-Z\._]+]] = fmul float %x, %y
14+
; CHECK: %[[B:[0-9a-zA-Z\._]+]] = fadd float %[[A]], %z
1415
ret float %1
15-
; CHECK: ret float %[[A]]
16+
; CHECK: ret float %[[B]]
1617
}
1718

1819
define spir_func double @test_fmuladd_f64(double %x, double %y, double %z) {
1920
%1 = call double @llvm.fmuladd.f64(double %x, double %y, double %z)
20-
; CHECK: %[[A:[0-9a-zA-Z\._]+]] = call double @llvm.fmuladd.f64(double %x, double %y, double %z)
21+
; CHECK: %[[A:[0-9a-zA-Z\._]+]] = fmul double %x, %y
22+
; CHECK: %[[B:[0-9a-zA-Z\._]+]] = fadd double %[[A]], %z
2123
ret double %1
22-
; CHECK: ret double %[[A]]
24+
; CHECK: ret double %[[B]]
2325
}
2426

‎sycl/include/sycl/ext/intel/esimd/detail/intrin.hpp‎

Lines changed: 0 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -13,13 +13,6 @@
1313

1414
/// @cond ESIMD_DETAIL
1515

16-
/// **************************** WARNING ************************************
17-
/// When declaring new SPIR-V intrinsics (functions starting with __spirv),
18-
/// it is imperitive to exactly follow the pattern of the existing SPIR-V
19-
/// intrinsics. If not followed, the declaration may conflict with
20-
/// the Clang-generated functions and cause compilation errors.
21-
/// **************************** WARNING ************************************
22-
2316
#include <sycl/ext/intel/esimd/common.hpp>
2417
#include <sycl/ext/intel/esimd/detail/types.hpp>
2518
#include <sycl/ext/intel/esimd/detail/util.hpp>

‎sycl/include/sycl/ext/intel/esimd/detail/math_intrin.hpp‎

Lines changed: 23 additions & 45 deletions
Original file line numberDiff line numberDiff line change
@@ -13,13 +13,6 @@
1313

1414
/// @cond ESIMD_DETAIL
1515

16-
/// **************************** WARNING ************************************
17-
/// When declaring new SPIR-V intrinsics (functions starting with __spirv),
18-
/// it is imperitive to exactly follow the pattern of the existing SPIR-V
19-
/// intrinsics. If not followed, the declaration may conflict with
20-
/// the Clang-generated functions and cause compilation errors.
21-
/// **************************** WARNING ************************************
22-
2316
#include <sycl/builtins.hpp>
2417
#include <sycl/ext/intel/esimd/common.hpp>
2518
#include <sycl/ext/intel/esimd/detail/elem_type_traits.hpp>
@@ -39,73 +32,45 @@
3932
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_native_exp2(T);
4033
template <typename T, int N>
4134
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
42-
__spirv_ocl_native_exp2(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
35+
__spirv_ocl_native_exp2(__ESIMD_raw_vec_t(T, N));
4336

4437
template <typename T>
4538
extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_native_recip(T);
4639
template <typename T, int N>
4740
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
48-
__spirv_ocl_native_recip(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
41+
__spirv_ocl_native_recip(__ESIMD_raw_vec_t(T, N));
4942

5043
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_native_cos(T);
5144
template <typename T, int N>
5245
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
53-
__spirv_ocl_native_cos(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
46+
__spirv_ocl_native_cos(__ESIMD_raw_vec_t(T, N));
5447

5548
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_native_log2(T);
5649
template <typename T, int N>
5750
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
58-
__spirv_ocl_native_log2(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
51+
__spirv_ocl_native_log2(__ESIMD_raw_vec_t(T, N));
5952

6053
template <typename T>
6154
extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_native_rsqrt(T);
6255
template <typename T, int N>
6356
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
64-
__spirv_ocl_native_rsqrt(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
57+
__spirv_ocl_native_rsqrt(__ESIMD_raw_vec_t(T, N));
6558

6659
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_native_sin(T);
6760
template <typename T, int N>
6861
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
69-
__spirv_ocl_native_sin(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
62+
__spirv_ocl_native_sin(__ESIMD_raw_vec_t(T, N));
7063

7164
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_native_sqrt(T);
7265
template <typename T, int N>
7366
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
74-
__spirv_ocl_native_sqrt(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
67+
__spirv_ocl_native_sqrt(__ESIMD_raw_vec_t(T, N));
7568

7669
template <typename T>
7770
extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_native_powr(T, T);
7871
template <typename T, int N>
7972
__ESIMD_INTRIN __ESIMD_raw_vec_t(T, N)
80-
__spirv_ocl_native_powr(__ESIMD_raw_vec_t(T, N),
81-
__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
82-
83-
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_fabs(T);
84-
template <typename T, int N>
85-
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
86-
__spirv_ocl_fabs(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
87-
88-
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_s_abs(T);
89-
template <typename T, int N>
90-
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
91-
__spirv_ocl_s_abs(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
92-
93-
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_fmin(T, T);
94-
template <typename T, int N>
95-
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
96-
__spirv_ocl_fmin(__ESIMD_raw_vec_t(T, N),
97-
__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
98-
99-
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_fmax(T, T);
100-
template <typename T, int N>
101-
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
102-
__spirv_ocl_fmax(__ESIMD_raw_vec_t(T, N),
103-
__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
104-
105-
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_rsqrt(T);
106-
template <typename T, int N>
107-
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
108-
__spirv_ocl_rsqrt(__ESIMD_raw_vec_t(T, N)) __ESIMD_INTRIN_END;
73+
__spirv_ocl_native_powr(__ESIMD_raw_vec_t(T, N), __ESIMD_raw_vec_t(T, N));
10974

11075
// saturation intrinsics
11176
template <typename T0, typename T1, int SZ>
@@ -136,7 +101,15 @@ template <typename T0, typename T1, int SZ>
136101
__ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
137102
__esimd_sstrunc_sat(__ESIMD_raw_vec_t(T1, SZ) src) __ESIMD_INTRIN_END;
138103

139-
/// 3 kinds of max, the missing fmax uses spir-v intrinsics above
104+
template <typename T, int SZ>
105+
__ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
106+
__esimd_abs(__ESIMD_raw_vec_t(T, SZ) src0) __ESIMD_INTRIN_END;
107+
108+
/// 3 kinds of max
109+
template <typename T, int SZ>
110+
__ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
111+
__esimd_fmax(__ESIMD_raw_vec_t(T, SZ) src0,
112+
__ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
140113
template <typename T, int SZ>
141114
__ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
142115
__esimd_umax(__ESIMD_raw_vec_t(T, SZ) src0,
@@ -146,7 +119,12 @@ __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
146119
__esimd_smax(__ESIMD_raw_vec_t(T, SZ) src0,
147120
__ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
148121

149-
/// 3 kinds of min, the missing fmin uses spir-v instrinsics above
122+
/// 3 kinds of min
123+
template <typename T, int SZ>
124+
__ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
125+
__esimd_fmin(__ESIMD_raw_vec_t(T, SZ) src0,
126+
__ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
127+
150128
template <typename T, int SZ>
151129
__ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
152130
__esimd_umin(__ESIMD_raw_vec_t(T, SZ) src0,

‎sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp‎

Lines changed: 0 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -11,13 +11,6 @@
1111

1212
/// @cond ESIMD_DETAIL
1313

14-
/// **************************** WARNING ************************************
15-
/// When declaring new SPIR-V intrinsics (functions starting with __spirv),
16-
/// it is imperitive to exactly follow the pattern of the existing SPIR-V
17-
/// intrinsics. If not followed, the declaration may conflict with
18-
/// the Clang-generated functions and cause compilation errors.
19-
/// **************************** WARNING ************************************
20-
2114
#pragma once
2215

2316
#include <sycl/accessor.hpp>

‎sycl/include/sycl/ext/intel/esimd/math.hpp‎

Lines changed: 9 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -100,11 +100,10 @@ __esimd_abs_common_internal(simd<TArg, SZ> src0) {
100100
simd<TArg, SZ> Result;
101101
if constexpr (detail::is_generic_floating_point_v<TArg>) {
102102
using CppT = __ESIMD_DNS::element_type_traits<TArg>::EnclosingCppT;
103-
Result =
104-
__ESIMD_DNS::convert_vector<TArg, CppT, SZ>(__spirv_ocl_fabs<CppT, SZ>(
105-
__ESIMD_DNS::convert_vector<CppT, TArg, SZ>(src0.data())));
103+
Result = __ESIMD_DNS::convert_vector<TArg, CppT, SZ>(__esimd_abs<CppT, SZ>(
104+
__ESIMD_DNS::convert_vector<CppT, TArg, SZ>(src0.data())));
106105
} else
107-
Result = simd<TArg, SZ>(__spirv_ocl_s_abs<TArg, SZ>(src0.data()));
106+
Result = simd<TArg, SZ>(__esimd_abs<TArg, SZ>(src0.data()));
108107
return convert<TRes>(Result);
109108
}
110109

@@ -190,7 +189,7 @@ __ESIMD_API simd<T, SZ>(max)(simd<T, SZ> src0, simd<T, SZ> src1, Sat sat = {}) {
190189
if constexpr (detail::is_generic_floating_point_v<T>) {
191190
using CppT = __ESIMD_DNS::element_type_traits<T>::EnclosingCppT;
192191
auto Result =
193-
__ESIMD_DNS::convert_vector<T, CppT, SZ>(__spirv_ocl_fmax<CppT, SZ>(
192+
__ESIMD_DNS::convert_vector<T, CppT, SZ>(__esimd_fmax<CppT, SZ>(
194193
__ESIMD_DNS::convert_vector<CppT, T, SZ>(src0.data()),
195194
__ESIMD_DNS::convert_vector<CppT, T, SZ>(src1.data())));
196195
if constexpr (is_sat)
@@ -279,7 +278,7 @@ __ESIMD_API simd<T, SZ>(min)(simd<T, SZ> src0, simd<T, SZ> src1, Sat sat = {}) {
279278
if constexpr (detail::is_generic_floating_point_v<T>) {
280279
using CppT = __ESIMD_DNS::element_type_traits<T>::EnclosingCppT;
281280
auto Result =
282-
__ESIMD_DNS::convert_vector<T, CppT, SZ>(__spirv_ocl_fmin<CppT, SZ>(
281+
__ESIMD_DNS::convert_vector<T, CppT, SZ>(__esimd_fmin<CppT, SZ>(
283282
__ESIMD_DNS::convert_vector<CppT, T, SZ>(src0.data()),
284283
__ESIMD_DNS::convert_vector<CppT, T, SZ>(src1.data())));
285284
if constexpr (is_sat)
@@ -456,24 +455,20 @@ __ESIMD_UNARY_INTRINSIC_DEF(__ESIMD_EMATH_SPIRV_COND, cos, cos)
456455
template <class T, int N, class Sat = saturation_off_tag>
457456
__ESIMD_API std::enable_if_t<std::is_same_v<T, double>, simd<double, N>>
458457
rsqrt(simd<T, N> src, Sat sat = {}) {
459-
__ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<double>, N> res =
460-
__spirv_ocl_rsqrt<__ESIMD_DNS::__raw_t<double>, N>(src.data());
461458
if constexpr (std::is_same_v<Sat, saturation_off_tag>)
462-
return res;
459+
return inv(sqrt(src));
463460
else
464-
return esimd::saturate<double>(simd<double, N>(res));
461+
return esimd::saturate<double>(inv(sqrt(src)));
465462
}
466463

467464
/** Scalar version. */
468465
template <class T, class Sat = saturation_off_tag>
469466
__ESIMD_API std::enable_if_t<std::is_same_v<T, double>, double>
470467
rsqrt(T src, Sat sat = {}) {
471-
__ESIMD_DNS::__raw_t<double> res =
472-
__spirv_ocl_rsqrt<__ESIMD_DNS::__raw_t<double>>(src);
473468
if constexpr (std::is_same_v<Sat, saturation_off_tag>)
474-
return res;
469+
return inv(sqrt(src));
475470
else
476-
return esimd::saturate<double>(simd<double, 1>(res))[0];
471+
return esimd::saturate<double>(inv(sqrt(src)));
477472
}
478473

479474
#undef __ESIMD_UNARY_INTRINSIC_DEF

‎sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp‎

Lines changed: 3 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -12,13 +12,6 @@
1212

1313
/// @cond ESIMD_DETAIL
1414

15-
/// **************************** WARNING ************************************
16-
/// When declaring new SPIR-V intrinsics (functions starting with __spirv),
17-
/// it is imperitive to exactly follow the pattern of the existing SPIR-V
18-
/// intrinsics. If not followed, the declaration may conflict with
19-
/// the Clang-generated functions and cause compilation errors.
20-
/// **************************** WARNING ************************************
21-
2215
#include <sycl/ext/intel/esimd/detail/defines_elementary.hpp>
2316
#include <sycl/ext/intel/esimd/detail/math_intrin.hpp>
2417
#include <sycl/ext/intel/esimd/detail/types.hpp>
@@ -105,26 +98,10 @@ __ESIMD_INTRIN __ESIMD_raw_vec_t(sycl::half, N)
10598
__ESIMD_DNS::vector_type_t<uint16_t, N> src2)
10699
__ESIMD_INTRIN_END;
107100

108-
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_fma(T, T, T);
109-
template <typename T, int N>
110-
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
111-
__spirv_ocl_fma(__ESIMD_raw_vec_t(T, N) a, __ESIMD_raw_vec_t(T, N) b,
112-
__ESIMD_raw_vec_t(T, N) c) __ESIMD_INTRIN_END;
113-
114-
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_popcount(T);
115-
template <typename T, int N>
116-
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
117-
__spirv_ocl_popcount(__ESIMD_raw_vec_t(T, N) src0) __ESIMD_INTRIN_END;
118-
119-
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_ctz(T);
120-
template <typename T, int N>
121-
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
122-
__spirv_ocl_ctz(__ESIMD_raw_vec_t(T, N) src0) __ESIMD_INTRIN_END;
123-
124-
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_clz(T);
125101
template <typename T, int N>
126-
extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
127-
__spirv_ocl_clz(__ESIMD_raw_vec_t(T, N) src0) __ESIMD_INTRIN_END;
102+
__ESIMD_INTRIN __ESIMD_raw_vec_t(T, N)
103+
__esimd_fmadd(__ESIMD_raw_vec_t(T, N) a, __ESIMD_raw_vec_t(T, N) b,
104+
__ESIMD_raw_vec_t(T, N) c) __ESIMD_INTRIN_END;
128105

129106
template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_FRem(T);
130107
template <typename T, int N>

‎sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp‎

Lines changed: 0 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -10,13 +10,6 @@
1010

1111
/// @cond ESIMD_DETAIL
1212

13-
/// **************************** WARNING ************************************
14-
/// When declaring new SPIR-V intrinsics (functions starting with __spirv),
15-
/// it is imperitive to exactly follow the pattern of the existing SPIR-V
16-
/// intrinsics. If not followed, the declaration may conflict with
17-
/// the Clang-generated functions and cause compilation errors.
18-
/// **************************** WARNING ************************************
19-
2013
#pragma once
2114

2215
#include <sycl/ext/intel/esimd/detail/defines_elementary.hpp>

‎sycl/include/sycl/ext/intel/experimental/esimd/math.hpp‎

Lines changed: 16 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,11 @@ template <typename T, int N>
3232
__ESIMD_API std::enable_if_t<std::is_integral_v<T> && sizeof(T) < 8,
3333
__ESIMD_NS::simd<T, N>>
3434
popcount(__ESIMD_NS::simd<T, N> vec) {
35-
return __spirv_ocl_popcount<T, N>(vec.data());
35+
#ifdef __SYCL_DEVICE_ONLY__
36+
return __spirv_ocl_popcount(vec.data());
37+
#else
38+
return vec;
39+
#endif
3640
}
3741

3842
/// Count the number of leading zeros.
@@ -44,7 +48,11 @@ template <typename T, int N>
4448
__ESIMD_API std::enable_if_t<std::is_integral_v<T> && sizeof(T) < 8,
4549
__ESIMD_NS::simd<T, N>>
4650
clz(__ESIMD_NS::simd<T, N> vec) {
47-
return __spirv_ocl_clz<T, N>(vec.data());
51+
#ifdef __SYCL_DEVICE_ONLY__
52+
return __spirv_ocl_clz(vec.data());
53+
#else
54+
return vec;
55+
#endif
4856
}
4957

5058
/// Count the number of trailing zeros.
@@ -55,7 +63,11 @@ template <typename T, int N>
5563
__ESIMD_API std::enable_if_t<std::is_integral_v<T> && sizeof(T) < 8,
5664
__ESIMD_NS::simd<T, N>>
5765
ctz(__ESIMD_NS::simd<T, N> vec) {
58-
return __spirv_ocl_ctz<T, N>(vec.data());
66+
#ifdef __SYCL_DEVICE_ONLY__
67+
return __spirv_ocl_ctz(vec.data());
68+
#else
69+
return vec;
70+
#endif
5971
}
6072

6173
/// @} sycl_esimd_bitmanip
@@ -752,7 +764,7 @@ ESIMD_INLINE __ESIMD_NS::simd<T, N> fma(__ESIMD_NS::simd<T, N> a,
752764
static_assert(__ESIMD_DNS::is_generic_floating_point_v<T>,
753765
"fma only supports floating point types");
754766
using CppT = __ESIMD_DNS::element_type_traits<T>::EnclosingCppT;
755-
auto Ret = __spirv_ocl_fma<__ESIMD_DNS::__raw_t<CppT>, N>(
767+
auto Ret = __esimd_fmadd<__ESIMD_DNS::__raw_t<CppT>, N>(
756768
__ESIMD_DNS::convert_vector<CppT, T, N>(a.data()),
757769
__ESIMD_DNS::convert_vector<CppT, T, N>(b.data()),
758770
__ESIMD_DNS::convert_vector<CppT, T, N>(c.data()));

0 commit comments

Comments
 (0)