Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
519 changes: 519 additions & 0 deletions include/sycl_ext_complex.hpp

Large diffs are not rendered by default.

2 changes: 1 addition & 1 deletion tests/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@ file(GLOB test_cases CONFIGURE_DEPENDS "*.cpp")
foreach(test_file IN LISTS test_cases)
if(EXISTS "${test_file}")
get_filename_component(exe_name "${test_file}" NAME_WE)

add_executable(${exe_name} ${test_file})
target_include_directories(${exe_name} PUBLIC ../include/)
target_link_libraries(${exe_name} PRIVATE
Expand Down
2 changes: 1 addition & 1 deletion tests/asin_complex.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -49,7 +49,7 @@ TEMPLATE_TEST_CASE("Test complex asin", "[asin]", double, float, sycl::half) {
}
Q.copy(d_cplx_out, &h_cplx_out, 1).wait();

check_results(h_cplx_out, std_out, /*tol_multiplier*/ 5);
check_results(h_cplx_out, std_out, /*tol_multiplier*/ 6);
}

// Check cplx::complex output from host
Expand Down
191 changes: 191 additions & 0 deletions tests/exclusive_scan_over_group.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,191 @@
#include <array>
#include <numeric>

#include "test_helper.hpp"

////////////////////////////////////////////////////////////////////////////////
// UTILITY FUNCTIONS
////////////////////////////////////////////////////////////////////////////////

template <typename T, typename BinaryOperation>
void test_exclusive_scan_over_group(sycl::queue q, T input,
BinaryOperation binary_op) {
using V = typename T::value_type;

constexpr size_t N = input.size();

auto init = sycl::ext::cplx::cplex::detail::get_init<V, BinaryOperation>();

auto *in = sycl::malloc_shared<V>(N, q);
auto *output_with_init = sycl::malloc_shared<V>(N, q);
auto *output_without_init = sycl::malloc_shared<V>(N, q);

for (std::size_t i = 0; i < N; i++) {
in[i] = input[i];
}

q.submit([&](sycl::handler &cgh) {
cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) {
auto gid = it.get_global_id(0);
auto lid = it.get_local_id(0);
auto g = it.get_group();

output_with_init[lid] = sycl::ext::cplx::exclusive_scan_over_group(
g, in[gid], init, binary_op);
output_without_init[lid] =
sycl::ext::cplx::exclusive_scan_over_group(g, in[gid], binary_op);
});
});

q.wait();

std::array<V, N> expected;
std::exclusive_scan(input.begin(), input.end(), expected.begin(), init,
binary_op);

std::array<V, N> result;
for (std::size_t i = 0; i < N; i++) {
result[i] = output_with_init[i];
}

check_results(result, expected);

for (std::size_t i = 0; i < N; i++) {
result[i] = output_without_init[i];
}

check_results(result, expected);

sycl::free(in, q);
sycl::free(output_with_init, q);
sycl::free(output_without_init, q);
}

////////////////////////////////////////////////////////////////////////////////
// COMPLEX TESTS
////////////////////////////////////////////////////////////////////////////////

TEMPLATE_TEST_CASE_SIG("Test complex exclusive_scan_over_group", "[scan]",
((typename T, std::size_t N, typename BinaryOperation),
T, N, BinaryOperation),
(double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
(sycl::half, 4, sycl::plus<>)) {

using Complex = typename sycl::ext::cplx::complex<T>;
using Array = typename std::array<Complex, N>;

sycl::queue q;

const auto test_cases = GENERATE(
// Basic value test
Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}},
// Random value test
Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
Complex{3.7, 3.7}},
// Repeated value test
Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
// Negative value test
Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
Complex{0, 0}},
// Large value test
Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
// Small value test
Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
// Edge case value test
Array{Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}});
const auto binary_op = BinaryOperation{};

if (is_type_supported<T>(q)) {
test_exclusive_scan_over_group(q, test_cases, binary_op);
}
}

////////////////////////////////////////////////////////////////////////////////
// MARRAY<COMPLEX> TESTS
////////////////////////////////////////////////////////////////////////////////

TEMPLATE_TEST_CASE_SIG("Test marray<complex> exclusive_scan_over_group",
"[scan]",
((typename T, std::size_t N, typename BinaryOperation),
T, N, BinaryOperation),
(double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
(sycl::half, 4, sycl::plus<>)) {

using Complex = typename sycl::ext::cplx::complex<T>;
using Marray = typename sycl::marray<Complex, N>;
using Array = typename std::array<Marray, N>;

sycl::queue q;

const auto test_cases = GENERATE(
// Basic value test
Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}},
Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}},
Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}},
Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}},
// Random value test
Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
Complex{3.7, 3.7}},
Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
Complex{3.7, 3.7}},
Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
Complex{3.7, 3.7}},
Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
Complex{3.7, 3.7}}},
// Repeated value test
Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}},
// Negative value test
Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
Complex{0, 0}},
Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
Complex{0, 0}},
Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
Complex{0, 0}},
Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
Complex{0, 0}}},
// Large value test
Array{
Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}},
// Small value test
Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}},
// Edge case value test
Array{
Marray{
Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
Marray{
Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
Marray{
Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
Marray{Complex{nan_val<T>, nan_val<T>},
Complex{inf_val<T>, inf_val<T>},
Complex{nan_val<T>, inf_val<T>},
Complex{inf_val<T>, nan_val<T>}}});
const auto binary_op = BinaryOperation{};

if (is_type_supported<T>(q)) {
test_exclusive_scan_over_group(q, test_cases, binary_op);
}
}
Loading