From bfca4d2307724ae87276a4b78c16db211d766fdc Mon Sep 17 00:00:00 2001
From: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>
Date: Fri, 28 May 2021 18:23:23 -0700
Subject: [PATCH] [SYCL][ESIMD] Host-compile simd.cpp test, fix errors &
 warnings.

Also removed deprecation message checking in the block_load_store.cpp,
as it leads to frequent test failures when new warnings appear in other parts
of the code base. This is unnecessary maintenance overhead.

Change-Id: I68437a46a4c17a6af6bef16ff45c804151bb956d
Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>
---
 .../ext/intel/experimental/esimd/math.hpp     |  3 +-
 .../ext/intel/experimental/esimd/simd.hpp     |  2 +
 .../intel/experimental/esimd/simd_view.hpp    |  1 +
 sycl/test/esimd/block_load_store.cpp          | 12 ++---
 sycl/test/esimd/simd.cpp                      | 54 +++++++++----------
 5 files changed, 35 insertions(+), 37 deletions(-)

diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
index 7383e756b5490..507132319368d 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/math.hpp
@@ -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.
@@ -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>
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp
index f74915385f971..bebb4fe8d39af 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp
@@ -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>
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp
index e0caea62af645..eaba47f098fab 100644
--- a/sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp
@@ -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) {
diff --git a/sycl/test/esimd/block_load_store.cpp b/sycl/test/esimd/block_load_store.cpp
index 5d0ec9aaa44f7..911dff6bedb4d 100644
--- a/sycl/test/esimd/block_load_store.cpp
+++ b/sycl/test/esimd/block_load_store.cpp
@@ -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>
@@ -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);
 }
diff --git a/sycl/test/esimd/simd.cpp b/sycl/test/esimd/simd.cpp
index e0d67cd44f9ad..dcbaafce8ce0c 100644
--- a/sycl/test/esimd/simd.cpp
+++ b/sycl/test/esimd/simd.cpp
@@ -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>
@@ -7,7 +7,7 @@
 
 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));
@@ -15,7 +15,7 @@ bool test_simd_ctors() __attribute__((sycl_device)) {
   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;
@@ -23,7 +23,7 @@ void test_conversion() __attribute__((sycl_device)) {
   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;
@@ -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>();
@@ -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;
@@ -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;
@@ -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;
@@ -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);
@@ -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
@@ -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);
@@ -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);
@@ -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)
@@ -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>();
 
@@ -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>();
 
@@ -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);