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

[SYCL] Add tests for marray support in common and some math functions #1656

Open
wants to merge 4 commits into
base: intel
Choose a base branch
from
Open

[SYCL] Add tests for marray support in common and some math functions #1656

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
040db97
[SYCL] Add tests for marray support in common and some math functions
dm-vodopyanov Mar 13, 2023
422d748
Fix some tests
dm-vodopyanov Mar 14, 2023
48adfe4
Fix test on windows
KornevNikita Mar 15, 2023
91b44ee
Check fp64 support
KornevNikita Mar 21, 2023
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
122 changes: 122 additions & 0 deletions SYCL/DeviceLib/built-ins/marray_common.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,122 @@
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

#ifdef _WIN32
#define _USE_MATH_DEFINES // To use math constants
#include <cmath>
#endif

#include <CL/sycl.hpp>

#define TEST(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...) \
{ \
{ \
MARRAY_ELEM_TYPE result[DIM]; \
{ \
sycl::buffer<MARRAY_ELEM_TYPE> b(result, sycl::range{DIM}); \
deviceQueue.submit([&](sycl::handler &cgh) { \
sycl::accessor res_access{b, cgh}; \
cgh.single_task([=]() { \
sycl::marray<MARRAY_ELEM_TYPE, DIM> res = FUNC(__VA_ARGS__); \
for (int i = 0; i < DIM; i++) \
res_access[i] = res[i]; \
}); \
}); \
} \
for (int i = 0; i < DIM; i++) \
assert(abs(result[i] - EXPECTED[i]) <= DELTA); \
} \
}

#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__})

int main() {
sycl::queue deviceQueue;
sycl::device dev = deviceQueue.get_device();

sycl::marray<float, 2> ma1{1.0f, 2.0f};
sycl::marray<float, 2> ma2{1.0f, 2.0f};
sycl::marray<float, 2> ma3{3.0f, 2.0f};
sycl::marray<double, 2> ma4{1.0, 2.0};
sycl::marray<float, 3> ma5{M_PI, M_PI, M_PI};
sycl::marray<double, 3> ma6{M_PI, M_PI, M_PI};
sycl::marray<sycl::half, 3> ma7{M_PI, M_PI, M_PI};
sycl::marray<float, 2> ma8{0.3f, 0.6f};
sycl::marray<double, 2> ma9{5.0, 8.0};
sycl::marray<float, 3> ma10{180, 180, 180};
sycl::marray<double, 3> ma11{180, 180, 180};
sycl::marray<sycl::half, 3> ma12{180, 180, 180};
sycl::marray<sycl::half, 3> ma13{181, 179, 181};
sycl::marray<float, 2> ma14{+0.0f, -0.6f};
sycl::marray<double, 2> ma15{-0.0, 0.6f};

// sycl::clamp
TEST(sycl::clamp, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma2, ma3);
TEST(sycl::clamp, float, 2, EXPECTED(float, 2.0f, 2.0f), 0, ma1, 3.0f, 2.0f);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::clamp, double, 2, EXPECTED(double, 2.0, 2.0), 0, ma4, 3.0, 2.0);
// sycl::degrees
TEST(sycl::degrees, float, 3, EXPECTED(float, 180, 180, 180), 0, ma5);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::degrees, double, 3, EXPECTED(double, 180, 180, 180), 0, ma6);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::degrees, sycl::half, 3, EXPECTED(sycl::half, 180, 180, 180), 0.2,
ma7);
// sycl::max
TEST(sycl::max, float, 2, EXPECTED(float, 3.0f, 2.0f), 0, ma1, ma3);
TEST(sycl::max, float, 2, EXPECTED(float, 1.5f, 2.0f), 0, ma1, 1.5f);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::max, double, 2, EXPECTED(double, 1.5, 2.0), 0, ma4, 1.5);
// sycl::min
TEST(sycl::min, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma3);
TEST(sycl::min, float, 2, EXPECTED(float, 1.0f, 1.5f), 0, ma1, 1.5f);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::min, double, 2, EXPECTED(double, 1.0, 1.5), 0, ma4, 1.5);
// sycl::mix
TEST(sycl::mix, float, 2, EXPECTED(float, 1.6f, 2.0f), 0, ma1, ma3, ma8);
TEST(sycl::mix, float, 2, EXPECTED(float, 1.4f, 2.0f), 0, ma1, ma3, 0.2);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::mix, double, 2, EXPECTED(double, 3.0, 5.0), 0, ma4, ma9, 0.5);
// sycl::radians
TEST(sycl::radians, float, 3, EXPECTED(float, M_PI, M_PI, M_PI), 0, ma10);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::radians, double, 3, EXPECTED(double, M_PI, M_PI, M_PI), 0, ma11);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::radians, sycl::half, 3, EXPECTED(sycl::half, M_PI, M_PI, M_PI),
0.002, ma12);
// sycl::step
TEST(sycl::step, float, 2, EXPECTED(float, 1.0f, 1.0f), 0, ma1, ma3);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::step, double, 2, EXPECTED(double, 1.0, 1.0), 0, ma4, ma9);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::step, sycl::half, 3, EXPECTED(sycl::half, 1.0, 0.0, 1.0), 0,
ma12, ma13);
TEST(sycl::step, float, 2, EXPECTED(float, 1.0f, 0.0f), 0, 2.5f, ma3);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::step, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f, ma9);
// sycl::smoothstep
TEST(sycl::smoothstep, float, 2, EXPECTED(float, 1.0f, 0.0f), 0, ma1, ma2,
ma3);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::smoothstep, double, 2, EXPECTED(double, 0.00147576, 0.00446826),
0.00000001, ma4, ma11, ma9);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::smoothstep, sycl::half, 3, EXPECTED(sycl::half, 1.0, 0.0, 1.0),
0, ma12, ma12, ma13);
TEST(sycl::smoothstep, float, 2, EXPECTED(float, 0.0553936f, 0.0f), 0.0000001,
2.5f, 6.0f, ma3);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::smoothstep, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f,
8.0f, ma9);
// sign
TEST(sycl::sign, float, 2, EXPECTED(float, +0.0f, -1.0f), 0, ma14);
if (dev.has(sycl::aspect::fp64))
TEST(sycl::sign, double, 2, EXPECTED(double, -0.0, 1.0), 0, ma15);
if (dev.has(sycl::aspect::fp16))
TEST(sycl::sign, sycl::half, 3, EXPECTED(sycl::half, 1.0, 1.0, 1.0), 0,
ma12);

return 0;
}
77 changes: 77 additions & 0 deletions SYCL/DeviceLib/built-ins/marray_math.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,77 @@
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
// RUN: %CPU_RUN_PLACEHOLDER %t.out
// RUN: %GPU_RUN_PLACEHOLDER %t.out
// RUN: %ACC_RUN_PLACEHOLDER %t.out

#include <CL/sycl.hpp>

#define TEST(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...) \
{ \
{ \
MARRAY_ELEM_TYPE result[DIM]; \
{ \
sycl::buffer<MARRAY_ELEM_TYPE> b(result, sycl::range{DIM}); \
deviceQueue.submit([&](sycl::handler &cgh) { \
sycl::accessor res_access{b, cgh}; \
cgh.single_task([=]() { \
sycl::marray<MARRAY_ELEM_TYPE, DIM> res = FUNC(__VA_ARGS__); \
for (int i = 0; i < DIM; i++) \
res_access[i] = res[i]; \
}); \
}); \
} \
for (int i = 0; i < DIM; i++) \
assert(abs(result[i] - EXPECTED[i]) <= DELTA); \
} \
}

#define TEST2(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...) \
{ \
{ \
int result[DIM]; \
{ \
sycl::buffer<int> b(result, sycl::range{DIM}); \
deviceQueue.submit([&](sycl::handler &cgh) { \
sycl::accessor res_access{b, cgh}; \
cgh.single_task([=]() { \
sycl::marray<int, DIM> res = FUNC(__VA_ARGS__); \
for (int i = 0; i < DIM; i++) \
res_access[i] = res[i]; \
}); \
}); \
} \
for (int i = 0; i < DIM; i++) \
assert(std::abs(result[i] - EXPECTED[i]) <= DELTA); \
} \
}

#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__})

int main() {
sycl::queue deviceQueue;

sycl::marray<float, 2> ma1{1.0f, 2.0f};
sycl::marray<float, 2> ma2{3.0f, 2.0f};
sycl::marray<float, 3> ma3{180, 180, 180};
sycl::marray<int, 3> ma4{1, 1, 1};
sycl::marray<float, 3> ma5{180, -180, -180};

TEST(sycl::fabs, float, 3, EXPECTED(float, 180, 180, 180), 0, ma5);
TEST2(sycl::ilogb, float, 3, EXPECTED(int, 7, 7, 7), 0, ma3);
Copy link

Choose a reason for hiding this comment

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

Unless I'm missing something, this seems simpler:

Suggested change
TEST2(sycl::ilogb, float, 3, EXPECTED(int, 7, 7, 7), 0, ma3);
TEST(sycl::ilogb, int, 3, EXPECTED(int, 7, 7, 7), 0, ma3);

Then we can remove TEST2. Alternatively TEST2 could be defined using TEST to avoid code duplication.

Copy link
Author

Choose a reason for hiding this comment

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

ilogb takes floats and returns ints, so this is not the usual case where function takes T type and returns T type, so we can't do this change. However, TEST2 can be merged to TEST by adding one new argument to the TEST - the type of return value

TEST(sycl::fmax, float, 2, EXPECTED(float, 3.0f, 2.0f), 0, ma1, ma2);
TEST(sycl::fmax, float, 2, EXPECTED(float, 5.0f, 5.0f), 0, ma1, 5.0f);
TEST(sycl::fmin, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma2);
TEST(sycl::fmin, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, 5.0f);
TEST(sycl::ldexp, float, 3, EXPECTED(float, 360, 360, 360), 0, ma3, ma4);
TEST(sycl::ldexp, float, 3, EXPECTED(float, 5760, 5760, 5760), 0, ma3, 5);
TEST(sycl::pown, float, 3, EXPECTED(float, 180, 180, 180), 0.1, ma3, ma4);
TEST(sycl::pown, float, 3, EXPECTED(float, 180, 180, 180), 0.1, ma3, 1);
TEST(sycl::rootn, float, 3, EXPECTED(float, 180, 180, 180), 0.1, ma3, ma4);
TEST(sycl::rootn, float, 3, EXPECTED(float, 2.82523, 2.82523, 2.82523),
0.00001, ma3, 5);

TEST(sycl::half_precision::exp10, float, 2, EXPECTED(float, 10, 100), 0.1,
ma1);

return 0;
}