Skip to content

[SYCL][ESIMD] Host-compile simd.cpp test, fix errors & warnings. #3846

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
Jun 2, 2021
Merged

[SYCL][ESIMD] Host-compile simd.cpp test, fix errors & warnings. #3846

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
3 changes: 2 additions & 1 deletion sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
//==-------------- math.hpp - DPC++ Explicit SIMD API --------------------==//
//==-------------- math.hpp - DPC++ Explicit SIMD API --------------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
Expand All @@ -11,6 +11,7 @@
#pragma once

#include <sycl/ext/intel/experimental/esimd/common.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/host_util.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/types.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/util.hpp>
Expand Down
2 changes: 2 additions & 0 deletions sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,8 @@

#pragma once

#include <sycl/ext/intel/experimental/esimd/simd_view.hpp>

#include <sycl/ext/intel/experimental/esimd/detail/intrin.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/sycl_util.hpp>
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,7 @@

#pragma once

#include <sycl/ext/intel/experimental/esimd/detail/intrin.hpp>
#include <sycl/ext/intel/experimental/esimd/detail/types.hpp>

__SYCL_INLINE_NAMESPACE(cl) {
Expand Down
12 changes: 3 additions & 9 deletions sycl/test/esimd/block_load_store.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,6 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
// RUN: %clangxx -fsycl -fsyntax-only %s

// This test checks that block_load/store API gets successfully compiled.

#include <CL/sycl.hpp>
#include <sycl/ext/intel/experimental/esimd.hpp>
Expand All @@ -12,22 +14,14 @@ SYCL_EXTERNAL void kernel1(
accessor<int, 1, access::mode::read_write, access::target::global_buffer>
&buf) SYCL_ESIMD_FUNCTION {
simd<int, 32> v1(0, 1);
// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/memory.hpp:* {{}}
auto v0 = block_load<int, 32>(buf, 0);
v0 = v0 + v1;
// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/memory.hpp:* {{}}
block_store<int, 32>(buf, 0, v0);
}

SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {
simd<int, 32> v1(0, 1);
// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/memory.hpp:* {{}}
auto v0 = block_load<int, 32>(ptr);
v0 = v0 + v1;
// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/memory.hpp:* {{}}
block_store<int, 32>(ptr, v0);
}
54 changes: 27 additions & 27 deletions sycl/test/esimd/simd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
// expected-no-diagnostics

#include <sycl/ext/intel/experimental/esimd.hpp>
Expand All @@ -7,23 +7,23 @@

using namespace sycl::ext::intel::experimental::esimd;

bool test_simd_ctors() __attribute__((sycl_device)) {
bool test_simd_ctors() SYCL_ESIMD_FUNCTION {
simd<int, 16> v0 = 1;
simd<int, 16> v1(v0);
simd<int, 16> v2(simd<int, 16>(0, 1));
const simd<int, 16> v3{0, 2, 4, 6, 1, 3, 5, 7};
return v0[0] + v1[1] + v2[2] + v3[3] == 1 + 1 + 2 + 6;
}

void test_conversion() __attribute__((sycl_device)) {
void test_conversion() SYCL_ESIMD_FUNCTION {
simd<int, 32> v = 3;
simd<float, 32> f = v;
simd<char, 32> c = f;
simd<char, 16> c1 = f.select<16, 1>(0);
f = v + static_cast<simd<int, 32>>(c);
}

bool test_1d_select() __attribute__((sycl_device)) {
bool test_1d_select() SYCL_ESIMD_FUNCTION {
simd<int, 32> v = 0;
v.select<8, 1>(0) = 1;
v.select<8, 1>(8) = 2;
Expand All @@ -32,7 +32,7 @@ bool test_1d_select() __attribute__((sycl_device)) {
return v[0] + v[8] + v[16] + v[24] == (1 + 2 + 3 + 4);
}

bool test_simd_format() __attribute__((sycl_device)) {
bool test_simd_format() SYCL_ESIMD_FUNCTION {
simd<int, 16> v{0, 1, 2, 3, 4, 5, 6, 7};
auto ref1 = v.format<short>();
auto ref2 = v.format<double>();
Expand All @@ -41,7 +41,7 @@ bool test_simd_format() __attribute__((sycl_device)) {
(decltype(ref3)::getSizeX() == 4) && (decltype(ref3)::getSizeY() == 8);
}

bool test_simd_select() __attribute__((sycl_device)) {
bool test_simd_select() SYCL_ESIMD_FUNCTION {
simd<int, 16> v(0, 1);
auto ref0 = v.select<4, 2>(1); // r{1, 3, 5, 7}
auto ref1 = v.format<int, 4, 4>(); // 0,1,2,3;
Expand All @@ -53,21 +53,21 @@ bool test_simd_select() __attribute__((sycl_device)) {
decltype(ref2)::getStrideY() == 1;
}

bool test_2d_offset() __attribute__((sycl_device)) {
bool test_2d_offset() SYCL_ESIMD_FUNCTION {
simd<int, 16> v = 0;
auto ref = v.format<short, 8, 4>();
return ref.select<2, 2, 2, 2>(2, 1).getOffsetX() == 1 &&
ref.select<2, 2, 2, 2>(2, 1).getOffsetY() == 2;
}

bool test_simd_bin_op_promotion() __attribute__((sycl_device)) {
bool test_simd_bin_op_promotion() SYCL_ESIMD_FUNCTION {
simd<short, 8> v0 = std::numeric_limits<short>::max();
simd<short, 8> v1 = 1;
simd<int, 8> v2 = v0 + v1;
return v2[0] == 32768;
}

bool test_simd_bin_ops() __attribute__((sycl_device)) {
bool test_simd_bin_ops() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0 = 1;
simd<int, 8> v1 = 2;
v0 += v1;
Expand All @@ -82,7 +82,7 @@ bool test_simd_bin_ops() __attribute__((sycl_device)) {
return v0[0] == 1;
}

bool test_simd_unary_ops() __attribute__((sycl_device)) {
bool test_simd_unary_ops() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0 = 1;
simd<int, 8> v1 = 2;
v0 <<= v1;
Expand All @@ -92,7 +92,7 @@ bool test_simd_unary_ops() __attribute__((sycl_device)) {
return v1[0] == 1;
}

bool test_nested_1d_select() __attribute__((sycl_device)) {
bool test_nested_1d_select() SYCL_ESIMD_FUNCTION {
simd<int, 8> r0(0, 1);

auto r1 = r0.select<4, 2>(0);
Expand All @@ -103,7 +103,7 @@ bool test_nested_1d_select() __attribute__((sycl_device)) {
return r0[4] == 37;
}

bool test_format_1d_read() __attribute__((sycl_device)) {
bool test_format_1d_read() SYCL_ESIMD_FUNCTION {
simd<int, 8> r = 0x0FF00F0F;
auto rl = r.format<short>();
auto rl2 = rl.select<8, 2>(0); // 0F0F
Expand All @@ -112,7 +112,7 @@ bool test_format_1d_read() __attribute__((sycl_device)) {
return rl2[0] == 0x0F0F && rh2[0] == 0x0FF0;
}

bool test_format_1d_write() __attribute__((sycl_device)) {
bool test_format_1d_write() SYCL_ESIMD_FUNCTION {
simd<int, 8> r;
auto rl = r.format<short>();
auto rl2 = rl.select<8, 2>(0);
Expand All @@ -122,7 +122,7 @@ bool test_format_1d_write() __attribute__((sycl_device)) {
return r[0] == 0x0FF0;
}

bool test_format_1d_read_write_nested() __attribute__((sycl_device)) {
bool test_format_1d_read_write_nested() SYCL_ESIMD_FUNCTION {
simd<int, 8> v = 0;
auto r1 = v.format<short>();
auto r11 = r1.select<8, 1>(0);
Expand All @@ -134,39 +134,39 @@ bool test_format_1d_read_write_nested() __attribute__((sycl_device)) {
return v[0] == 1 && v[4] == 2;
}

bool test_format_2d_read() __attribute__((sycl_device)) {
bool test_format_2d_read() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0(0, 1);
auto r1 = v0.format<int, 2, 4>();
simd<int, 4> v1 = r1.select<1, 0, 4, 1>(1, 0).read(); // second row
return v1[0] == 4;
}

bool test_format_2d_write() __attribute__((sycl_device)) {
bool test_format_2d_write() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0(0, 1);
auto r1 = v0.format<int, 2, 4>();
r1.select<1, 0, 4, 1>(1, 0) = 37;
return v0[4] == 37;
}

bool test_select_rvalue() __attribute__((sycl_device)) {
bool test_select_rvalue() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0(0, 1);
v0.select<4, 2>(1).select<2, 2>(0) = 37;
return v0[5] == 37;
}

auto test_format_1d_write_rvalue() __attribute__((sycl_device)) {
auto test_format_1d_write_rvalue() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0 = 0x0F0F0F0F;
v0.format<short>().select<8, 2>(0) = 0x0E0E;
return v0[2] == 0x0E0E0E0E;
}

bool test_format_2d_write_rvalue() __attribute__((sycl_device)) {
bool test_format_2d_write_rvalue() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0(0, 1);
v0.format<int, 2, 4>().select<1, 0, 4, 1>(0, 0) = 37;
return v0[3] == 37;
}

auto test_format_2d_read_rvalue() __attribute__((sycl_device)) {
auto test_format_2d_read_rvalue() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0(0, 1);
auto r1 = v0.format<int, 2, 4>()
.select<1, 0, 4, 1>(1, 0)
Expand All @@ -175,7 +175,7 @@ auto test_format_2d_read_rvalue() __attribute__((sycl_device)) {
return r1[0] == 5;
}

bool test_row_read_write() __attribute__((sycl_device)) {
bool test_row_read_write() SYCL_ESIMD_FUNCTION {
simd<int, 16> v0(0, 1);
auto m = v0.format<int, 4, 4>();

Expand All @@ -190,7 +190,7 @@ bool test_row_read_write() __attribute__((sycl_device)) {
return r0[0] == 8 && r1[0] == 16;
}

bool test_column_read_write() __attribute__((sycl_device)) {
bool test_column_read_write() SYCL_ESIMD_FUNCTION {
simd<int, 4> v0(0, 1);
auto m = v0.format<int, 2, 2>();

Expand All @@ -203,35 +203,35 @@ bool test_column_read_write() __attribute__((sycl_device)) {
return v0[0] == 1 && v0[3] == 4;
}

bool test_replicate() __attribute__((sycl_device)) {
bool test_replicate() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0(0, 1);
auto v0_rep = v0.replicate<1>();

return v0[0] == v0_rep[0] && v0[7] == v0_rep[7];
}

bool test_replicate1() __attribute__((sycl_device)) {
bool test_replicate1() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0(0, 1);
auto v0_rep = v0.replicate<4, 2>(2);

return v0[2] == v0_rep[2] && v0[3] == v0_rep[5];
}

bool test_replicate2() __attribute__((sycl_device)) {
bool test_replicate2() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0(0, 1);
auto v0_rep = v0.replicate<2, 4, 2>(1);

return v0_rep[0] == v0[1] && v0_rep[1] == v0[2] && v0_rep[2] == v0[5];
}

bool test_replicate3() __attribute__((sycl_device)) {
bool test_replicate3() SYCL_ESIMD_FUNCTION {
simd<int, 8> v0(0, 1);
auto v0_rep = v0.replicate<2, 4, 2, 2>(1);

return v0_rep[0] == v0[1] && v0_rep[1] == v0[3] && v0_rep[2] == v0[5];
}

bool test_simd_iselect() __attribute__((sycl_device)) {
bool test_simd_iselect() SYCL_ESIMD_FUNCTION {
simd<int, 16> v(0, 1);
simd<ushort, 8> a(0, 2);
auto data = v.iselect(a);
Expand Down