From e2bc253d0106e68f80930080c29b257c9b64df3e Mon Sep 17 00:00:00 2001 From: gregory Date: Wed, 2 Nov 2022 15:54:52 -0700 Subject: [PATCH 1/2] Add test to validate tanh --- SYCL/ESIMD/regression/tanh_fix_test.cpp | 118 ++++++++++++++++++++++++ 1 file changed, 118 insertions(+) create mode 100644 SYCL/ESIMD/regression/tanh_fix_test.cpp diff --git a/SYCL/ESIMD/regression/tanh_fix_test.cpp b/SYCL/ESIMD/regression/tanh_fix_test.cpp new file mode 100644 index 0000000000..796a12b9fc --- /dev/null +++ b/SYCL/ESIMD/regression/tanh_fix_test.cpp @@ -0,0 +1,118 @@ +// REQUIRES: gpu +// UNSUPPORTED: cuda || hip +// TODO online compiler check fails for esimd_emulator +// XFAIL: esimd_emulator +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +//==- tanh_compatibility_test.cpp - Test for tanh -==// +// +// 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 +#include +#include + +constexpr auto sycl_write = sycl::access::mode::write; +#define SIMD 16 + +int test_tanh(float x) { + std::vector out(SIMD); + std::vector out1(SIMD); + + float ha = x; + float scalar_result = 0; + float scalar_result1 = 0; + + { + sycl::queue queue; + + sycl::buffer vector_buffer(out.data(), out.size()); + sycl::buffer scalar_buffer(&scalar_result, sycl::range<1>(1)); + sycl::buffer vector_buffer1(out1.data(), out1.size()); + sycl::buffer scalar_buffer1(&scalar_result1, sycl::range<1>(1)); + + auto e = queue.submit([&](sycl::handler &cgh) { + sycl::accessor vector_out = + vector_buffer.get_access(cgh); + sycl::accessor scalar_out = + scalar_buffer.get_access(cgh); + sycl::accessor vector_out1 = + vector_buffer1.get_access(cgh); + sycl::accessor scalar_out1 = + scalar_buffer1.get_access(cgh); + + auto kernel = ([=]() [[intel::sycl_explicit_simd]] { + using namespace sycl::ext::intel::esimd; + + simd a = ha; + + simd vector_result = + sycl::ext::intel::experimental::esimd::tanh(a); + simd scalar_result = + sycl::ext::intel::experimental::esimd::tanh(ha); + simd vector_result1 = + sycl::ext::intel::experimental::esimd::tanh_cody_waite(a); + simd scalar_result1 = + sycl::ext::intel::experimental::esimd::tanh_cody_waite(ha); + + vector_result.copy_to(vector_out, 0); + scalar_result.copy_to(scalar_out, 0); + vector_result1.copy_to(vector_out1, 0); + scalar_result1.copy_to(scalar_out1, 0); + }); + + cgh.single_task(kernel); + }); + queue.wait(); + } + + float std_result = std::tanh(ha); + + if (std::fabs(std_result - out[0]) > 0.000001f) { + std::cout << "Vector test failed for " << x << std::hex << " (" + << *(int *)&std_result << " | " << *(int *)&out[0] << ")." + << std::dec << std::endl; + return 1; + } + + if (std::fabs(std_result - scalar_result) > 0.000001f) { + std::cout << "Scalar test failed for " << x << "." << std::endl; + return 1; + } + + if (std::fabs(std_result - out1[0]) > 0.000001f) { + std::cout << "Vector test failed for cody waite implementation for " << x + << "." << std::endl; + return 1; + } + + if (std::fabs(std_result - scalar_result1) > 0.000001f) { + std::cout << "Scalar test failed for cody waite implementation for " << x + << "." << std::endl; + return 1; + } + + return 0; +} + +int main() { + + int test_result = 0; + + for (float x = -100.f; x < 100.f; x += 0.1f) { + test_result |= test_tanh(x); + } + + if (!test_result) { + std::cout << "Pass" << std::endl; + } + return test_result; +} From 685940e23b98184451aab11fdd26a6dcf7436f3a Mon Sep 17 00:00:00 2001 From: gregory Date: Thu, 3 Nov 2022 13:37:28 -0700 Subject: [PATCH 2/2] Address PR comments --- SYCL/ESIMD/regression/tanh_fix_test.cpp | 30 +++++++++++++------------ 1 file changed, 16 insertions(+), 14 deletions(-) diff --git a/SYCL/ESIMD/regression/tanh_fix_test.cpp b/SYCL/ESIMD/regression/tanh_fix_test.cpp index 796a12b9fc..81ebd2c8d7 100644 --- a/SYCL/ESIMD/regression/tanh_fix_test.cpp +++ b/SYCL/ESIMD/regression/tanh_fix_test.cpp @@ -5,7 +5,7 @@ // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out -//==- tanh_compatibility_test.cpp - Test for tanh -==// +//==- tanh_fix_test.cpp - Test for tanh -==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -76,30 +76,32 @@ int test_tanh(float x) { float std_result = std::tanh(ha); - if (std::fabs(std_result - out[0]) > 0.000001f) { - std::cout << "Vector test failed for " << x << std::hex << " (" - << *(int *)&std_result << " | " << *(int *)&out[0] << ")." - << std::dec << std::endl; - return 1; - } - if (std::fabs(std_result - scalar_result) > 0.000001f) { std::cout << "Scalar test failed for " << x << "." << std::endl; return 1; } - if (std::fabs(std_result - out1[0]) > 0.000001f) { - std::cout << "Vector test failed for cody waite implementation for " << x - << "." << std::endl; - return 1; - } - if (std::fabs(std_result - scalar_result1) > 0.000001f) { std::cout << "Scalar test failed for cody waite implementation for " << x << "." << std::endl; return 1; } + for (int i = 0; i < SIMD; ++i) { + if (scalar_result != out[i] && + !(std::isnan(scalar_result) && std::isnan(out[i]))) { + std::cout << "Vector test failed for " << x << "." << std::endl; + return 1; + } + + if (scalar_result1 != out1[i] && + !(std::isnan(scalar_result1) && std::isnan(out1[i]))) { + std::cout << "Vector test failed for cody waite implementation for " << x + << "." << std::endl; + return 1; + } + } + return 0; }