diff --git a/sycl/test-e2e/ESIMD/fp_call_recursive.cpp b/sycl/test-e2e/ESIMD/fp_call_recursive.cpp deleted file mode 100644 index 5869bbc7ae553..0000000000000 --- a/sycl/test-e2e/ESIMD/fp_call_recursive.cpp +++ /dev/null @@ -1,79 +0,0 @@ -//==--------------- fp_call_recursive.cpp - DPC++ ESIMD on-device test ----==// -// -// 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 -// -//===----------------------------------------------------------------------===// -// Recursion is not supported in ESIMD (intel/llvm PR#3390) -// REQUIRES: TEMPORARY_DISBLED -// RUN: %{build} -Xclang -fsycl-allow-func-ptr -o %t.out -// RUN: %{run} %t.out -// -// The test checks that ESIMD kernels support use of function pointers -// recursively. - -#include "esimd_test_utils.hpp" - -class KernelID; - -ESIMD_NOINLINE unsigned add(unsigned A, unsigned B, unsigned C) { - if (B == 0) - return A; - - auto foo = &add; - return (B % C == 0) ? foo(A + 1, B - 1, C) : foo(A - C, B - 2, C); -} - -int main(int argc, char **argv) { - queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler()); - - auto dev = q.get_device(); - std::cout << "Running on " << dev.get_info() << "\n"; - - unsigned result = 0; - unsigned *output = &result; - - unsigned in1 = 233; - unsigned in2 = 21; - unsigned in3 = 3; - - try { - buffer buf(output, range<1>(1)); - - q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); - - cgh.parallel_for(sycl::range<1>{1}, - [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::ext::intel::esimd; - - auto foo = &add; - auto res = foo(in1, in2, in3); - - scalar_store(acc, 0, res); - }); - }); - } catch (sycl::exception const &e) { - std::cout << "SYCL exception caught: " << e.what() << std::endl; - return e.code().value(); - } - - int etalon = in1; - while (in2 > 0) { - if (in2 % in3 == 0) { - etalon += 1; - in2 -= 1; - } else { - etalon -= in3; - in2 -= 2; - } - } - - if (result != etalon) { - std::cout << "Failed with result: " << result << std::endl; - return 1; - } - - return 0; -} diff --git a/sycl/test-e2e/ESIMD/noinline_call_recursive.cpp b/sycl/test-e2e/ESIMD/noinline_call_recursive.cpp deleted file mode 100644 index d11026688c602..0000000000000 --- a/sycl/test-e2e/ESIMD/noinline_call_recursive.cpp +++ /dev/null @@ -1,77 +0,0 @@ -//===------ noinline_call_recursive.cpp - DPC++ ESIMD on-device test -----===// -// -// 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 -// -//===----------------------------------------------------------------------===// -// Recursion is not supported in ESIMD (intel/llvm PR#3390) -// REQUIRES: TEMPORARY_DISBLED -// RUN: %{build} -o %t.out -// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %{run} %t.out -// -// The test checks that ESIMD kernels support recursive call of noinline -// functions. - -#include "esimd_test_utils.hpp" - -class KernelID; - -ESIMD_NOINLINE unsigned add(unsigned A, unsigned B, unsigned C) { - if (B == 0) - return A; - - return (B % C == 0) ? add(A + 1, B - 1, C) : add(A - C, B - 2, C); -} - -int main(int argc, char **argv) { - queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler()); - - auto dev = q.get_device(); - std::cout << "Running on " << dev.get_info() << "\n"; - - unsigned result = 0; - unsigned *output = &result; - - unsigned in1 = 233; - unsigned in2 = 21; - unsigned in3 = 3; - - try { - buffer buf(output, range<1>(1)); - - q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); - - cgh.parallel_for(sycl::range<1>{1}, - [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::ext::intel::esimd; - - auto res = add(in1, in2, in3); - - scalar_store(acc, 0, res); - }); - }); - } catch (sycl::exception const &e) { - std::cout << "SYCL exception caught: " << e.what() << std::endl; - return e.code().value(); - } - - int etalon = in1; - while (in2 > 0) { - if (in2 % in3 == 0) { - etalon += 1; - in2 -= 1; - } else { - etalon -= in3; - in2 -= 2; - } - } - - if (result != etalon) { - std::cout << "Failed: " << result << std::endl; - return 1; - } - - return 0; -}