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..9a1bd93098 100644 --- a/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp +++ b/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp @@ -7,8 +7,6 @@ //===----------------------------------------------------------------------===// // 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: %GPU_RUN_PLACEHOLDER %t.out // 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 diff --git a/SYCL/ESIMD/regression/half_conversion_test.cpp b/SYCL/ESIMD/regression/half_conversion_test.cpp new file mode 100644 index 0000000000..7be3dce9d1 --- /dev/null +++ b/SYCL/ESIMD/regression/half_conversion_test.cpp @@ -0,0 +1,96 @@ +// 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 +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; +}