From 4f9c63a043882f7a53da9810b49537417a67310e Mon Sep 17 00:00:00 2001 From: gregory Date: Tue, 16 Aug 2022 22:13:17 -0700 Subject: [PATCH 1/4] Test half type conversion under ESIMD emulator --- .../ESIMD/regression/half_conversion_test.cpp | 97 +++++++++++++++++++ 1 file changed, 97 insertions(+) create mode 100644 SYCL/ESIMD/regression/half_conversion_test.cpp diff --git a/SYCL/ESIMD/regression/half_conversion_test.cpp b/SYCL/ESIMD/regression/half_conversion_test.cpp new file mode 100644 index 0000000000..dcd60cbe01 --- /dev/null +++ b/SYCL/ESIMD/regression/half_conversion_test.cpp @@ -0,0 +1,97 @@ +// REQUIRES: gpu +// UNSUPPORTED: cuda || hip +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +//==- half_conversion_test.cpp - Test for half conversion under ESIMD_EMULATOR +// backend -==/ +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include + +#include + +using namespace ::sycl; +using namespace ::sycl::ext; +using namespace sycl::ext::intel; +using namespace sycl::ext::intel::esimd; + +template struct test_id; +template +using int_type_t = std::conditional_t< + N == 1, int8_t, + std::conditional_t< + N == 2, int16_t, + std::conditional_t>>>; + +template bool test(queue q, int inc) { + Ty *data = new Ty[1]; + + data[0] = (Ty)0; + Ty VAL = (Ty)inc; + + try { + buffer buf(data, range<1>(1)); + q.submit([&](handler &cgh) { + std::cout << "Running on " + << q.get_device().get_info<::sycl::info::device::name>() + << "\n"; + auto acc = buf.template get_access(cgh); + cgh.single_task>([=]() SYCL_ESIMD_KERNEL { + simd offsets(0); + simd vec = gather(acc, offsets); + vec[0] += (Ty)inc; + scalar_store(acc, 0, vec[0]); + }); + }); + } catch (::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + delete[] data; + return false; + } + + using Tint = int_type_t; + Tint ResBits = *(Tint *)&data[0]; + Tint GoldBits = *(Tint *)&VAL; + + std::cout << "Comparison of representation '" << inc << "' of Type " + << typeid(Ty).name() << std::endl; + std::cout << "Bits(data[0]) = 0x" << std::hex << ResBits << " / " + << "Bits(GOLD) = 0x" << GoldBits << std::dec << std::endl; + + if (VAL == data[0]) { + std::cout << "Pass"; + } else { + std::cout << "Fail"; + } + + return ((Ty)inc == data[0]); +} + +int main(int argc, char *argv[]) { + bool passed = true; + queue q; + + std::cout << "\n===================" << std::endl; + passed &= test(q, 1); + std::cout << "\n===================" << std::endl; + passed &= test(q, 1); + std::cout << "\n===================" << std::endl; + passed &= test(q, 1); + std::cout << "\n===================" << std::endl; + + if (passed) { + std::cout << "Pass!!" << std::endl; + } else { + std::cout << "Fail!!" << std::endl; + } + + return passed ? 0 : -1; +} From 592c8b425fb6034f87bc2de40b40606047e739ad Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 17 Aug 2022 09:29:02 -0700 Subject: [PATCH 2/4] Enable failing tests that used to fail in emulator but succeeding now --- SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp | 2 -- SYCL/ESIMD/api/simd_memory_access.cpp | 2 -- SYCL/ESIMD/api/simd_view_select_2d_fp.cpp | 2 -- SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp | 6 ++---- SYCL/ESIMD/api/unary_ops_heavy.cpp | 2 -- 5 files changed, 2 insertions(+), 12 deletions(-) diff --git a/SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp b/SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp index e0503e8597..2beaffee85 100644 --- a/SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp +++ b/SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp @@ -7,8 +7,6 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip -// TODO: esimd_emulator fails due to unimplemented 'half' type -// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/api/simd_memory_access.cpp b/SYCL/ESIMD/api/simd_memory_access.cpp index 6f715cb002..4d42011169 100644 --- a/SYCL/ESIMD/api/simd_memory_access.cpp +++ b/SYCL/ESIMD/api/simd_memory_access.cpp @@ -7,8 +7,6 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip -// TODO: esimd_emulator fails due to unimplemented 'half' type -// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/simd_view_select_2d_fp.cpp b/SYCL/ESIMD/api/simd_view_select_2d_fp.cpp index 7eca9af2de..c397945845 100644 --- a/SYCL/ESIMD/api/simd_view_select_2d_fp.cpp +++ b/SYCL/ESIMD/api/simd_view_select_2d_fp.cpp @@ -7,8 +7,6 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip -// TODO: esimd_emulator fails due to unimplemented 'single_task()' method -// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp b/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp index 1d1372cbfd..7967444f5a 100644 --- a/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp +++ b/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp @@ -7,9 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip -// TODO: esimd_emulator fails due to outdated memory intrinsic -// XFAIL: esimd_emulator -// RUN: %clangxx -fsycl %s -o %t.out +``````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````` // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // // The test checks functionality of the slm gather/scatter ESIMD intrinsics. @@ -33,7 +31,7 @@ #include #include -using namespace sycl; + using namespace sycl; template using Acc = accessor; diff --git a/SYCL/ESIMD/api/unary_ops_heavy.cpp b/SYCL/ESIMD/api/unary_ops_heavy.cpp index 2a4940fe9a..22c4f85c13 100644 --- a/SYCL/ESIMD/api/unary_ops_heavy.cpp +++ b/SYCL/ESIMD/api/unary_ops_heavy.cpp @@ -7,8 +7,6 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip -// TODO: esimd_emulator fails due to unimplemented 'half' type -// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out From e1f8aac8da633e13f34b650f7add70e8df7d3a32 Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 17 Aug 2022 12:28:17 -0700 Subject: [PATCH 3/4] Fix a typo --- SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp b/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp index 7967444f5a..9a1bd93098 100644 --- a/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp +++ b/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip -``````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````````` // RUN: %clangxx -fsycl %s -o %t.out +// RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // // The test checks functionality of the slm gather/scatter ESIMD intrinsics. @@ -31,7 +31,7 @@ #include #include - using namespace sycl; +using namespace sycl; template using Acc = accessor; From ba28796133df2c11e3473b457328f1b67f5eb2b1 Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 17 Aug 2022 23:42:45 -0700 Subject: [PATCH 4/4] Address PR comments --- SYCL/ESIMD/regression/half_conversion_test.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/SYCL/ESIMD/regression/half_conversion_test.cpp b/SYCL/ESIMD/regression/half_conversion_test.cpp index dcd60cbe01..7be3dce9d1 100644 --- a/SYCL/ESIMD/regression/half_conversion_test.cpp +++ b/SYCL/ESIMD/regression/half_conversion_test.cpp @@ -11,9 +11,9 @@ // //===----------------------------------------------------------------------===// -#include #include #include +#include #include @@ -22,7 +22,6 @@ using namespace ::sycl::ext; using namespace sycl::ext::intel; using namespace sycl::ext::intel::esimd; -template struct test_id; template using int_type_t = std::conditional_t< N == 1, int8_t, @@ -44,7 +43,7 @@ template bool test(queue q, int inc) { << q.get_device().get_info<::sycl::info::device::name>() << "\n"; auto acc = buf.template get_access(cgh); - cgh.single_task>([=]() SYCL_ESIMD_KERNEL { + cgh.single_task([=]() SYCL_ESIMD_KERNEL { simd offsets(0); simd vec = gather(acc, offsets); vec[0] += (Ty)inc;