diff --git a/SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp b/SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp new file mode 100644 index 0000000000..8a908497c0 --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp @@ -0,0 +1,104 @@ +//===------ noinline_args_size_common.hpp - 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 +// +//===----------------------------------------------------------------------===// +// +// The test checks that ESIMD kernels support call of noinline function from +// main function with different total arguments size and retval size. Cases: +// Total arguments size < %arg register size (32 GRFs) +// Total arguments size == %arg register size +// Total arguments size > %arg register size (i.e. stack mem is required) +// Return value size < %retval register size (12 GRFs) +// Return value size == %retval register size +// Return value size > %retval register size + +#include "esimd_test_utils.hpp" + +#include +#include +#include + +static_assert(SIZE >= VL, "Size must greater than or equal to VL"); +static_assert(SIZE % VL == 0, "Size must be multiple of VL"); +constexpr unsigned ROWS = SIZE / VL; + +using namespace cl::sycl; + +class KernelID; + +template +ESIMD_NOINLINE TC add(TA A, TB B) { + return (TC)A + (TC)B; +} + +int main(void) { + queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()); + + auto dev = q.get_device(); + std::cout << "Running on " << dev.get_info() << "\n"; + auto ctx = q.get_context(); + + a_data_t *A = static_cast( + sycl::malloc_shared(SIZE * sizeof(a_data_t), dev, ctx)); + for (int i = 0; i < SIZE; i++) + A[i] = (a_data_t)1; + + b_data_t *B = static_cast( + sycl::malloc_shared(SIZE * sizeof(b_data_t), dev, ctx)); + for (int i = 0; i < SIZE; i++) + B[i] = (b_data_t)i; + + c_data_t *C = static_cast( + sycl::malloc_shared(SIZE * sizeof(c_data_t), dev, ctx)); + memset(C, 0, SIZE * sizeof(c_data_t)); + + try { + auto qq = q.submit([&](handler &cgh) { + cgh.parallel_for( + sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + + simd va(0); + simd vb(0); + for (int j = 0; j < ROWS; j++) { + va.select(j * VL) = block_load(A + j * VL); + vb.select(j * VL) = block_load(B + j * VL); + } + + auto vc = add, simd, + simd>(va, vb); + + for (int j = 0; j < ROWS; j++) + block_store(C + j * VL, vc.select(j * VL)); + }); + }); + + qq.wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + sycl::free(A, ctx); + sycl::free(B, ctx); + sycl::free(C, ctx); + return e.get_cl_code(); + } + + unsigned err_cnt = 0; + for (int i = 0; i < SIZE; i++) + if (C[i] != A[i] + B[i]) + err_cnt++; + + sycl::free(A, ctx); + sycl::free(B, ctx); + sycl::free(C, ctx); + + if (err_cnt > 0) { + std::cout << "FAILED" << std::endl; + return 1; + } + + std::cout << "passed" << std::endl; + return 0; +} diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_192.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_192.cpp new file mode 100644 index 0000000000..e152604bb0 --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_192.cpp @@ -0,0 +1,22 @@ +//===-- noinline_args_char_int_size_192.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 192; + +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_256.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_256.cpp new file mode 100644 index 0000000000..ff06f8d4fb --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_256.cpp @@ -0,0 +1,22 @@ +//===-- noinline_args_char_int_size_256.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 256; + +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_512.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_512.cpp new file mode 100644 index 0000000000..1894de3c5a --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_512.cpp @@ -0,0 +1,22 @@ +//===-- noinline_args_char_int_size_512.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 512; + +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_64.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_64.cpp new file mode 100644 index 0000000000..5785b3e48b --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_64.cpp @@ -0,0 +1,22 @@ +//===--- noinline_args_char_int_size_64.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 64; + +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_96.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_96.cpp new file mode 100644 index 0000000000..cbd084e192 --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_96.cpp @@ -0,0 +1,22 @@ +//===--- noinline_args_char_int_size_96.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 96; + +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp new file mode 100644 index 0000000000..2cb5ca0361 --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp @@ -0,0 +1,22 @@ +//===-- noinline_args_char_int_size_192.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 192; + +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp new file mode 100644 index 0000000000..ade1d5d91b --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp @@ -0,0 +1,22 @@ +//===-- noinline_args_char_int_size_256.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 256; + +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp new file mode 100644 index 0000000000..22e3b14fe8 --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp @@ -0,0 +1,22 @@ +//===-- noinline_args_char_int_size_512.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 512; + +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp new file mode 100644 index 0000000000..69fcf80022 --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp @@ -0,0 +1,22 @@ +//===--- noinline_args_char_int_size_64.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 64; + +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp new file mode 100644 index 0000000000..7c9c198cea --- /dev/null +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp @@ -0,0 +1,22 @@ +//===--- noinline_args_char_int_size_96.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out + +#include + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 96; + +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; + +#include "Inputs/noinline_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_call_from_func.cpp b/SYCL/ESIMD/noinline_call_from_func.cpp new file mode 100644 index 0000000000..6484bfee6d --- /dev/null +++ b/SYCL/ESIMD/noinline_call_from_func.cpp @@ -0,0 +1,71 @@ +//===------ noinline_call_from_func.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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// Test currently timeouts on Windows Level Zero +// UNSUPPORTED: windows && level_zero +// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda +// +// The test checks that ESIMD kernels support call of noinline function from +// within other functions. + +#include "esimd_test_utils.hpp" + +#include +#include + +#include + +class KernelID; + +ESIMD_NOINLINE int add(int A, int B) { return A + B; } + +template ESIMD_NOINLINE void test(AccTy acc, int A, int B) { + using namespace sycl::INTEL::gpu; + + auto res = add(A, B); + + scalar_store(acc, 0, res); +} + +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"; + + int result = 0; + int *output = &result; + + int in1 = 100; + int in2 = 233; + + 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 { test(acc, in1, in2); }); + }); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return e.get_cl_code(); + } + + if (result != (in1 + in2)) { + std::cout << "Failed: with result: " << result << std::endl; + return 1; + } + + std::cout << "passed" << std::endl; + return 0; +} diff --git a/SYCL/ESIMD/noinline_call_recursive.cpp b/SYCL/ESIMD/noinline_call_recursive.cpp new file mode 100644 index 0000000000..60d879ee88 --- /dev/null +++ b/SYCL/ESIMD/noinline_call_recursive.cpp @@ -0,0 +1,82 @@ +//===------ 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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda +// +// The test checks that ESIMD kernels support recursive call of noinline +// functions. + +#include "esimd_test_utils.hpp" + +#include +#include + +#include + +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::INTEL::gpu; + + auto res = add(in1, in2, in3); + + scalar_store(acc, 0, res); + }); + }); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return e.get_cl_code(); + } + + 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; +}