Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

[ESIMD] Enable math functions: cos,sin,exp,log #312

Merged
merged 4 commits into from
Jul 20, 2021
Merged

[ESIMD] Enable math functions: cos,sin,exp,log #312

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
b428b17
[ESIMD] Enable math functions: cos,sin,exp,log
DenisBakhvalov Jun 4, 2021
8bb0992
Updated according to code review comments
DenisBakhvalov Jun 25, 2021
ac0fd45
Clang-frmat
DenisBakhvalov Jun 25, 2021
09293af
Add XFAIL on windows
DenisBakhvalov Jul 19, 2021
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
44 changes: 31 additions & 13 deletions SYCL/ESIMD/ext_math.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,13 +9,16 @@
// UNSUPPORTED: cuda
// RUN: %clangxx -fsycl %s -o %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// Enable when driver fixes will be propagated into the official release
// XFAIL: windows

// This test checks extended math operations.

#include "esimd_test_utils.hpp"

#include <CL/sycl.hpp>
#include <CL/sycl/INTEL/esimd.hpp>
#include <CL/sycl/builtins_esimd.hpp>
#include <iostream>

using namespace cl::sycl;
Expand All @@ -35,7 +38,16 @@ struct InitDataFuncWide {
struct InitDataFuncNarrow {
void operator()(float *In, float *Out, size_t Size) const {
for (auto I = 0; I < Size; ++I) {
In[I] = 2.0f + 16.0f * ((float)I / (float)(Size - 1)); // in [2..16] range
In[I] = 2.0f + 16.0f * ((float)I / (float)(Size - 1)); // in [2..18] range
Out[I] = (float)0.0;
}
}
};

struct InitDataInRange0_5 {
void operator()(float *In, float *Out, size_t Size) const {
for (auto I = 0; I < Size; ++I) {
In[I] = 5.0f * ((float)I / (float)(Size - 1)); // in [0..5] range
Out[I] = (float)0.0;
}
}
Expand All @@ -52,7 +64,7 @@ template <MathOp Op> float HostMathFunc(float X);

// --- Specializations per each extended math operation

#define DEFINE_OP(Op, HostOp) \
#define DEFINE_ESIMD_OP(Op, HostOp) \
template <> float HostMathFunc<MathOp::Op>(float X) { return HostOp(X); } \
template <int VL> struct DeviceMathFunc<VL, MathOp::Op> { \
simd<float, VL> \
Expand All @@ -61,13 +73,22 @@ template <MathOp Op> float HostMathFunc(float X);
} \
}

DEFINE_OP(sin, sin);
DEFINE_OP(cos, cos);
DEFINE_OP(exp, exp);
DEFINE_OP(log, log);
DEFINE_OP(inv, 1.0f /);
DEFINE_OP(sqrt, sqrt);
DEFINE_OP(rsqrt, 1.0f / sqrt);
#define DEFINE_SIMD_OVERLOADED_STD_SYCL_OP(Op, HostOp) \
template <> float HostMathFunc<MathOp::Op>(float X) { return HostOp(X); } \
template <int VL> struct DeviceMathFunc<VL, MathOp::Op> { \
simd<float, VL> \
operator()(const simd<float, VL> &X) const SYCL_ESIMD_FUNCTION { \
return sycl::Op<VL>(X); \
} \
}

DEFINE_SIMD_OVERLOADED_STD_SYCL_OP(sin, sin);
DEFINE_SIMD_OVERLOADED_STD_SYCL_OP(cos, cos);
DEFINE_SIMD_OVERLOADED_STD_SYCL_OP(exp, exp);
DEFINE_SIMD_OVERLOADED_STD_SYCL_OP(log, log);
DEFINE_ESIMD_OP(inv, 1.0f /);
DEFINE_ESIMD_OP(sqrt, sqrt);
DEFINE_ESIMD_OP(rsqrt, 1.0f / sqrt);

// --- Generic kernel calculating an extended math operation on array elements

Expand Down Expand Up @@ -159,13 +180,10 @@ template <int VL> bool test(queue &Q) {
Pass &= test<MathOp::sqrt, VL>(Q, "sqrt", InitDataFuncWide{});
Pass &= test<MathOp::inv, VL>(Q, "inv");
Pass &= test<MathOp::rsqrt, VL>(Q, "rsqrt");
// TODO enable these tests after the implementation is fixed
#if ENABLE_SIN_COS_EXP_LOG
Pass &= test<MathOp::sin, VL>(Q, "sin", InitDataFuncWide{});
Pass &= test<MathOp::cos, VL>(Q, "cos", InitDataFuncWide{});
Pass &= test<MathOp::exp, VL>(Q, "exp");
Pass &= test<MathOp::exp, VL>(Q, "exp", InitDataInRange0_5{});
Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The precision of FP operations gets lower towards bigger numbers, that's why we use the smaller value range.

Pass &= test<MathOp::log, VL>(Q, "log", InitDataFuncWide{});
#endif
return Pass;
}

Expand Down