From 5d55ee256063201d5703249a6e072e5f22069445 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 4 Dec 2020 01:03:44 +0300 Subject: [PATCH 01/32] ESIMD: add TPM tests --- SYCL/ESIMD/tpm_basic.cpp | 134 +++++++++++++++++++++ SYCL/ESIMD/tpm_pointer.cpp | 165 ++++++++++++++++++++++++++ SYCL/ESIMD/tpm_pointer_v2.cpp | 215 ++++++++++++++++++++++++++++++++++ 3 files changed, 514 insertions(+) create mode 100644 SYCL/ESIMD/tpm_basic.cpp create mode 100644 SYCL/ESIMD/tpm_pointer.cpp create mode 100644 SYCL/ESIMD/tpm_pointer_v2.cpp diff --git a/SYCL/ESIMD/tpm_basic.cpp b/SYCL/ESIMD/tpm_basic.cpp new file mode 100644 index 0000000000..702a629745 --- /dev/null +++ b/SYCL/ESIMD/tpm_basic.cpp @@ -0,0 +1,134 @@ +//==---------------- basic_tpm.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 +// +//===----------------------------------------------------------------------===// +// TODO enable on Windows and Level Zero +// REQUIRES: linux && gpu && opencl +// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %HOST_RUN_PLACEHOLDER %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out + +#include "esimd_test_utils.hpp" + +#include +#include +#include + +using namespace cl::sycl; + +int main(void) { + constexpr unsigned VL = 8; + constexpr unsigned SZ = 800; // big enough to use TPM + + 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(); + + int *output = static_cast( + malloc_shared(VL * sizeof(int), dev, ctx)); + memset(output, 0, VL * sizeof(int)); + + int off1 = 16; + int off2 = 128; + int base1 = 500; + int base2 = 100; + int divisor = 4; + + { + auto e = q.submit([&](handler &cgh) { + cgh.parallel_for( + sycl::range<1> {1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + off1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + off2) % SZ; + x2[idx] = base2 << (j % 32); + } + + // some work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; + } + + // some work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) > base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + int o = 0; + for (int j = 0; j < SZ; ++j) { + if (j % 3 == 0) + o += x1[j]; + else + o -= x2[j]; + } + + simd inc(0, 1); + block_store(output, inc + o); + }); + }); + e.wait(); + } + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + off1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + // same work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + off2) % SZ; + x2[idx] = base2 << (j % 32); + } + + // same work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) > base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + int o = 0; + for (int j = 0; j < SZ; ++j) { + if (j % 3 == 0) + o += x1[j]; + else + o -= x2[j]; + } + + int err_cnt = 0; + for (int j = 0; j < VL; ++j) { + if (output[j] != (o + j)) + err_cnt += 1; + } + + if (err_cnt > 0) { + std::cout << "GPU: " << output[0] << " vs CPU: " << o << "\n"; + std::cout << "FAILED.\n"; + return 1; + } + + std::cout << "Passed\n"; + return 0; +} diff --git a/SYCL/ESIMD/tpm_pointer.cpp b/SYCL/ESIMD/tpm_pointer.cpp new file mode 100644 index 0000000000..28440bce40 --- /dev/null +++ b/SYCL/ESIMD/tpm_pointer.cpp @@ -0,0 +1,165 @@ +//==---------------- pointer_tpm.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 +// +//===----------------------------------------------------------------------===// +// TODO enable on Windows and Level Zero +// REQUIRES: linux && gpu && opencl +// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %HOST_RUN_PLACEHOLDER %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out + +#include "esimd_test_utils.hpp" + +#include +#include +#include + +using namespace cl::sycl; + +int main(void) { + constexpr unsigned VL = 8; + constexpr unsigned SZ = 800; // big enough to use TPM + + 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(); + + int *output = static_cast( + malloc_shared(VL * sizeof(int), dev, ctx)); + memset(output, 0, VL * sizeof(int)); + + int offx1 = 55; + int offx2 = 11; + int offy = 111; + int base1 = 500; + int base2 = 100; + int divisor = 4; + + { + auto e = q.submit([&](handler &cgh) { + cgh.parallel_for( + sycl::range<1> {1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx2) % SZ; + x2[idx] = base2 << (j % 32); + } + + int* y[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy) % SZ; + y[j] = j % 6 == 0 ? x1 + idx : x2 + idx; + } + + // some work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; + } + + // some work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) > base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + // some work with Y + for (int j = 0; j < SZ; j += 2) { + if ((j % 6 != 0) && (y[j] > y[j + 1])) { + auto temp = y[j]; + y[j] = y[j + 1]; + y[j + 1] = temp; + } + if (*(y[j]) > *(y[j + 1])) + *(y[j]) = *(y[j + 1]) - *(y[j]); + } + + int o = 0; + for (int j = 0; j < SZ; ++j) { + if (j % 3 == 0) + o += *(y[j]); + } + + simd inc(0, 1); + block_store(output, inc + o); + }); + }); + e.wait(); + } + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + // same work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx2) % SZ; + x2[idx] = base2 << (j % 32); + } + + // same work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) > base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + int* y[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy) % SZ; + y[j] = j % 6 == 0 ? x1 + idx : x2 + idx; + } + + // same work with Y + for (int j = 0; j < SZ; j += 2) { + if ((j % 6 != 0) && (y[j] > y[j + 1])) { + auto temp = y[j]; + y[j] = y[j + 1]; + y[j + 1] = temp; + } + if (*(y[j]) > *(y[j + 1])) + *(y[j]) = *(y[j + 1]) - *(y[j]); + } + + int o = 0; + for (int j = 0; j < SZ; ++j) { + if (j % 3 == 0) + o += *(y[j]); + } + + int err_cnt = 0; + for (int j = 0; j < VL; ++j) { + if (output[j] != (o + j)) + err_cnt += 1; + } + + if (err_cnt > 0) { + std::cout << "GPU: " << output[0] << " vs CPU: " << o << "\n"; + std::cout << "FAILED.\n"; + return 1; + } + + std::cout << "Passed\n"; + return 0; +} diff --git a/SYCL/ESIMD/tpm_pointer_v2.cpp b/SYCL/ESIMD/tpm_pointer_v2.cpp new file mode 100644 index 0000000000..ee8c996992 --- /dev/null +++ b/SYCL/ESIMD/tpm_pointer_v2.cpp @@ -0,0 +1,215 @@ +//==---------------- pointer_tpm_v2.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 +// +//===----------------------------------------------------------------------===// +// TODO enable on Windows and Level Zero +// REQUIRES: linux && gpu && opencl +// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %HOST_RUN_PLACEHOLDER %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out + +#include "esimd_test_utils.hpp" + +#include +#include +#include + +using namespace cl::sycl; + +int main(void) { + constexpr unsigned VL = 8; + constexpr unsigned SZ = 800; // big enough to use TPM + + 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(); + + int *output = static_cast( + malloc_shared(VL * sizeof(int), dev, ctx)); + memset(output, 0, VL * sizeof(int)); + + int offx1 = 111; + int offx2 = 55; + int offy1 = 499; + int offy2 = 223; + int offz = 99; + int base1 = 500; + int base2 = 100; + int divisor = 4; + + { + auto e = q.submit([&](handler &cgh) { + cgh.parallel_for( + sycl::range<1> {1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx2) % SZ; + x2[idx] = base2 << (j % 32); + } + + int* y1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy1) % SZ; + y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; + } + + int* y2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy2) % SZ; + y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; + } + + int** z[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offz) % SZ; + z[j] = y1 + idx; + } + + // some work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; + } + + // some work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) < base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + // some work with Y1 + for (int j = 0; j < SZ; j += 2) { + if (*(y1[j]) > *(y1[j + 1])) + *(y1[j]) = *(y1[j + 1]) - *(y1[j]); + } + + // some work with Y2 + for (int j = 1; j < SZ - 1; j += 2) { + if ((*(y2[j]) <= *(y2[j + 1]))) { + auto temp = y2[j]; + y2[j] = y2[j + 1]; + y2[j + 1] = temp; + } + } + + // some work with Z + for (int j = 0; j < SZ - 1; ++j) { + if ( *(*(z[j])) < *(*(z[j + 1])) ) + z[j] = y2 + j; + if (j % 18 == 0) + (*(*(z[j])))++; + } + + int o = 0; + for (int j = 0; j < SZ; ++j) { + if (j % 3 == 0) + o += *(*(z[j])); + } + + simd inc(0, 1); + block_store(output, inc + o); + }); + }); + e.wait(); + } + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + // same work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx2) % SZ; + x2[idx] = base2 << (j % 32); + } + + // same work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) < base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + int* y1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy1) % SZ; + y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; + } + + // same work with Y1 + for (int j = 0; j < SZ; j += 2) { + if (*(y1[j]) > *(y1[j + 1])) + *(y1[j]) = *(y1[j + 1]) - *(y1[j]); + } + + int* y2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy2) % SZ; + y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; + } + + // same work with Y2 + for (int j = 1; j < SZ - 1; j += 2) { + if ((*(y2[j]) <= *(y2[j + 1]))) { + auto temp = y2[j]; + y2[j] = y2[j + 1]; + y2[j + 1] = temp; + } + } + + int** z[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offz) % SZ; + z[j] = y1 + idx; + } + + // same work with Z + for (int j = 0; j < SZ - 1; ++j) { + if ( *(*(z[j])) < *(*(z[j + 1])) ) + z[j] = y2 + j; + if (j % 18 == 0) + (*(*(z[j])))++; + } + + int o = 0; + for (int j = 0; j < SZ; ++j) { + if (j % 3 == 0) + o += *(*(z[j])); + } + + int err_cnt = 0; + for (int j = 0; j < VL; ++j) { + if (output[j] != (o + j)) + err_cnt += 1; + } + + if (err_cnt > 0) { + std::cout << "GPU: " << output[0] << " vs CPU: " << o << "\n"; + std::cout << "FAILED.\n"; + return 1; + } + + std::cout << "Passed\n"; + return 0; +} From bb7498da5050991e73c5161105da3c3d07719c04 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 4 Dec 2020 01:03:44 +0300 Subject: [PATCH 02/32] [SYCL][ESIMD] TPM tests stylecheck fix --- SYCL/ESIMD/tpm_pointer_v2.cpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/SYCL/ESIMD/tpm_pointer_v2.cpp b/SYCL/ESIMD/tpm_pointer_v2.cpp index ee8c996992..3a5f2aa092 100644 --- a/SYCL/ESIMD/tpm_pointer_v2.cpp +++ b/SYCL/ESIMD/tpm_pointer_v2.cpp @@ -108,7 +108,11 @@ int main(void) { // some work with Z for (int j = 0; j < SZ - 1; ++j) { +<<<<<<< HEAD if ( *(*(z[j])) < *(*(z[j + 1])) ) +======= + if (*(*(z[j])) < *(*(z[j + 1]))) +>>>>>>> fe7b8274 (ESIMD: add TPM tests) z[j] = y2 + j; if (j % 18 == 0) (*(*(z[j])))++; @@ -186,7 +190,11 @@ int main(void) { // same work with Z for (int j = 0; j < SZ - 1; ++j) { +<<<<<<< HEAD if ( *(*(z[j])) < *(*(z[j + 1])) ) +======= + if (*(*(z[j])) < *(*(z[j + 1]))) +>>>>>>> fe7b8274 (ESIMD: add TPM tests) z[j] = y2 + j; if (j % 18 == 0) (*(*(z[j])))++; From 59999ef43b987acc76ba01a69132430b92dab33e Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 4 Dec 2020 18:32:11 +0300 Subject: [PATCH 03/32] [SYCL][ESIMD] clang-format patch --- SYCL/ESIMD/tpm_basic.cpp | 79 ++++++++-------- SYCL/ESIMD/tpm_pointer.cpp | 111 ++++++++++++----------- SYCL/ESIMD/tpm_pointer_v2.cpp | 163 +++++++++++++++++----------------- 3 files changed, 175 insertions(+), 178 deletions(-) diff --git a/SYCL/ESIMD/tpm_basic.cpp b/SYCL/ESIMD/tpm_basic.cpp index 702a629745..cf2bffcb27 100644 --- a/SYCL/ESIMD/tpm_basic.cpp +++ b/SYCL/ESIMD/tpm_basic.cpp @@ -30,8 +30,7 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto ctx = q.get_context(); - int *output = static_cast( - malloc_shared(VL * sizeof(int), dev, ctx)); + int *output = static_cast(malloc_shared(VL * sizeof(int), dev, ctx)); memset(output, 0, VL * sizeof(int)); int off1 = 16; @@ -43,44 +42,44 @@ int main(void) { { auto e = q.submit([&](handler &cgh) { cgh.parallel_for( - sycl::range<1> {1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + off1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + off2) % SZ; - x2[idx] = base2 << (j % 32); - } - - // some work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - // some work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) > base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - int o = 0; - for (int j = 0; j < SZ; ++j) { - if (j % 3 == 0) - o += x1[j]; - else - o -= x2[j]; - } - - simd inc(0, 1); - block_store(output, inc + o); - }); + sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + off1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + off2) % SZ; + x2[idx] = base2 << (j % 32); + } + + // some work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; + } + + // some work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) > base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + int o = 0; + for (int j = 0; j < SZ; ++j) { + if (j % 3 == 0) + o += x1[j]; + else + o -= x2[j]; + } + + simd inc(0, 1); + block_store(output, inc + o); + }); }); e.wait(); } diff --git a/SYCL/ESIMD/tpm_pointer.cpp b/SYCL/ESIMD/tpm_pointer.cpp index 28440bce40..b19ad1d319 100644 --- a/SYCL/ESIMD/tpm_pointer.cpp +++ b/SYCL/ESIMD/tpm_pointer.cpp @@ -30,13 +30,12 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto ctx = q.get_context(); - int *output = static_cast( - malloc_shared(VL * sizeof(int), dev, ctx)); + int *output = static_cast(malloc_shared(VL * sizeof(int), dev, ctx)); memset(output, 0, VL * sizeof(int)); int offx1 = 55; int offx2 = 11; - int offy = 111; + int offy = 111; int base1 = 500; int base2 = 100; int divisor = 4; @@ -44,59 +43,59 @@ int main(void) { { auto e = q.submit([&](handler &cgh) { cgh.parallel_for( - sycl::range<1> {1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx2) % SZ; - x2[idx] = base2 << (j % 32); - } - - int* y[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy) % SZ; - y[j] = j % 6 == 0 ? x1 + idx : x2 + idx; - } - - // some work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - // some work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) > base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - // some work with Y - for (int j = 0; j < SZ; j += 2) { - if ((j % 6 != 0) && (y[j] > y[j + 1])) { - auto temp = y[j]; - y[j] = y[j + 1]; - y[j + 1] = temp; + sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx2) % SZ; + x2[idx] = base2 << (j % 32); + } + + int *y[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy) % SZ; + y[j] = j % 6 == 0 ? x1 + idx : x2 + idx; + } + + // some work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; } - if (*(y[j]) > *(y[j + 1])) - *(y[j]) = *(y[j + 1]) - *(y[j]); - } - - int o = 0; - for (int j = 0; j < SZ; ++j) { - if (j % 3 == 0) - o += *(y[j]); - } - - simd inc(0, 1); - block_store(output, inc + o); - }); + + // some work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) > base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + // some work with Y + for (int j = 0; j < SZ; j += 2) { + if ((j % 6 != 0) && (y[j] > y[j + 1])) { + auto temp = y[j]; + y[j] = y[j + 1]; + y[j + 1] = temp; + } + if (*(y[j]) > *(y[j + 1])) + *(y[j]) = *(y[j + 1]) - *(y[j]); + } + + int o = 0; + for (int j = 0; j < SZ; ++j) { + if (j % 3 == 0) + o += *(y[j]); + } + + simd inc(0, 1); + block_store(output, inc + o); + }); }); e.wait(); } @@ -125,7 +124,7 @@ int main(void) { x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; } - int* y[SZ]; + int *y[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offy) % SZ; y[j] = j % 6 == 0 ? x1 + idx : x2 + idx; diff --git a/SYCL/ESIMD/tpm_pointer_v2.cpp b/SYCL/ESIMD/tpm_pointer_v2.cpp index 3a5f2aa092..d2843c7243 100644 --- a/SYCL/ESIMD/tpm_pointer_v2.cpp +++ b/SYCL/ESIMD/tpm_pointer_v2.cpp @@ -30,15 +30,14 @@ int main(void) { std::cout << "Running on " << dev.get_info() << "\n"; auto ctx = q.get_context(); - int *output = static_cast( - malloc_shared(VL * sizeof(int), dev, ctx)); + int *output = static_cast(malloc_shared(VL * sizeof(int), dev, ctx)); memset(output, 0, VL * sizeof(int)); int offx1 = 111; int offx2 = 55; int offy1 = 499; int offy2 = 223; - int offz = 99; + int offz = 99; int base1 = 500; int base2 = 100; int divisor = 4; @@ -46,87 +45,87 @@ int main(void) { { auto e = q.submit([&](handler &cgh) { cgh.parallel_for( - sycl::range<1> {1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx2) % SZ; - x2[idx] = base2 << (j % 32); - } - - int* y1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy1) % SZ; - y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; - } - - int* y2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy2) % SZ; - y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; - } - - int** z[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offz) % SZ; - z[j] = y1 + idx; - } - - // some work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - // some work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) < base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - // some work with Y1 - for (int j = 0; j < SZ; j += 2) { - if (*(y1[j]) > *(y1[j + 1])) - *(y1[j]) = *(y1[j + 1]) - *(y1[j]); - } - - // some work with Y2 - for (int j = 1; j < SZ - 1; j += 2) { - if ((*(y2[j]) <= *(y2[j + 1]))) { - auto temp = y2[j]; - y2[j] = y2[j + 1]; - y2[j + 1] = temp; + sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx2) % SZ; + x2[idx] = base2 << (j % 32); + } + + int *y1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy1) % SZ; + y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; + } + + int *y2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy2) % SZ; + y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; + } + + int **z[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offz) % SZ; + z[j] = y1 + idx; + } + + // some work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; } - } - // some work with Z - for (int j = 0; j < SZ - 1; ++j) { + // some work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) < base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + // some work with Y1 + for (int j = 0; j < SZ; j += 2) { + if (*(y1[j]) > *(y1[j + 1])) + *(y1[j]) = *(y1[j + 1]) - *(y1[j]); + } + + // some work with Y2 + for (int j = 1; j < SZ - 1; j += 2) { + if ((*(y2[j]) <= *(y2[j + 1]))) { + auto temp = y2[j]; + y2[j] = y2[j + 1]; + y2[j + 1] = temp; + } + } + + // some work with Z + for (int j = 0; j < SZ - 1; ++j) { <<<<<<< HEAD - if ( *(*(z[j])) < *(*(z[j + 1])) ) + if (*(*(z[j])) < *(*(z[j + 1]))) ======= if (*(*(z[j])) < *(*(z[j + 1]))) >>>>>>> fe7b8274 (ESIMD: add TPM tests) - z[j] = y2 + j; - if (j % 18 == 0) - (*(*(z[j])))++; - } - - int o = 0; - for (int j = 0; j < SZ; ++j) { - if (j % 3 == 0) - o += *(*(z[j])); - } - - simd inc(0, 1); - block_store(output, inc + o); - }); + z[j] = y2 + j; + if (j % 18 == 0) + (*(*(z[j])))++; + } + + int o = 0; + for (int j = 0; j < SZ; ++j) { + if (j % 3 == 0) + o += *(*(z[j])); + } + + simd inc(0, 1); + block_store(output, inc + o); + }); }); e.wait(); } @@ -155,7 +154,7 @@ int main(void) { x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; } - int* y1[SZ]; + int *y1[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offy1) % SZ; y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; @@ -167,7 +166,7 @@ int main(void) { *(y1[j]) = *(y1[j + 1]) - *(y1[j]); } - int* y2[SZ]; + int *y2[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offy2) % SZ; y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; @@ -182,7 +181,7 @@ int main(void) { } } - int** z[SZ]; + int **z[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offz) % SZ; z[j] = y1 + idx; @@ -191,7 +190,7 @@ int main(void) { // same work with Z for (int j = 0; j < SZ - 1; ++j) { <<<<<<< HEAD - if ( *(*(z[j])) < *(*(z[j + 1])) ) + if (*(*(z[j])) < *(*(z[j + 1]))) ======= if (*(*(z[j])) < *(*(z[j + 1]))) >>>>>>> fe7b8274 (ESIMD: add TPM tests) From 68ee5ad623598b61f55b3985c3a776803beff9b6 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 8 Dec 2020 15:51:34 +0300 Subject: [PATCH 04/32] [SYCL][ESIMD] typo fix --- SYCL/ESIMD/tpm_pointer_v2.cpp | 8 -------- 1 file changed, 8 deletions(-) diff --git a/SYCL/ESIMD/tpm_pointer_v2.cpp b/SYCL/ESIMD/tpm_pointer_v2.cpp index d2843c7243..d1c4429e29 100644 --- a/SYCL/ESIMD/tpm_pointer_v2.cpp +++ b/SYCL/ESIMD/tpm_pointer_v2.cpp @@ -107,11 +107,7 @@ int main(void) { // some work with Z for (int j = 0; j < SZ - 1; ++j) { -<<<<<<< HEAD if (*(*(z[j])) < *(*(z[j + 1]))) -======= - if (*(*(z[j])) < *(*(z[j + 1]))) ->>>>>>> fe7b8274 (ESIMD: add TPM tests) z[j] = y2 + j; if (j % 18 == 0) (*(*(z[j])))++; @@ -189,11 +185,7 @@ int main(void) { // same work with Z for (int j = 0; j < SZ - 1; ++j) { -<<<<<<< HEAD if (*(*(z[j])) < *(*(z[j + 1]))) -======= - if (*(*(z[j])) < *(*(z[j + 1]))) ->>>>>>> fe7b8274 (ESIMD: add TPM tests) z[j] = y2 + j; if (j % 18 == 0) (*(*(z[j])))++; From 740856830cd475575f4f949f24232c43844bc9fc Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 8 Dec 2020 18:11:08 +0300 Subject: [PATCH 05/32] [SYCL][ESIMD] improve TPM tests self-check --- SYCL/ESIMD/tpm_basic.cpp | 30 +++++++++--------------------- SYCL/ESIMD/tpm_pointer.cpp | 26 +++++++++----------------- SYCL/ESIMD/tpm_pointer_v2.cpp | 26 +++++++++----------------- 3 files changed, 27 insertions(+), 55 deletions(-) diff --git a/SYCL/ESIMD/tpm_basic.cpp b/SYCL/ESIMD/tpm_basic.cpp index cf2bffcb27..a818cd519f 100644 --- a/SYCL/ESIMD/tpm_basic.cpp +++ b/SYCL/ESIMD/tpm_basic.cpp @@ -69,16 +69,10 @@ int main(void) { x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; } - int o = 0; - for (int j = 0; j < SZ; ++j) { - if (j % 3 == 0) - o += x1[j]; - else - o -= x2[j]; - } - - simd inc(0, 1); - block_store(output, inc + o); + simd val(0); + for (int j = 0; j < SZ; ++j) + val.select<1, 1>(j % VL) += x1[j] - x2[j]; + block_store(output, val); }); }); e.wait(); @@ -108,22 +102,16 @@ int main(void) { x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; } - int o = 0; - for (int j = 0; j < SZ; ++j) { - if (j % 3 == 0) - o += x1[j]; - else - o -= x2[j]; - } + int o[VL] = {0}; + for (int j = 0; j < SZ; ++j) + o[j % VL] += x1[j] - x2[j]; int err_cnt = 0; - for (int j = 0; j < VL; ++j) { - if (output[j] != (o + j)) + for (int j = 0; j < VL; ++j) + if (output[j] != o[j]) err_cnt += 1; - } if (err_cnt > 0) { - std::cout << "GPU: " << output[0] << " vs CPU: " << o << "\n"; std::cout << "FAILED.\n"; return 1; } diff --git a/SYCL/ESIMD/tpm_pointer.cpp b/SYCL/ESIMD/tpm_pointer.cpp index b19ad1d319..98799177bc 100644 --- a/SYCL/ESIMD/tpm_pointer.cpp +++ b/SYCL/ESIMD/tpm_pointer.cpp @@ -87,14 +87,10 @@ int main(void) { *(y[j]) = *(y[j + 1]) - *(y[j]); } - int o = 0; - for (int j = 0; j < SZ; ++j) { - if (j % 3 == 0) - o += *(y[j]); - } - - simd inc(0, 1); - block_store(output, inc + o); + simd val(0); + for (int j = 0; j < SZ; ++j) + val.select<1, 1>(j % VL) += *(y[j]); + block_store(output, val); }); }); e.wait(); @@ -141,20 +137,16 @@ int main(void) { *(y[j]) = *(y[j + 1]) - *(y[j]); } - int o = 0; - for (int j = 0; j < SZ; ++j) { - if (j % 3 == 0) - o += *(y[j]); - } + int o[VL] = {0}; + for (int j = 0; j < SZ; ++j) + o[j % VL] += *(y[j]); int err_cnt = 0; - for (int j = 0; j < VL; ++j) { - if (output[j] != (o + j)) + for (int j = 0; j < VL; ++j) + if (output[j] != o[j]) err_cnt += 1; - } if (err_cnt > 0) { - std::cout << "GPU: " << output[0] << " vs CPU: " << o << "\n"; std::cout << "FAILED.\n"; return 1; } diff --git a/SYCL/ESIMD/tpm_pointer_v2.cpp b/SYCL/ESIMD/tpm_pointer_v2.cpp index d1c4429e29..b30760150a 100644 --- a/SYCL/ESIMD/tpm_pointer_v2.cpp +++ b/SYCL/ESIMD/tpm_pointer_v2.cpp @@ -113,14 +113,10 @@ int main(void) { (*(*(z[j])))++; } - int o = 0; - for (int j = 0; j < SZ; ++j) { - if (j % 3 == 0) - o += *(*(z[j])); - } - - simd inc(0, 1); - block_store(output, inc + o); + simd val(0); + for (int j = 0; j < SZ; ++j) + val.select<1, 1>(j % VL) += *(*(z[j])); + block_store(output, val); }); }); e.wait(); @@ -191,20 +187,16 @@ int main(void) { (*(*(z[j])))++; } - int o = 0; - for (int j = 0; j < SZ; ++j) { - if (j % 3 == 0) - o += *(*(z[j])); - } + int o[VL] = {0}; + for (int j = 0; j < SZ; ++j) + o[j % VL] += *(*(z[j])); int err_cnt = 0; - for (int j = 0; j < VL; ++j) { - if (output[j] != (o + j)) + for (int j = 0; j < VL; ++j) + if (output[j] != o[j]) err_cnt += 1; - } if (err_cnt > 0) { - std::cout << "GPU: " << output[0] << " vs CPU: " << o << "\n"; std::cout << "FAILED.\n"; return 1; } From 6c817eeb563d83fc8176cf79633b97e1e985171c Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Wed, 16 Dec 2020 19:32:08 +0300 Subject: [PATCH 06/32] [SYCL][ESIMD] add description to TPM tests; cosmetic changes --- SYCL/ESIMD/tpm_basic.cpp | 10 +++++++--- SYCL/ESIMD/tpm_pointer.cpp | 10 +++++++--- SYCL/ESIMD/tpm_pointer_v2.cpp | 12 +++++++++--- 3 files changed, 23 insertions(+), 9 deletions(-) diff --git a/SYCL/ESIMD/tpm_basic.cpp b/SYCL/ESIMD/tpm_basic.cpp index a818cd519f..d982be17f6 100644 --- a/SYCL/ESIMD/tpm_basic.cpp +++ b/SYCL/ESIMD/tpm_basic.cpp @@ -1,5 +1,4 @@ -//==---------------- basic_tpm.cpp - DPC++ ESIMD on-device test -//------------==// +//==--------------- tpm_basic.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. @@ -9,9 +8,12 @@ // TODO enable on Windows and Level Zero // REQUIRES: linux && gpu && opencl // RUN: %clangxx-esimd -fsycl %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// This test is intended to use Thread Private Memory (TPM) to support +// implementation in ESIMD backend. In order to force using of TPM need to +// allocate 96x32 bytes or more. + #include "esimd_test_utils.hpp" #include @@ -111,6 +113,8 @@ int main(void) { if (output[j] != o[j]) err_cnt += 1; + free(output, ctx); + if (err_cnt > 0) { std::cout << "FAILED.\n"; return 1; diff --git a/SYCL/ESIMD/tpm_pointer.cpp b/SYCL/ESIMD/tpm_pointer.cpp index 98799177bc..e42acecae7 100644 --- a/SYCL/ESIMD/tpm_pointer.cpp +++ b/SYCL/ESIMD/tpm_pointer.cpp @@ -1,5 +1,4 @@ -//==---------------- pointer_tpm.cpp - DPC++ ESIMD on-device test -//------------==// +//==--------------- tpm_pointer.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. @@ -9,9 +8,12 @@ // TODO enable on Windows and Level Zero // REQUIRES: linux && gpu && opencl // RUN: %clangxx-esimd -fsycl %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// This test is intended to use pointer to Thread Private Memory (TPM) located +// in TPM to support implementation in ESIMD backend. In order to force using +// of TPM need to allocate 96x32 bytes or more. + #include "esimd_test_utils.hpp" #include @@ -146,6 +148,8 @@ int main(void) { if (output[j] != o[j]) err_cnt += 1; + free(output, ctx); + if (err_cnt > 0) { std::cout << "FAILED.\n"; return 1; diff --git a/SYCL/ESIMD/tpm_pointer_v2.cpp b/SYCL/ESIMD/tpm_pointer_v2.cpp index b30760150a..19f77d4888 100644 --- a/SYCL/ESIMD/tpm_pointer_v2.cpp +++ b/SYCL/ESIMD/tpm_pointer_v2.cpp @@ -1,5 +1,4 @@ -//==---------------- pointer_tpm_v2.cpp - DPC++ ESIMD on-device test -//------------==// +//==--------------- tpm_pointer_v2.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. @@ -9,9 +8,14 @@ // TODO enable on Windows and Level Zero // REQUIRES: linux && gpu && opencl // RUN: %clangxx-esimd -fsycl %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// This test is intended to use pointer to Thread Private Memory (TPM) located +// in TPM to support implementation in ESIMD backend. In order to force using +// of TPM need to allocate 96x32 bytes or more. +// This test uses pointer to pointer in TPM to trigger some thresholds in +// backend. + #include "esimd_test_utils.hpp" #include @@ -196,6 +200,8 @@ int main(void) { if (output[j] != o[j]) err_cnt += 1; + free(output, ctx); + if (err_cnt > 0) { std::cout << "FAILED.\n"; return 1; From 8823c79e8136e25fcee413b9d9f87ac72f9fee7c Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Thu, 17 Dec 2020 23:17:11 +0300 Subject: [PATCH 07/32] [SYCL][ESIMD] merged tests to one with 3 cases --- SYCL/ESIMD/tpm_basic.cpp | 125 ----------------- SYCL/ESIMD/tpm_pointer.cpp | 160 ---------------------- SYCL/ESIMD/tpm_pointer_v2.cpp | 212 ----------------------------- SYCL/ESIMD/tpm_tests.cpp | 244 ++++++++++++++++++++++++++++++++++ 4 files changed, 244 insertions(+), 497 deletions(-) delete mode 100644 SYCL/ESIMD/tpm_basic.cpp delete mode 100644 SYCL/ESIMD/tpm_pointer.cpp delete mode 100644 SYCL/ESIMD/tpm_pointer_v2.cpp create mode 100644 SYCL/ESIMD/tpm_tests.cpp diff --git a/SYCL/ESIMD/tpm_basic.cpp b/SYCL/ESIMD/tpm_basic.cpp deleted file mode 100644 index d982be17f6..0000000000 --- a/SYCL/ESIMD/tpm_basic.cpp +++ /dev/null @@ -1,125 +0,0 @@ -//==--------------- tpm_basic.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 -// -//===----------------------------------------------------------------------===// -// TODO enable on Windows and Level Zero -// REQUIRES: linux && gpu && opencl -// RUN: %clangxx-esimd -fsycl %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out - -// This test is intended to use Thread Private Memory (TPM) to support -// implementation in ESIMD backend. In order to force using of TPM need to -// allocate 96x32 bytes or more. - -#include "esimd_test_utils.hpp" - -#include -#include -#include - -using namespace cl::sycl; - -int main(void) { - constexpr unsigned VL = 8; - constexpr unsigned SZ = 800; // big enough to use TPM - - 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(); - - int *output = static_cast(malloc_shared(VL * sizeof(int), dev, ctx)); - memset(output, 0, VL * sizeof(int)); - - int off1 = 16; - int off2 = 128; - int base1 = 500; - int base2 = 100; - int divisor = 4; - - { - auto e = q.submit([&](handler &cgh) { - cgh.parallel_for( - sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + off1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + off2) % SZ; - x2[idx] = base2 << (j % 32); - } - - // some work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - // some work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) > base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - simd val(0); - for (int j = 0; j < SZ; ++j) - val.select<1, 1>(j % VL) += x1[j] - x2[j]; - block_store(output, val); - }); - }); - e.wait(); - } - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + off1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - // same work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + off2) % SZ; - x2[idx] = base2 << (j % 32); - } - - // same work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) > base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - int o[VL] = {0}; - for (int j = 0; j < SZ; ++j) - o[j % VL] += x1[j] - x2[j]; - - int err_cnt = 0; - for (int j = 0; j < VL; ++j) - if (output[j] != o[j]) - err_cnt += 1; - - free(output, ctx); - - if (err_cnt > 0) { - std::cout << "FAILED.\n"; - return 1; - } - - std::cout << "Passed\n"; - return 0; -} diff --git a/SYCL/ESIMD/tpm_pointer.cpp b/SYCL/ESIMD/tpm_pointer.cpp deleted file mode 100644 index e42acecae7..0000000000 --- a/SYCL/ESIMD/tpm_pointer.cpp +++ /dev/null @@ -1,160 +0,0 @@ -//==--------------- tpm_pointer.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 -// -//===----------------------------------------------------------------------===// -// TODO enable on Windows and Level Zero -// REQUIRES: linux && gpu && opencl -// RUN: %clangxx-esimd -fsycl %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out - -// This test is intended to use pointer to Thread Private Memory (TPM) located -// in TPM to support implementation in ESIMD backend. In order to force using -// of TPM need to allocate 96x32 bytes or more. - -#include "esimd_test_utils.hpp" - -#include -#include -#include - -using namespace cl::sycl; - -int main(void) { - constexpr unsigned VL = 8; - constexpr unsigned SZ = 800; // big enough to use TPM - - 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(); - - int *output = static_cast(malloc_shared(VL * sizeof(int), dev, ctx)); - memset(output, 0, VL * sizeof(int)); - - int offx1 = 55; - int offx2 = 11; - int offy = 111; - int base1 = 500; - int base2 = 100; - int divisor = 4; - - { - auto e = q.submit([&](handler &cgh) { - cgh.parallel_for( - sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx2) % SZ; - x2[idx] = base2 << (j % 32); - } - - int *y[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy) % SZ; - y[j] = j % 6 == 0 ? x1 + idx : x2 + idx; - } - - // some work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - // some work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) > base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - // some work with Y - for (int j = 0; j < SZ; j += 2) { - if ((j % 6 != 0) && (y[j] > y[j + 1])) { - auto temp = y[j]; - y[j] = y[j + 1]; - y[j + 1] = temp; - } - if (*(y[j]) > *(y[j + 1])) - *(y[j]) = *(y[j + 1]) - *(y[j]); - } - - simd val(0); - for (int j = 0; j < SZ; ++j) - val.select<1, 1>(j % VL) += *(y[j]); - block_store(output, val); - }); - }); - e.wait(); - } - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - // same work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx2) % SZ; - x2[idx] = base2 << (j % 32); - } - - // same work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) > base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - int *y[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy) % SZ; - y[j] = j % 6 == 0 ? x1 + idx : x2 + idx; - } - - // same work with Y - for (int j = 0; j < SZ; j += 2) { - if ((j % 6 != 0) && (y[j] > y[j + 1])) { - auto temp = y[j]; - y[j] = y[j + 1]; - y[j + 1] = temp; - } - if (*(y[j]) > *(y[j + 1])) - *(y[j]) = *(y[j + 1]) - *(y[j]); - } - - int o[VL] = {0}; - for (int j = 0; j < SZ; ++j) - o[j % VL] += *(y[j]); - - int err_cnt = 0; - for (int j = 0; j < VL; ++j) - if (output[j] != o[j]) - err_cnt += 1; - - free(output, ctx); - - if (err_cnt > 0) { - std::cout << "FAILED.\n"; - return 1; - } - - std::cout << "Passed\n"; - return 0; -} diff --git a/SYCL/ESIMD/tpm_pointer_v2.cpp b/SYCL/ESIMD/tpm_pointer_v2.cpp deleted file mode 100644 index 19f77d4888..0000000000 --- a/SYCL/ESIMD/tpm_pointer_v2.cpp +++ /dev/null @@ -1,212 +0,0 @@ -//==--------------- tpm_pointer_v2.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 -// -//===----------------------------------------------------------------------===// -// TODO enable on Windows and Level Zero -// REQUIRES: linux && gpu && opencl -// RUN: %clangxx-esimd -fsycl %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out - -// This test is intended to use pointer to Thread Private Memory (TPM) located -// in TPM to support implementation in ESIMD backend. In order to force using -// of TPM need to allocate 96x32 bytes or more. -// This test uses pointer to pointer in TPM to trigger some thresholds in -// backend. - -#include "esimd_test_utils.hpp" - -#include -#include -#include - -using namespace cl::sycl; - -int main(void) { - constexpr unsigned VL = 8; - constexpr unsigned SZ = 800; // big enough to use TPM - - 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(); - - int *output = static_cast(malloc_shared(VL * sizeof(int), dev, ctx)); - memset(output, 0, VL * sizeof(int)); - - int offx1 = 111; - int offx2 = 55; - int offy1 = 499; - int offy2 = 223; - int offz = 99; - int base1 = 500; - int base2 = 100; - int divisor = 4; - - { - auto e = q.submit([&](handler &cgh) { - cgh.parallel_for( - sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx2) % SZ; - x2[idx] = base2 << (j % 32); - } - - int *y1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy1) % SZ; - y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; - } - - int *y2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy2) % SZ; - y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; - } - - int **z[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offz) % SZ; - z[j] = y1 + idx; - } - - // some work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - // some work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) < base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - // some work with Y1 - for (int j = 0; j < SZ; j += 2) { - if (*(y1[j]) > *(y1[j + 1])) - *(y1[j]) = *(y1[j + 1]) - *(y1[j]); - } - - // some work with Y2 - for (int j = 1; j < SZ - 1; j += 2) { - if ((*(y2[j]) <= *(y2[j + 1]))) { - auto temp = y2[j]; - y2[j] = y2[j + 1]; - y2[j + 1] = temp; - } - } - - // some work with Z - for (int j = 0; j < SZ - 1; ++j) { - if (*(*(z[j])) < *(*(z[j + 1]))) - z[j] = y2 + j; - if (j % 18 == 0) - (*(*(z[j])))++; - } - - simd val(0); - for (int j = 0; j < SZ; ++j) - val.select<1, 1>(j % VL) += *(*(z[j])); - block_store(output, val); - }); - }); - e.wait(); - } - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - // same work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx2) % SZ; - x2[idx] = base2 << (j % 32); - } - - // same work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) < base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - int *y1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy1) % SZ; - y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; - } - - // same work with Y1 - for (int j = 0; j < SZ; j += 2) { - if (*(y1[j]) > *(y1[j + 1])) - *(y1[j]) = *(y1[j + 1]) - *(y1[j]); - } - - int *y2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy2) % SZ; - y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; - } - - // same work with Y2 - for (int j = 1; j < SZ - 1; j += 2) { - if ((*(y2[j]) <= *(y2[j + 1]))) { - auto temp = y2[j]; - y2[j] = y2[j + 1]; - y2[j + 1] = temp; - } - } - - int **z[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offz) % SZ; - z[j] = y1 + idx; - } - - // same work with Z - for (int j = 0; j < SZ - 1; ++j) { - if (*(*(z[j])) < *(*(z[j + 1]))) - z[j] = y2 + j; - if (j % 18 == 0) - (*(*(z[j])))++; - } - - int o[VL] = {0}; - for (int j = 0; j < SZ; ++j) - o[j % VL] += *(*(z[j])); - - int err_cnt = 0; - for (int j = 0; j < VL; ++j) - if (output[j] != o[j]) - err_cnt += 1; - - free(output, ctx); - - if (err_cnt > 0) { - std::cout << "FAILED.\n"; - return 1; - } - - std::cout << "Passed\n"; - return 0; -} diff --git a/SYCL/ESIMD/tpm_tests.cpp b/SYCL/ESIMD/tpm_tests.cpp new file mode 100644 index 0000000000..8795e3f2b5 --- /dev/null +++ b/SYCL/ESIMD/tpm_tests.cpp @@ -0,0 +1,244 @@ +//==--------------- tpm_pointer_v2.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 +// +//===----------------------------------------------------------------------===// +// TODO enable on Windows and Level Zero +// REQUIRES: linux && gpu && opencl +// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 1 +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 2 +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 3 + +// Since in ESIMD a single WI occupies entire underlying H/W thread, SYCL +// private memory maps to what's known as 'thread private memory' in CM. +// This test is intended to use TPM to support implementation in ESIMD +// backend. In order to force using of TPM need to allocate 96x32 bytes or more. + +#include "esimd_test_utils.hpp" + +#include +#include +#include + +using namespace cl::sycl; + +int main(int argc, char **argv) { + constexpr unsigned VL = 8; + constexpr unsigned SZ = 800; // big enough to use TPM + + if (argc != 2) { + std::cout << "Skipped! Specify case number" << std::endl; + return 1; + } + + 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(); + + int *output = static_cast( + sycl::malloc_shared(VL * sizeof(int), dev, ctx)); + memset(output, 0, VL * sizeof(int)); + + int case_num = atoi(argv[1]); + std::cout << "CASE NUM: " << case_num << std::endl; + + int offx1 = 111; + int offx2 = 55; + int offy1 = 499; + int offy2 = 223; + int offz = 99; + int base1 = 500; + int base2 = 100; + int divisor = 4; + + { + auto e = q.submit([&](handler &cgh) { + cgh.parallel_for( + sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + simd val(0); + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx2) % SZ; + x2[idx] = base2 << (j % 32); + } + + // some work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; + } + + // some work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) < base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + if (case_num == 1) { + for (int j = 0; j < SZ; ++j) + val.select<1, 1>(j % VL) += x1[j] - x2[j]; + } else { + int *y1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy1) % SZ; + y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; + } + + int *y2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy2) % SZ; + y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; + } + + // some work with Y1 + for (int j = 0; j < SZ; j += 2) { + if (*(y1[j]) > *(y1[j + 1])) + *(y1[j]) = *(y1[j + 1]) - *(y1[j]); + } + + // some work with Y2 + for (int j = 1; j < SZ - 1; j += 2) { + if ((*(y2[j]) <= *(y2[j + 1]))) { + auto temp = y2[j]; + y2[j] = y2[j + 1]; + y2[j + 1] = temp; + } + } + + if (case_num == 2) { + for (int j = 0; j < SZ; ++j) + val.select<1, 1>(j % VL) += *(y1[j]) - *(y2[j]); + } else { // case_num == 3 + int **z[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offz) % SZ; + z[j] = y1 + idx; + } + + // some work with Z + for (int j = 0; j < SZ - 1; ++j) { + if (*(*(z[j])) < *(*(z[j + 1]))) + z[j] = y2 + j; + if (j % 18 == 0) + (*(*(z[j])))++; + } + + for (int j = 0; j < SZ; ++j) + val.select<1, 1>(j % VL) += *(*(z[j])); + } + } + + block_store(output, val); + }); + }); + e.wait(); + } + + int o[VL] = {0}; + + int x1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx1) % SZ; + x1[idx] = (idx % 2) == 0 ? j : base1; + } + + int x2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offx2) % SZ; + x2[idx] = base2 << (j % 32); + } + + // some work with X1 + for (int j = 1; j < SZ; ++j) { + if ((x1[j] + j) > base1) + x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; + } + + // some work with X2 + for (int j = 1; j < SZ; ++j) { + if ((x2[j] + j) < base2) + x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; + } + + if (case_num == 1) { + for (int j = 0; j < SZ; ++j) + o[j % VL] += x1[j] - x2[j]; + } else { + int *y1[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy1) % SZ; + y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; + } + + int *y2[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offy2) % SZ; + y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; + } + + // some work with Y1 + for (int j = 0; j < SZ; j += 2) { + if (*(y1[j]) > *(y1[j + 1])) + *(y1[j]) = *(y1[j + 1]) - *(y1[j]); + } + + // some work with Y2 + for (int j = 1; j < SZ - 1; j += 2) { + if ((*(y2[j]) <= *(y2[j + 1]))) { + auto temp = y2[j]; + y2[j] = y2[j + 1]; + y2[j + 1] = temp; + } + } + + if (case_num == 2) { + for (int j = 0; j < SZ; ++j) + o[j % VL] += *(y1[j]) - *(y2[j]); + } else { // case_num == 3 + int **z[SZ]; + for (int j = 0; j < SZ; ++j) { + int idx = (j + offz) % SZ; + z[j] = y1 + idx; + } + + // some work with Z + for (int j = 0; j < SZ - 1; ++j) { + if (*(*(z[j])) < *(*(z[j + 1]))) + z[j] = y2 + j; + if (j % 18 == 0) + (*(*(z[j])))++; + } + + for (int j = 0; j < SZ; ++j) + o[j % VL] += *(*(z[j])); + } + } + + int err_cnt = 0; + for (int j = 0; j < VL; ++j) + if (output[j] != o[j]) + err_cnt += 1; + + sycl::free(output, ctx); + + if (err_cnt > 0) { + std::cout << "FAILED.\n"; + return 1; + } + + std::cout << "Passed\n"; + return 0; +} From f62acbc49fad9d38e99d61d28191a619515376e3 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Thu, 17 Dec 2020 23:31:49 +0300 Subject: [PATCH 08/32] clang-format patch --- SYCL/ESIMD/tpm_tests.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/SYCL/ESIMD/tpm_tests.cpp b/SYCL/ESIMD/tpm_tests.cpp index 8795e3f2b5..478eaffcf2 100644 --- a/SYCL/ESIMD/tpm_tests.cpp +++ b/SYCL/ESIMD/tpm_tests.cpp @@ -40,8 +40,8 @@ int main(int argc, char **argv) { std::cout << "Running on " << dev.get_info() << "\n"; auto ctx = q.get_context(); - int *output = static_cast( - sycl::malloc_shared(VL * sizeof(int), dev, ctx)); + int *output = + static_cast(sycl::malloc_shared(VL * sizeof(int), dev, ctx)); memset(output, 0, VL * sizeof(int)); int case_num = atoi(argv[1]); @@ -121,7 +121,7 @@ int main(int argc, char **argv) { if (case_num == 2) { for (int j = 0; j < SZ; ++j) val.select<1, 1>(j % VL) += *(y1[j]) - *(y2[j]); - } else { // case_num == 3 + } else { // case_num == 3 int **z[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offz) % SZ; @@ -207,7 +207,7 @@ int main(int argc, char **argv) { if (case_num == 2) { for (int j = 0; j < SZ; ++j) o[j % VL] += *(y1[j]) - *(y2[j]); - } else { // case_num == 3 + } else { // case_num == 3 int **z[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offz) % SZ; From ad93c1a3c05169213937b2b8586cff9cd0ffe4ab Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Thu, 17 Dec 2020 23:34:38 +0300 Subject: [PATCH 09/32] cosmetic changes --- SYCL/ESIMD/tpm_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/ESIMD/tpm_tests.cpp b/SYCL/ESIMD/tpm_tests.cpp index 478eaffcf2..2f6afc72ed 100644 --- a/SYCL/ESIMD/tpm_tests.cpp +++ b/SYCL/ESIMD/tpm_tests.cpp @@ -1,4 +1,4 @@ -//==--------------- tpm_pointer_v2.cpp - DPC++ ESIMD on-device test --------==// +//==--------------- tpm_tests.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. From 51319c8f47eca79ec0943bb2b46e3ad6ed0d5c86 Mon Sep 17 00:00:00 2001 From: Fedor Veselovskiy Date: Fri, 18 Dec 2020 16:26:33 +0300 Subject: [PATCH 10/32] Update SYCL/ESIMD/tpm_tests.cpp Co-authored-by: kbobrovs --- SYCL/ESIMD/tpm_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/ESIMD/tpm_tests.cpp b/SYCL/ESIMD/tpm_tests.cpp index 2f6afc72ed..cf7b22159a 100644 --- a/SYCL/ESIMD/tpm_tests.cpp +++ b/SYCL/ESIMD/tpm_tests.cpp @@ -14,7 +14,7 @@ // Since in ESIMD a single WI occupies entire underlying H/W thread, SYCL // private memory maps to what's known as 'thread private memory' in CM. -// This test is intended to use TPM to support implementation in ESIMD +// This test is intended to check TPM support implementation in ESIMD // backend. In order to force using of TPM need to allocate 96x32 bytes or more. #include "esimd_test_utils.hpp" From b14db0781dc06af35e5214a18bd08e01af23f82c Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 18 Dec 2020 20:01:54 +0300 Subject: [PATCH 11/32] reworked and renamed --- ...pm_tests.cpp => private_memory_access.cpp} | 212 +++++++----------- 1 file changed, 77 insertions(+), 135 deletions(-) rename SYCL/ESIMD/{tpm_tests.cpp => private_memory_access.cpp} (60%) diff --git a/SYCL/ESIMD/tpm_tests.cpp b/SYCL/ESIMD/private_memory_access.cpp similarity index 60% rename from SYCL/ESIMD/tpm_tests.cpp rename to SYCL/ESIMD/private_memory_access.cpp index cf7b22159a..9e440cce1f 100644 --- a/SYCL/ESIMD/tpm_tests.cpp +++ b/SYCL/ESIMD/private_memory_access.cpp @@ -1,4 +1,4 @@ -//==--------------- tpm_tests.cpp - DPC++ ESIMD on-device test -------------==// +//==--------------- private_memory_access.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. @@ -25,154 +25,43 @@ using namespace cl::sycl; -int main(int argc, char **argv) { - constexpr unsigned VL = 8; - constexpr unsigned SZ = 800; // big enough to use TPM - - if (argc != 2) { - std::cout << "Skipped! Specify case number" << std::endl; - return 1; - } - - 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(); - - int *output = - static_cast(sycl::malloc_shared(VL * sizeof(int), dev, ctx)); - memset(output, 0, VL * sizeof(int)); - - int case_num = atoi(argv[1]); - std::cout << "CASE NUM: " << case_num << std::endl; - - int offx1 = 111; - int offx2 = 55; - int offy1 = 499; - int offy2 = 223; - int offz = 99; - int base1 = 500; - int base2 = 100; - int divisor = 4; - - { - auto e = q.submit([&](handler &cgh) { - cgh.parallel_for( - sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; - simd val(0); - - int x1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx1) % SZ; - x1[idx] = (idx % 2) == 0 ? j : base1; - } - - int x2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offx2) % SZ; - x2[idx] = base2 << (j % 32); - } - - // some work with X1 - for (int j = 1; j < SZ; ++j) { - if ((x1[j] + j) > base1) - x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; - } - - // some work with X2 - for (int j = 1; j < SZ; ++j) { - if ((x2[j] + j) < base2) - x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; - } - - if (case_num == 1) { - for (int j = 0; j < SZ; ++j) - val.select<1, 1>(j % VL) += x1[j] - x2[j]; - } else { - int *y1[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy1) % SZ; - y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; - } - - int *y2[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offy2) % SZ; - y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; - } - - // some work with Y1 - for (int j = 0; j < SZ; j += 2) { - if (*(y1[j]) > *(y1[j + 1])) - *(y1[j]) = *(y1[j + 1]) - *(y1[j]); - } - - // some work with Y2 - for (int j = 1; j < SZ - 1; j += 2) { - if ((*(y2[j]) <= *(y2[j + 1]))) { - auto temp = y2[j]; - y2[j] = y2[j + 1]; - y2[j + 1] = temp; - } - } - - if (case_num == 2) { - for (int j = 0; j < SZ; ++j) - val.select<1, 1>(j % VL) += *(y1[j]) - *(y2[j]); - } else { // case_num == 3 - int **z[SZ]; - for (int j = 0; j < SZ; ++j) { - int idx = (j + offz) % SZ; - z[j] = y1 + idx; - } - - // some work with Z - for (int j = 0; j < SZ - 1; ++j) { - if (*(*(z[j])) < *(*(z[j + 1]))) - z[j] = y2 + j; - if (j % 18 == 0) - (*(*(z[j])))++; - } - - for (int j = 0; j < SZ; ++j) - val.select<1, 1>(j % VL) += *(*(z[j])); - } - } - - block_store(output, val); - }); - }); - e.wait(); - } - - int o[VL] = {0}; - +constexpr unsigned VL = 8; +constexpr unsigned SZ = 800; // big enough to use TPM + +ESIMD_INLINE void work(int *o, + int case_num, + int offx1, + int offx2, + int offy1, + int offy2, + int offz, + int base1, + int base2, + int divisor) { int x1[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offx1) % SZ; x1[idx] = (idx % 2) == 0 ? j : base1; } - + int x2[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offx2) % SZ; x2[idx] = base2 << (j % 32); } - + // some work with X1 for (int j = 1; j < SZ; ++j) { if ((x1[j] + j) > base1) x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; } - + // some work with X2 for (int j = 1; j < SZ; ++j) { if ((x2[j] + j) < base2) x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; } - + if (case_num == 1) { for (int j = 0; j < SZ; ++j) o[j % VL] += x1[j] - x2[j]; @@ -182,19 +71,19 @@ int main(int argc, char **argv) { int idx = (j + offy1) % SZ; y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; } - + int *y2[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offy2) % SZ; y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; } - + // some work with Y1 for (int j = 0; j < SZ; j += 2) { if (*(y1[j]) > *(y1[j + 1])) *(y1[j]) = *(y1[j + 1]) - *(y1[j]); } - + // some work with Y2 for (int j = 1; j < SZ - 1; j += 2) { if ((*(y2[j]) <= *(y2[j + 1]))) { @@ -203,7 +92,7 @@ int main(int argc, char **argv) { y2[j + 1] = temp; } } - + if (case_num == 2) { for (int j = 0; j < SZ; ++j) o[j % VL] += *(y1[j]) - *(y2[j]); @@ -213,7 +102,7 @@ int main(int argc, char **argv) { int idx = (j + offz) % SZ; z[j] = y1 + idx; } - + // some work with Z for (int j = 0; j < SZ - 1; ++j) { if (*(*(z[j])) < *(*(z[j + 1]))) @@ -221,11 +110,64 @@ int main(int argc, char **argv) { if (j % 18 == 0) (*(*(z[j])))++; } - + for (int j = 0; j < SZ; ++j) o[j % VL] += *(*(z[j])); } } +} + +int main(int argc, char **argv) { + if (argc != 2) { + std::cout << "Skipped! Specify case number" << std::endl; + return 1; + } + + 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(); + + int *output = + static_cast(sycl::malloc_shared(VL * sizeof(int), dev, ctx)); + memset(output, 0, VL * sizeof(int)); + + int case_num = atoi(argv[1]); + std::cout << "CASE NUM: " << case_num << std::endl; + + int offx1 = 111; + int offx2 = 55; + int offy1 = 499; + int offy2 = 223; + int offz = 99; + int base1 = 500; + int base2 = 100; + int divisor = 4; + + { + auto e = q.submit([&](handler &cgh) { + cgh.parallel_for( + sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + + int o[VL] = {0}; + + work(o, case_num, offx1, offx2, offy1, offy2, offz, base1, base2, divisor); + + simd val(0); + for (int j = 0; j < VL; j++) + val.select<1, 1>(j) += o[j]; + + block_store(output, val); + }); + }); + e.wait(); + } + + int o[VL] = {0}; + + work(o, case_num, offx1, offx2, offy1, offy2, offz, base1, base2, divisor); int err_cnt = 0; for (int j = 0; j < VL; ++j) From 9fe01516b2459a14ccc8c6ebcb9c91e7e262558d Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 18 Dec 2020 20:12:56 +0300 Subject: [PATCH 12/32] clang-format patch --- SYCL/ESIMD/private_memory_access.cpp | 53 ++++++++++++---------------- 1 file changed, 23 insertions(+), 30 deletions(-) diff --git a/SYCL/ESIMD/private_memory_access.cpp b/SYCL/ESIMD/private_memory_access.cpp index 9e440cce1f..d1c5b7f3ae 100644 --- a/SYCL/ESIMD/private_memory_access.cpp +++ b/SYCL/ESIMD/private_memory_access.cpp @@ -28,40 +28,32 @@ using namespace cl::sycl; constexpr unsigned VL = 8; constexpr unsigned SZ = 800; // big enough to use TPM -ESIMD_INLINE void work(int *o, - int case_num, - int offx1, - int offx2, - int offy1, - int offy2, - int offz, - int base1, - int base2, - int divisor) { +ESIMD_INLINE void work(int *o, int case_num, int offx1, int offx2, int offy1, + int offy2, int offz, int base1, int base2, int divisor) { int x1[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offx1) % SZ; x1[idx] = (idx % 2) == 0 ? j : base1; } - + int x2[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offx2) % SZ; x2[idx] = base2 << (j % 32); } - + // some work with X1 for (int j = 1; j < SZ; ++j) { if ((x1[j] + j) > base1) x1[j] = (j * (x1[j] + x1[j - 1]) / divisor) - base2; } - + // some work with X2 for (int j = 1; j < SZ; ++j) { if ((x2[j] + j) < base2) x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; } - + if (case_num == 1) { for (int j = 0; j < SZ; ++j) o[j % VL] += x1[j] - x2[j]; @@ -71,19 +63,19 @@ ESIMD_INLINE void work(int *o, int idx = (j + offy1) % SZ; y1[j] = j % 6 == 0 ? x1 + idx : x2 + idx; } - + int *y2[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offy2) % SZ; y2[j] = j % 2 == 0 ? x2 + idx : x1 + idx; } - + // some work with Y1 for (int j = 0; j < SZ; j += 2) { if (*(y1[j]) > *(y1[j + 1])) *(y1[j]) = *(y1[j + 1]) - *(y1[j]); } - + // some work with Y2 for (int j = 1; j < SZ - 1; j += 2) { if ((*(y2[j]) <= *(y2[j + 1]))) { @@ -92,7 +84,7 @@ ESIMD_INLINE void work(int *o, y2[j + 1] = temp; } } - + if (case_num == 2) { for (int j = 0; j < SZ; ++j) o[j % VL] += *(y1[j]) - *(y2[j]); @@ -102,7 +94,7 @@ ESIMD_INLINE void work(int *o, int idx = (j + offz) % SZ; z[j] = y1 + idx; } - + // some work with Z for (int j = 0; j < SZ - 1; ++j) { if (*(*(z[j])) < *(*(z[j + 1]))) @@ -110,7 +102,7 @@ ESIMD_INLINE void work(int *o, if (j % 18 == 0) (*(*(z[j])))++; } - + for (int j = 0; j < SZ; ++j) o[j % VL] += *(*(z[j])); } @@ -147,20 +139,21 @@ int main(int argc, char **argv) { { auto e = q.submit([&](handler &cgh) { - cgh.parallel_for( - sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; + cgh.parallel_for(sycl::range<1>{1}, + [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; - int o[VL] = {0}; + int o[VL] = {0}; - work(o, case_num, offx1, offx2, offy1, offy2, offz, base1, base2, divisor); + work(o, case_num, offx1, offx2, offy1, + offy2, offz, base1, base2, divisor); - simd val(0); - for (int j = 0; j < VL; j++) - val.select<1, 1>(j) += o[j]; + simd val(0); + for (int j = 0; j < VL; j++) + val.select<1, 1>(j) += o[j]; - block_store(output, val); - }); + block_store(output, val); + }); }); e.wait(); } From f019b28c67dc56e589ce189f0bf5b8d0feecf640 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 12 Jan 2021 18:36:07 +0300 Subject: [PATCH 13/32] [SYCL][ESIMD] evaluate condition on compile-time --- .../Inputs/pm_common.cpp} | 69 ++++++++++--------- SYCL/ESIMD/private_memory/pm_access_1.cpp | 12 ++++ SYCL/ESIMD/private_memory/pm_access_2.cpp | 12 ++++ SYCL/ESIMD/private_memory/pm_access_3.cpp | 12 ++++ 4 files changed, 73 insertions(+), 32 deletions(-) rename SYCL/ESIMD/{private_memory_access.cpp => private_memory/Inputs/pm_common.cpp} (71%) create mode 100644 SYCL/ESIMD/private_memory/pm_access_1.cpp create mode 100644 SYCL/ESIMD/private_memory/pm_access_2.cpp create mode 100644 SYCL/ESIMD/private_memory/pm_access_3.cpp diff --git a/SYCL/ESIMD/private_memory_access.cpp b/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp similarity index 71% rename from SYCL/ESIMD/private_memory_access.cpp rename to SYCL/ESIMD/private_memory/Inputs/pm_common.cpp index d1c5b7f3ae..fc15a153ee 100644 --- a/SYCL/ESIMD/private_memory_access.cpp +++ b/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp @@ -1,16 +1,10 @@ -//==--------------- private_memory_access.cpp - DPC++ ESIMD on-device test -==// +//==--------------- pm_common.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 // //===----------------------------------------------------------------------===// -// TODO enable on Windows and Level Zero -// REQUIRES: linux && gpu && opencl -// RUN: %clangxx-esimd -fsycl %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 1 -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 2 -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 3 // Since in ESIMD a single WI occupies entire underlying H/W thread, SYCL // private memory maps to what's known as 'thread private memory' in CM. @@ -28,8 +22,9 @@ using namespace cl::sycl; constexpr unsigned VL = 8; constexpr unsigned SZ = 800; // big enough to use TPM -ESIMD_INLINE void work(int *o, int case_num, int offx1, int offx2, int offy1, - int offy2, int offz, int base1, int base2, int divisor) { +template +ESIMD_INLINE void work(int *o, int offx1, int offx2, int offy1, + int offy2, int offz, int base1, int base2, int divisor) { int x1[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offx1) % SZ; @@ -54,7 +49,7 @@ ESIMD_INLINE void work(int *o, int case_num, int offx1, int offx2, int offy1, x2[j] = (divisor * (x2[j] - x2[j - 1]) / j) + base1; } - if (case_num == 1) { + if constexpr (CASE_NUM == 1) { for (int j = 0; j < SZ; ++j) o[j % VL] += x1[j] - x2[j]; } else { @@ -85,10 +80,12 @@ ESIMD_INLINE void work(int *o, int case_num, int offx1, int offx2, int offy1, } } - if (case_num == 2) { + if constexpr (CASE_NUM == 2) { for (int j = 0; j < SZ; ++j) o[j % VL] += *(y1[j]) - *(y2[j]); - } else { // case_num == 3 + } else { + static_assert(CASE_NUM == 3, "invalid CASE_NUM"); + int **z[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offz) % SZ; @@ -109,12 +106,9 @@ ESIMD_INLINE void work(int *o, int case_num, int offx1, int offx2, int offy1, } } -int main(int argc, char **argv) { - if (argc != 2) { - std::cout << "Skipped! Specify case number" << std::endl; - return 1; - } +template class KernelID; +template int test() { queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()); auto dev = q.get_device(); @@ -125,9 +119,6 @@ int main(int argc, char **argv) { static_cast(sycl::malloc_shared(VL * sizeof(int), dev, ctx)); memset(output, 0, VL * sizeof(int)); - int case_num = atoi(argv[1]); - std::cout << "CASE NUM: " << case_num << std::endl; - int offx1 = 111; int offx2 = 55; int offy1 = 499; @@ -139,28 +130,26 @@ int main(int argc, char **argv) { { auto e = q.submit([&](handler &cgh) { - cgh.parallel_for(sycl::range<1>{1}, - [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; + cgh.parallel_for>(sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; - int o[VL] = {0}; + int o[VL] = {0}; - work(o, case_num, offx1, offx2, offy1, - offy2, offz, base1, base2, divisor); + work(o, offx1, offx2, offy1, offy2, offz, base1, base2, divisor); - simd val(0); - for (int j = 0; j < VL; j++) - val.select<1, 1>(j) += o[j]; + simd val(0); + for (int j = 0; j < VL; j++) + val.select<1, 1>(j) += o[j]; - block_store(output, val); - }); + block_store(output, val); + }); }); e.wait(); } int o[VL] = {0}; - work(o, case_num, offx1, offx2, offy1, offy2, offz, base1, base2, divisor); + work(o, offx1, offx2, offy1, offy2, offz, base1, base2, divisor); int err_cnt = 0; for (int j = 0; j < VL; ++j) @@ -177,3 +166,19 @@ int main(int argc, char **argv) { std::cout << "Passed\n"; return 0; } + +int main(int argc, char **argv) { + if (argc != 2) { + std::cout << "Skipped! Specify case number" << std::endl; + return 1; + } + + int case_num = atoi(argv[1]); + switch (case_num) { + case 1: return test<1>(); + case 2: return test<2>(); + case 3: return test<3>(); + } + std::cerr << "Invalid case number: " << case_num << "\n"; + return 1; +} diff --git a/SYCL/ESIMD/private_memory/pm_access_1.cpp b/SYCL/ESIMD/private_memory/pm_access_1.cpp new file mode 100644 index 0000000000..60eb0914ae --- /dev/null +++ b/SYCL/ESIMD/private_memory/pm_access_1.cpp @@ -0,0 +1,12 @@ +//==--------------- pm_access_1.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 +// +//===----------------------------------------------------------------------===// + +// TODO enable on Windows and Level Zero +// REQUIRES: linux && gpu && opencl +// RUN: %clangxx-esimd -fsycl -I%S/.. %S/Inputs/pm_common.cpp -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 1 diff --git a/SYCL/ESIMD/private_memory/pm_access_2.cpp b/SYCL/ESIMD/private_memory/pm_access_2.cpp new file mode 100644 index 0000000000..3051b245a2 --- /dev/null +++ b/SYCL/ESIMD/private_memory/pm_access_2.cpp @@ -0,0 +1,12 @@ +//==--------------- pm_access_2.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 +// +//===----------------------------------------------------------------------===// + +// TODO enable on Windows and Level Zero +// REQUIRES: linux && gpu && opencl +// RUN: %clangxx-esimd -fsycl -I%S/.. %S/Inputs/pm_common.cpp -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 2 diff --git a/SYCL/ESIMD/private_memory/pm_access_3.cpp b/SYCL/ESIMD/private_memory/pm_access_3.cpp new file mode 100644 index 0000000000..d18745068c --- /dev/null +++ b/SYCL/ESIMD/private_memory/pm_access_3.cpp @@ -0,0 +1,12 @@ +//==--------------- pm_access_3.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 +// +//===----------------------------------------------------------------------===// + +// TODO enable on Windows and Level Zero +// REQUIRES: linux && gpu && opencl +// RUN: %clangxx-esimd -fsycl -I%S/.. %S/Inputs/pm_common.cpp -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 3 From d958a9173ddfc08a158cccc3691a61c641e55432 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 12 Jan 2021 18:47:56 +0300 Subject: [PATCH 14/32] clang-format patch --- .../ESIMD/private_memory/Inputs/pm_common.cpp | 35 +++++++++++-------- 1 file changed, 20 insertions(+), 15 deletions(-) diff --git a/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp b/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp index fc15a153ee..830aa89367 100644 --- a/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp +++ b/SYCL/ESIMD/private_memory/Inputs/pm_common.cpp @@ -22,9 +22,9 @@ using namespace cl::sycl; constexpr unsigned VL = 8; constexpr unsigned SZ = 800; // big enough to use TPM -template -ESIMD_INLINE void work(int *o, int offx1, int offx2, int offy1, - int offy2, int offz, int base1, int base2, int divisor) { +template +ESIMD_INLINE void work(int *o, int offx1, int offx2, int offy1, int offy2, + int offz, int base1, int base2, int divisor) { int x1[SZ]; for (int j = 0; j < SZ; ++j) { int idx = (j + offx1) % SZ; @@ -130,19 +130,21 @@ template int test() { { auto e = q.submit([&](handler &cgh) { - cgh.parallel_for>(sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; + cgh.parallel_for>( + sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; - int o[VL] = {0}; + int o[VL] = {0}; - work(o, offx1, offx2, offy1, offy2, offz, base1, base2, divisor); + work(o, offx1, offx2, offy1, offy2, offz, base1, base2, + divisor); - simd val(0); - for (int j = 0; j < VL; j++) - val.select<1, 1>(j) += o[j]; + simd val(0); + for (int j = 0; j < VL; j++) + val.select<1, 1>(j) += o[j]; - block_store(output, val); - }); + block_store(output, val); + }); }); e.wait(); } @@ -175,9 +177,12 @@ int main(int argc, char **argv) { int case_num = atoi(argv[1]); switch (case_num) { - case 1: return test<1>(); - case 2: return test<2>(); - case 3: return test<3>(); + case 1: + return test<1>(); + case 2: + return test<2>(); + case 3: + return test<3>(); } std::cerr << "Invalid case number: " << case_num << "\n"; return 1; From cdf05660454babb15f7fec7f3be877ff41ff5f5f Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 02:54:42 +0300 Subject: [PATCH 15/32] initial commit of ESIMD function pointer tests --- .../Inputs/fp_args_size_common.hpp | 94 ++++++++++++++++++ .../fp_args_char_int_size_192.cpp | 20 ++++ .../fp_args_char_int_size_256.cpp | 20 ++++ .../fp_args_char_int_size_512.cpp | 20 ++++ .../fp_args_char_int_size_592.cpp | 20 ++++ .../fp_args_size/fp_args_char_int_size_64.cpp | 20 ++++ .../fp_args_size/fp_args_char_int_size_96.cpp | 20 ++++ .../fp_args_size/fp_args_int_size_192.cpp | 20 ++++ .../fp_args_size/fp_args_int_size_256.cpp | 20 ++++ .../fp_args_size/fp_args_int_size_512.cpp | 20 ++++ .../fp_args_size/fp_args_int_size_592.cpp | 20 ++++ .../fp_args_size/fp_args_int_size_64.cpp | 20 ++++ .../fp_args_size/fp_args_int_size_96.cpp | 20 ++++ SYCL/ESIMD/fp_call_from_func.cpp | 70 ++++++++++++++ SYCL/ESIMD/fp_call_recursive.cpp | 81 ++++++++++++++++ SYCL/ESIMD/fp_in_phi.cpp | 95 +++++++++++++++++++ SYCL/ESIMD/fp_in_select.cpp | 80 ++++++++++++++++ 17 files changed, 660 insertions(+) create mode 100644 SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp create mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp create mode 100644 SYCL/ESIMD/fp_call_from_func.cpp create mode 100644 SYCL/ESIMD/fp_call_recursive.cpp create mode 100644 SYCL/ESIMD/fp_in_phi.cpp create mode 100644 SYCL/ESIMD/fp_in_select.cpp diff --git a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp new file mode 100644 index 0000000000..d133b58650 --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp @@ -0,0 +1,94 @@ +//==------- fp_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 use of function pointers 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)); + + { + auto e = 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 foo = &add, simd, simd>; + auto vc = foo(va, vb); + + for(int j = 0; j < ROWS; j++) + block_store(C + j * VL, vc.select(j * VL)); + }); + }); + e.wait(); + } + + 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/fp_args_size/fp_args_char_int_size_192.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp new file mode 100644 index 0000000000..e00d186486 --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 192; + +typedef char a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp new file mode 100644 index 0000000000..8acc6f7837 --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 256; + +typedef char a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp new file mode 100644 index 0000000000..51b9c023da --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 512; + +typedef char a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp new file mode 100644 index 0000000000..ea5beb6643 --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp @@ -0,0 +1,20 @@ +//==------- fp_args_char_int_size_592.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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 592; + +typedef char a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp new file mode 100644 index 0000000000..767a324346 --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 64; + +typedef char a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp new file mode 100644 index 0000000000..681bc3f064 --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 96; + +typedef char a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp new file mode 100644 index 0000000000..54034ab39e --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 192; + +typedef int a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp new file mode 100644 index 0000000000..0414d6348e --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 256; + +typedef int a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp new file mode 100644 index 0000000000..9f9cd122ca --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 512; + +typedef int a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp new file mode 100644 index 0000000000..d0a4010ef2 --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp @@ -0,0 +1,20 @@ +//==------- fp_args_char_int_size_592.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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 592; + +typedef int a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp new file mode 100644 index 0000000000..fd6ba6bed0 --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 64; + +typedef int a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp new file mode 100644 index 0000000000..c620968e57 --- /dev/null +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp @@ -0,0 +1,20 @@ +//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +constexpr unsigned VL = 16; +constexpr unsigned SIZE = 96; + +typedef int a_data_t; +typedef int b_data_t; +typedef int c_data_t; + +#include"Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_call_from_func.cpp b/SYCL/ESIMD/fp_call_from_func.cpp new file mode 100644 index 0000000000..5bc998b95a --- /dev/null +++ b/SYCL/ESIMD/fp_call_from_func.cpp @@ -0,0 +1,70 @@ +//==--------------- fp_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 +// XFAIL: * +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// +// The test checks that ESIMD kernels support use of function pointers from +// 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 int test(AccTy acc, int A, int B) { + using namespace sycl::INTEL::gpu; + + auto foo = &add; + auto res = foo(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; + + { + buffer buf(output, range<1>(1)); + + auto e = 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); + }); + }); + e.wait(); + } + + if (result != (in1 + in2)) { + std::cout << "Failed" << std::endl; + return 1; + } + + return 0; +} diff --git a/SYCL/ESIMD/fp_call_recursive.cpp b/SYCL/ESIMD/fp_call_recursive.cpp new file mode 100644 index 0000000000..9bd1fd16e4 --- /dev/null +++ b/SYCL/ESIMD/fp_call_recursive.cpp @@ -0,0 +1,81 @@ +//==--------------- 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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// +// The test checks that ESIMD kernels support use of function pointers +// recursively. + +#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; + + 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; + + { + buffer buf(output, range<1>(1)); + + auto e = 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 foo = &add; + auto res = foo(in1, in2, in3); + + scalar_store(acc, 0, res); + }); + }); + e.wait(); + } + + 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; +} diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp new file mode 100644 index 0000000000..c063522874 --- /dev/null +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -0,0 +1,95 @@ +//==--------------- fp_in_phi.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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// +// The test checks that ESIMD kernels correctly handle function pointers as +// arguments of LLVM's PHI function. + +#include "esimd_test_utils.hpp" + +#include +#include + +#include +#include + +class KernelID; + +ESIMD_NOINLINE int f1(int x) { + return x + 1; +} + +ESIMD_NOINLINE int f2(int x) { + return x + 2; +} + +ESIMD_NOINLINE int f3(int x) { + return x + 3; +} + +bool test(queue q, bool flag) { + int result = 0; + int *output = &result; + + std::vector Y = {0, 1}; + + int in1 = 233; + int in2 = 1; + + { + buffer o_buf(output, range<1>(1)); + buffer y_buf(Y.data(), Y.size()); + + auto e = q.submit([&](handler &cgh) { + auto o_acc = o_buf.get_access(cgh); + auto y_acc = y_buf.get_access(cgh); + + cgh.parallel_for( + sycl::range<1> {1}, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + using f = int(*)(int); + + f a[] = {f1, f2}; + if (flag) { + a[0] = f3; + scalar_store(y_acc, 0, 2); + } + + auto res = a[0](in1) + a[1](in2); + + scalar_store(o_acc, 0, res); + }); + }); + e.wait(); + } + + int etalon = in1 + (flag ? 3 : 1) + in2 + 2; + + if (result != etalon) { + std::cout << "Failed with result: " << result << std::endl; + return false; + } + + return true; +} + +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"; + + bool passed = true; + passed &= test(q, true); + passed &= test(q, false); + + return passed ? 0 : 1; +} diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp new file mode 100644 index 0000000000..9e0017e6e1 --- /dev/null +++ b/SYCL/ESIMD/fp_in_select.cpp @@ -0,0 +1,80 @@ +//==--------------- fp_in_select.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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// +// The test checks that ESIMD kernels correctly handle function pointers as +// arguments of select function. + +#include "esimd_test_utils.hpp" + +#include +#include + +#include + +class KernelID; + +ESIMD_NOINLINE int add(int a, int b) { + return a + b; +} + +ESIMD_NOINLINE int sub(int a, int b) { + return a - b; +} + +bool test(queue q, bool flag) { + int result = 0; + int *output = &result; + + int in1 = 233; + int in2 = 100; + + { + buffer buf(output, range<1>(1)); + + auto e = 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 foo = flag ? &add : ⊂ + auto res = foo(in1, in2); + + scalar_store(acc, 0, res); + }); + }); + e.wait(); + } + + int etalon = flag ? in1 + in2 : in1 - in2; + + if (result != etalon) { + std::cout << "Failed with result: " << result << std::endl; + return false; + } + + return true; +} + +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"; + + bool passed = true; + passed &= test(q, true); + passed &= test(q, false); + + return passed ? 0 : 1; +} From fb11440ff1d3b27541f5705f023424e62e26d1e8 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 03:13:07 +0300 Subject: [PATCH 16/32] clang-format patch --- .../Inputs/fp_args_size_common.hpp | 32 +++++++++-------- .../fp_args_char_int_size_192.cpp | 2 +- .../fp_args_char_int_size_256.cpp | 2 +- .../fp_args_char_int_size_512.cpp | 2 +- .../fp_args_char_int_size_592.cpp | 2 +- .../fp_args_size/fp_args_char_int_size_64.cpp | 2 +- .../fp_args_size/fp_args_char_int_size_96.cpp | 2 +- .../fp_args_size/fp_args_int_size_192.cpp | 2 +- .../fp_args_size/fp_args_int_size_256.cpp | 2 +- .../fp_args_size/fp_args_int_size_512.cpp | 2 +- .../fp_args_size/fp_args_int_size_592.cpp | 2 +- .../fp_args_size/fp_args_int_size_64.cpp | 2 +- .../fp_args_size/fp_args_int_size_96.cpp | 2 +- SYCL/ESIMD/fp_call_from_func.cpp | 12 +++---- SYCL/ESIMD/fp_call_recursive.cpp | 19 +++++----- SYCL/ESIMD/fp_in_phi.cpp | 36 ++++++++----------- SYCL/ESIMD/fp_in_select.cpp | 22 +++++------- 17 files changed, 69 insertions(+), 76 deletions(-) diff --git a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp index d133b58650..b33b0a50a5 100644 --- a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp +++ b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp @@ -30,7 +30,9 @@ using namespace cl::sycl; class KernelID; template -ESIMD_NOINLINE TC add(TA A, TB B) { return (TC)A + (TC)B; } +ESIMD_NOINLINE TC add(TA A, TB B) { + return (TC)A + (TC)B; +} int main(void) { queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()); @@ -39,18 +41,18 @@ int main(void) { 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_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_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)); + c_data_t *C = static_cast( + sycl::malloc_shared(SIZE * sizeof(c_data_t), dev, ctx)); memset(C, 0, SIZE * sizeof(c_data_t)); { @@ -61,15 +63,16 @@ int main(void) { simd va(0); simd vb(0); - for(int j = 0; j < ROWS; j++) { + 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 foo = &add, simd, simd>; + auto foo = &add, simd, + simd>; auto vc = foo(va, vb); - for(int j = 0; j < ROWS; j++) + for (int j = 0; j < ROWS; j++) block_store(C + j * VL, vc.select(j * VL)); }); }); @@ -77,8 +80,9 @@ int main(void) { } unsigned err_cnt = 0; - for(int i = 0; i < SIZE; i++) - if (C[i] != A[i] + B[i]) err_cnt++; + for (int i = 0; i < SIZE; i++) + if (C[i] != A[i] + B[i]) + err_cnt++; sycl::free(A, ctx); sycl::free(B, ctx); diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp index e00d186486..144c0d62c4 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp @@ -17,4 +17,4 @@ typedef char a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp index 8acc6f7837..03f853fbe8 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp @@ -17,4 +17,4 @@ typedef char a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp index 51b9c023da..a0f4787b4b 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp @@ -17,4 +17,4 @@ typedef char a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp index ea5beb6643..1a33b1c9fb 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp @@ -17,4 +17,4 @@ typedef char a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp index 767a324346..40ee81393b 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp @@ -17,4 +17,4 @@ typedef char a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp index 681bc3f064..f8eb2c0c4c 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp @@ -17,4 +17,4 @@ typedef char a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp index 54034ab39e..1a2f6c2452 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp @@ -17,4 +17,4 @@ typedef int a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp index 0414d6348e..63589134c0 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp @@ -17,4 +17,4 @@ typedef int a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp index 9f9cd122ca..69bf678b65 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp @@ -17,4 +17,4 @@ typedef int a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp index d0a4010ef2..b779fc23a5 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp @@ -17,4 +17,4 @@ typedef int a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp index fd6ba6bed0..eb24a669b0 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp @@ -17,4 +17,4 @@ typedef int a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp index c620968e57..3247026439 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp @@ -17,4 +17,4 @@ typedef int a_data_t; typedef int b_data_t; typedef int c_data_t; -#include"Inputs/fp_args_size_common.hpp" +#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_call_from_func.cpp b/SYCL/ESIMD/fp_call_from_func.cpp index 5bc998b95a..8ff07ac50d 100644 --- a/SYCL/ESIMD/fp_call_from_func.cpp +++ b/SYCL/ESIMD/fp_call_from_func.cpp @@ -25,13 +25,12 @@ class KernelID; ESIMD_NOINLINE int add(int A, int B) { return A + B; } -template -ESIMD_NOINLINE int test(AccTy acc, int A, int B) { +template ESIMD_NOINLINE int test(AccTy acc, int A, int B) { using namespace sycl::INTEL::gpu; - + auto foo = &add; auto res = foo(A, B); - + scalar_store(acc, 0, res); } @@ -54,9 +53,8 @@ int main(int argc, char **argv) { auto acc = buf.get_access(cgh); cgh.parallel_for( - sycl::range<1> {1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - test(acc, in1, in2); - }); + sycl::range<1>{1}, + [=](id<1> i) SYCL_ESIMD_KERNEL { test(acc, in1, in2); }); }); e.wait(); } diff --git a/SYCL/ESIMD/fp_call_recursive.cpp b/SYCL/ESIMD/fp_call_recursive.cpp index 9bd1fd16e4..bc47aeecab 100644 --- a/SYCL/ESIMD/fp_call_recursive.cpp +++ b/SYCL/ESIMD/fp_call_recursive.cpp @@ -23,7 +23,8 @@ class KernelID; ESIMD_NOINLINE unsigned add(unsigned A, unsigned B, unsigned C) { - if (B == 0) return A; + if (B == 0) + return A; auto foo = &add; return (B % C == 0) ? foo(A + 1, B - 1, C) : foo(A - C, B - 2, C); @@ -48,21 +49,21 @@ int main(int argc, char **argv) { auto e = 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; + cgh.parallel_for(sycl::range<1>{1}, + [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; - auto foo = &add; - auto res = foo(in1, in2, in3); + auto foo = &add; + auto res = foo(in1, in2, in3); - scalar_store(acc, 0, res); - }); + scalar_store(acc, 0, res); + }); }); e.wait(); } int etalon = in1; - while(in2 > 0) { + while (in2 > 0) { if (in2 % in3 == 0) { etalon += 1; in2 -= 1; diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp index c063522874..4b37eb3915 100644 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -23,17 +23,11 @@ class KernelID; -ESIMD_NOINLINE int f1(int x) { - return x + 1; -} +ESIMD_NOINLINE int f1(int x) { return x + 1; } -ESIMD_NOINLINE int f2(int x) { - return x + 2; -} +ESIMD_NOINLINE int f2(int x) { return x + 2; } -ESIMD_NOINLINE int f3(int x) { - return x + 3; -} +ESIMD_NOINLINE int f3(int x) { return x + 3; } bool test(queue q, bool flag) { int result = 0; @@ -52,21 +46,21 @@ bool test(queue q, bool flag) { auto o_acc = o_buf.get_access(cgh); auto y_acc = y_buf.get_access(cgh); - cgh.parallel_for( - sycl::range<1> {1}, [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; - using f = int(*)(int); + cgh.parallel_for(sycl::range<1>{1}, + [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; + using f = int (*)(int); - f a[] = {f1, f2}; - if (flag) { - a[0] = f3; - scalar_store(y_acc, 0, 2); - } + f a[] = {f1, f2}; + if (flag) { + a[0] = f3; + scalar_store(y_acc, 0, 2); + } - auto res = a[0](in1) + a[1](in2); + auto res = a[0](in1) + a[1](in2); - scalar_store(o_acc, 0, res); - }); + scalar_store(o_acc, 0, res); + }); }); e.wait(); } diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp index 9e0017e6e1..04af9db691 100644 --- a/SYCL/ESIMD/fp_in_select.cpp +++ b/SYCL/ESIMD/fp_in_select.cpp @@ -22,13 +22,9 @@ class KernelID; -ESIMD_NOINLINE int add(int a, int b) { - return a + b; -} +ESIMD_NOINLINE int add(int a, int b) { return a + b; } -ESIMD_NOINLINE int sub(int a, int b) { - return a - b; -} +ESIMD_NOINLINE int sub(int a, int b) { return a - b; } bool test(queue q, bool flag) { int result = 0; @@ -43,15 +39,15 @@ bool test(queue q, bool flag) { auto e = 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; + cgh.parallel_for(sycl::range<1>{1}, + [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::INTEL::gpu; - auto foo = flag ? &add : ⊂ - auto res = foo(in1, in2); + auto foo = flag ? &add : ⊂ + auto res = foo(in1, in2); - scalar_store(acc, 0, res); - }); + scalar_store(acc, 0, res); + }); }); e.wait(); } From d176d94698b22771dfba6f5da609cc8f9f8df32b Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 12:45:29 +0300 Subject: [PATCH 17/32] remove some tests --- .../fp_args_char_int_size_592.cpp | 20 ------------------- .../fp_args_size/fp_args_int_size_592.cpp | 20 ------------------- 2 files changed, 40 deletions(-) delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp deleted file mode 100644 index 1a33b1c9fb..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_592.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==------- fp_args_char_int_size_592.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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out - -constexpr unsigned VL = 16; -constexpr unsigned SIZE = 592; - -typedef char a_data_t; -typedef int b_data_t; -typedef int c_data_t; - -#include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp deleted file mode 100644 index b779fc23a5..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_592.cpp +++ /dev/null @@ -1,20 +0,0 @@ -//==------- fp_args_char_int_size_592.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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out - -constexpr unsigned VL = 16; -constexpr unsigned SIZE = 592; - -typedef int a_data_t; -typedef int b_data_t; -typedef int c_data_t; - -#include "Inputs/fp_args_size_common.hpp" From fb1a80740d112d91392c5399f9dd6adefd3151b0 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 17:17:32 +0300 Subject: [PATCH 18/32] handling synchronous SYCL exceptions --- .../ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp | 11 ++++++++--- SYCL/ESIMD/fp_call_from_func.cpp | 11 ++++++----- SYCL/ESIMD/fp_call_recursive.cpp | 8 +++++--- SYCL/ESIMD/fp_in_phi.cpp | 8 +++++--- SYCL/ESIMD/fp_in_select.cpp | 8 +++++--- 5 files changed, 29 insertions(+), 17 deletions(-) diff --git a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp index b33b0a50a5..aac404fcd8 100644 --- a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp +++ b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp @@ -55,8 +55,8 @@ int main(void) { sycl::malloc_shared(SIZE * sizeof(c_data_t), dev, ctx)); memset(C, 0, SIZE * sizeof(c_data_t)); - { - auto e = q.submit([&](handler &cgh) { + try { + q.submit([&](handler &cgh) { cgh.parallel_for( sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::INTEL::gpu; @@ -76,7 +76,12 @@ int main(void) { block_store(C + j * VL, vc.select(j * VL)); }); }); - e.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; diff --git a/SYCL/ESIMD/fp_call_from_func.cpp b/SYCL/ESIMD/fp_call_from_func.cpp index 8ff07ac50d..8890553f8d 100644 --- a/SYCL/ESIMD/fp_call_from_func.cpp +++ b/SYCL/ESIMD/fp_call_from_func.cpp @@ -6,7 +6,6 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: * // UNSUPPORTED: cuda // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out @@ -25,7 +24,7 @@ class KernelID; ESIMD_NOINLINE int add(int A, int B) { return A + B; } -template ESIMD_NOINLINE int test(AccTy acc, int A, int B) { +template ESIMD_NOINLINE void test(AccTy acc, int A, int B) { using namespace sycl::INTEL::gpu; auto foo = &add; @@ -46,17 +45,19 @@ int main(int argc, char **argv) { int in1 = 100; int in2 = 233; - { + try { buffer buf(output, range<1>(1)); - auto e = q.submit([&](handler &cgh) { + 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); }); }); - e.wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return e.get_cl_code(); } if (result != (in1 + in2)) { diff --git a/SYCL/ESIMD/fp_call_recursive.cpp b/SYCL/ESIMD/fp_call_recursive.cpp index bc47aeecab..eb3b34abd1 100644 --- a/SYCL/ESIMD/fp_call_recursive.cpp +++ b/SYCL/ESIMD/fp_call_recursive.cpp @@ -43,10 +43,10 @@ int main(int argc, char **argv) { unsigned in2 = 21; unsigned in3 = 3; - { + try { buffer buf(output, range<1>(1)); - auto e = q.submit([&](handler &cgh) { + q.submit([&](handler &cgh) { auto acc = buf.get_access(cgh); cgh.parallel_for(sycl::range<1>{1}, @@ -59,7 +59,9 @@ int main(int argc, char **argv) { scalar_store(acc, 0, res); }); }); - e.wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return e.get_cl_code(); } int etalon = in1; diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp index 4b37eb3915..ff8f7df991 100644 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -38,11 +38,11 @@ bool test(queue q, bool flag) { int in1 = 233; int in2 = 1; - { + try { buffer o_buf(output, range<1>(1)); buffer y_buf(Y.data(), Y.size()); - auto e = q.submit([&](handler &cgh) { + q.submit([&](handler &cgh) { auto o_acc = o_buf.get_access(cgh); auto y_acc = y_buf.get_access(cgh); @@ -62,7 +62,9 @@ bool test(queue q, bool flag) { scalar_store(o_acc, 0, res); }); }); - e.wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return e.get_cl_code(); } int etalon = in1 + (flag ? 3 : 1) + in2 + 2; diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp index 04af9db691..8c9152557f 100644 --- a/SYCL/ESIMD/fp_in_select.cpp +++ b/SYCL/ESIMD/fp_in_select.cpp @@ -33,10 +33,10 @@ bool test(queue q, bool flag) { int in1 = 233; int in2 = 100; - { + try { buffer buf(output, range<1>(1)); - auto e = q.submit([&](handler &cgh) { + q.submit([&](handler &cgh) { auto acc = buf.get_access(cgh); cgh.parallel_for(sycl::range<1>{1}, @@ -49,7 +49,9 @@ bool test(queue q, bool flag) { scalar_store(acc, 0, res); }); }); - e.wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return false; } int etalon = flag ? in1 + in2 : in1 - in2; From 0c9aad6451dc3f1ac97049beef65be0074043553 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 19:15:53 +0300 Subject: [PATCH 19/32] minor fixes --- SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp | 4 +++- SYCL/ESIMD/fp_call_from_func.cpp | 3 ++- SYCL/ESIMD/fp_in_phi.cpp | 1 + SYCL/ESIMD/fp_in_select.cpp | 1 + 4 files changed, 7 insertions(+), 2 deletions(-) diff --git a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp index aac404fcd8..2afcdc19f1 100644 --- a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp +++ b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp @@ -56,7 +56,7 @@ int main(void) { memset(C, 0, SIZE * sizeof(c_data_t)); try { - q.submit([&](handler &cgh) { + auto qq = q.submit([&](handler &cgh) { cgh.parallel_for( sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::INTEL::gpu; @@ -76,6 +76,8 @@ int main(void) { 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); diff --git a/SYCL/ESIMD/fp_call_from_func.cpp b/SYCL/ESIMD/fp_call_from_func.cpp index 8890553f8d..483f66b17d 100644 --- a/SYCL/ESIMD/fp_call_from_func.cpp +++ b/SYCL/ESIMD/fp_call_from_func.cpp @@ -5,8 +5,9 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// +// TODO: Support Windows Level Zero - this test timeouts // REQUIRES: gpu -// UNSUPPORTED: cuda +// UNSUPPORTED: cuda || (windows && level_zero) // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp index ff8f7df991..cbf4678468 100644 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// XFAIL: windows // UNSUPPORTED: cuda // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp index 8c9152557f..e53b11d9cd 100644 --- a/SYCL/ESIMD/fp_in_select.cpp +++ b/SYCL/ESIMD/fp_in_select.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// XFAIL: windows // UNSUPPORTED: cuda // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out From 43b80bdab449913aad8ccda688c20c8e429a2f4c Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 19:25:35 +0300 Subject: [PATCH 20/32] cosmetic change --- SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp index 2afcdc19f1..a3cd76b192 100644 --- a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp +++ b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp @@ -5,6 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// +// // The test checks that ESIMD kernels support use of function pointers from main // function with different total arguments size and retval size. // Cases: From 32be755e47672a4482eab57461ddc8af522e68a4 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Wed, 10 Feb 2021 16:18:24 +0300 Subject: [PATCH 21/32] handle general exception --- SYCL/ESIMD/fp_in_phi.cpp | 4 +++- SYCL/ESIMD/fp_in_select.cpp | 4 +++- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp index cbf4678468..b2cc890aa4 100644 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -6,7 +6,6 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: windows // UNSUPPORTED: cuda // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out @@ -66,6 +65,9 @@ bool test(queue q, bool flag) { } catch (cl::sycl::exception const &e) { std::cout << "SYCL exception caught: " << e.what() << std::endl; return e.get_cl_code(); + } catch (std::exception const &e) { + std::cout << "General exception caught: " << e.what() << std::endl; + return 2; } int etalon = in1 + (flag ? 3 : 1) + in2 + 2; diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp index e53b11d9cd..34b4b8696d 100644 --- a/SYCL/ESIMD/fp_in_select.cpp +++ b/SYCL/ESIMD/fp_in_select.cpp @@ -6,7 +6,6 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: windows // UNSUPPORTED: cuda // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out @@ -53,6 +52,9 @@ bool test(queue q, bool flag) { } catch (cl::sycl::exception const &e) { std::cout << "SYCL exception caught: " << e.what() << std::endl; return false; + } catch (std::exception const &e) { + std::cout << "General exception caught: " << e.what() << std::endl; + return 2; } int etalon = flag ? in1 + in2 : in1 - in2; From e978829352e46dbe5186b9d104028a260e127984 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Wed, 10 Feb 2021 17:03:48 +0300 Subject: [PATCH 22/32] revert --- SYCL/ESIMD/fp_in_phi.cpp | 3 --- SYCL/ESIMD/fp_in_select.cpp | 3 --- 2 files changed, 6 deletions(-) diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp index b2cc890aa4..ff8f7df991 100644 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -65,9 +65,6 @@ bool test(queue q, bool flag) { } catch (cl::sycl::exception const &e) { std::cout << "SYCL exception caught: " << e.what() << std::endl; return e.get_cl_code(); - } catch (std::exception const &e) { - std::cout << "General exception caught: " << e.what() << std::endl; - return 2; } int etalon = in1 + (flag ? 3 : 1) + in2 + 2; diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp index 34b4b8696d..8c9152557f 100644 --- a/SYCL/ESIMD/fp_in_select.cpp +++ b/SYCL/ESIMD/fp_in_select.cpp @@ -52,9 +52,6 @@ bool test(queue q, bool flag) { } catch (cl::sycl::exception const &e) { std::cout << "SYCL exception caught: " << e.what() << std::endl; return false; - } catch (std::exception const &e) { - std::cout << "General exception caught: " << e.what() << std::endl; - return 2; } int etalon = flag ? in1 + in2 : in1 - in2; From 96d312155f5789952b794f98927bf8abf12fd366 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Wed, 10 Feb 2021 22:19:03 +0300 Subject: [PATCH 23/32] cosmetic fixes --- .../ESIMD/fp_args_size/fp_args_char_int_size_192.cpp | 8 +++++--- .../ESIMD/fp_args_size/fp_args_char_int_size_256.cpp | 8 +++++--- .../ESIMD/fp_args_size/fp_args_char_int_size_512.cpp | 8 +++++--- SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp | 8 +++++--- SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp | 8 +++++--- SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp | 8 +++++--- SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp | 8 +++++--- SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp | 8 +++++--- SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp | 8 +++++--- SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp | 8 +++++--- SYCL/ESIMD/fp_in_phi.cpp | 12 ++++++++++-- SYCL/ESIMD/fp_in_select.cpp | 12 ++++++++++-- 12 files changed, 70 insertions(+), 34 deletions(-) diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp index 144c0d62c4..39f38df192 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 192; -typedef char a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp index 03f853fbe8..cfe71a18bb 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 256; -typedef char a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp index a0f4787b4b..2fee0d351c 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 512; -typedef char a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp index 40ee81393b..b33e13e475 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 64; -typedef char a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp index f8eb2c0c4c..02f4fecb28 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 96; -typedef char a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int8_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp index 1a2f6c2452..73c7c2f83b 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 192; -typedef int a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp index 63589134c0..c098139a48 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 256; -typedef int a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp index 69bf678b65..67c4e63c35 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 512; -typedef int a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp index eb24a669b0..6ed389d0fd 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 64; -typedef int a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp index 3247026439..164025654c 100644 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp +++ b/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp @@ -10,11 +10,13 @@ // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +#include + constexpr unsigned VL = 16; constexpr unsigned SIZE = 96; -typedef int a_data_t; -typedef int b_data_t; -typedef int c_data_t; +using a_data_t = int32_t; +using b_data_t = int32_t; +using c_data_t = int32_t; #include "Inputs/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp index ff8f7df991..0c867422aa 100644 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -84,8 +84,16 @@ int main(int argc, char **argv) { std::cout << "Running on " << dev.get_info() << "\n"; bool passed = true; - passed &= test(q, true); - passed &= test(q, false); + try { + passed &= test(q, true); + passed &= test(q, false); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return e.get_cl_code(); + } catch (std::exception const &e) { + std::cout << "General exception caught: " << e.what() << std::endl; + return -1; + } return passed ? 0 : 1; } diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp index 8c9152557f..ea3d2de8b7 100644 --- a/SYCL/ESIMD/fp_in_select.cpp +++ b/SYCL/ESIMD/fp_in_select.cpp @@ -71,8 +71,16 @@ int main(int argc, char **argv) { std::cout << "Running on " << dev.get_info() << "\n"; bool passed = true; - passed &= test(q, true); - passed &= test(q, false); + try { + passed &= test(q, true); + passed &= test(q, false); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << std::endl; + return e.get_cl_code(); + } catch (std::exception const &e) { + std::cout << "General exception caught: " << e.what() << std::endl; + return -1; + } return passed ? 0 : 1; } From 117ae42b727db8c58529bfe2491a03d606a1d01d Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Thu, 11 Feb 2021 14:27:12 +0300 Subject: [PATCH 24/32] revert --- SYCL/ESIMD/fp_call_from_func.cpp | 7 ++++--- SYCL/ESIMD/fp_in_phi.cpp | 12 ++---------- SYCL/ESIMD/fp_in_select.cpp | 14 +++----------- 3 files changed, 9 insertions(+), 24 deletions(-) diff --git a/SYCL/ESIMD/fp_call_from_func.cpp b/SYCL/ESIMD/fp_call_from_func.cpp index 483f66b17d..64cdf8d8a1 100644 --- a/SYCL/ESIMD/fp_call_from_func.cpp +++ b/SYCL/ESIMD/fp_call_from_func.cpp @@ -5,14 +5,15 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO: Support Windows Level Zero - this test timeouts // REQUIRES: gpu -// UNSUPPORTED: cuda || (windows && level_zero) +// Test timeouts on Windows Level Zero +// UNSUPPORTED: windows && level_zero // RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda // // The test checks that ESIMD kernels support use of function pointers from -// functions. +// within other functions. #include "esimd_test_utils.hpp" diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp index 0c867422aa..ff8f7df991 100644 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -84,16 +84,8 @@ int main(int argc, char **argv) { std::cout << "Running on " << dev.get_info() << "\n"; bool passed = true; - try { - passed &= test(q, true); - passed &= test(q, false); - } catch (cl::sycl::exception const &e) { - std::cout << "SYCL exception caught: " << e.what() << std::endl; - return e.get_cl_code(); - } catch (std::exception const &e) { - std::cout << "General exception caught: " << e.what() << std::endl; - return -1; - } + passed &= test(q, true); + passed &= test(q, false); return passed ? 0 : 1; } diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp index ea3d2de8b7..c63a2ff31f 100644 --- a/SYCL/ESIMD/fp_in_select.cpp +++ b/SYCL/ESIMD/fp_in_select.cpp @@ -11,7 +11,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // // The test checks that ESIMD kernels correctly handle function pointers as -// arguments of select function. +// arguments of LLVM's select function. #include "esimd_test_utils.hpp" @@ -71,16 +71,8 @@ int main(int argc, char **argv) { std::cout << "Running on " << dev.get_info() << "\n"; bool passed = true; - try { - passed &= test(q, true); - passed &= test(q, false); - } catch (cl::sycl::exception const &e) { - std::cout << "SYCL exception caught: " << e.what() << std::endl; - return e.get_cl_code(); - } catch (std::exception const &e) { - std::cout << "General exception caught: " << e.what() << std::endl; - return -1; - } + passed &= test(q, true); + passed &= test(q, false); return passed ? 0 : 1; } From ded60bdeac9e28410545473805ce8a3398e85a18 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 12 Feb 2021 14:21:41 +0300 Subject: [PATCH 25/32] improve fail message --- SYCL/ESIMD/fp_call_from_func.cpp | 2 +- SYCL/ESIMD/fp_call_recursive.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/SYCL/ESIMD/fp_call_from_func.cpp b/SYCL/ESIMD/fp_call_from_func.cpp index 64cdf8d8a1..c6b4c800c1 100644 --- a/SYCL/ESIMD/fp_call_from_func.cpp +++ b/SYCL/ESIMD/fp_call_from_func.cpp @@ -63,7 +63,7 @@ int main(int argc, char **argv) { } if (result != (in1 + in2)) { - std::cout << "Failed" << std::endl; + std::cout << "Failed with result: " << result << std::endl; return 1; } diff --git a/SYCL/ESIMD/fp_call_recursive.cpp b/SYCL/ESIMD/fp_call_recursive.cpp index eb3b34abd1..3d4cc36f56 100644 --- a/SYCL/ESIMD/fp_call_recursive.cpp +++ b/SYCL/ESIMD/fp_call_recursive.cpp @@ -76,7 +76,7 @@ int main(int argc, char **argv) { } if (result != etalon) { - std::cout << "Failed: " << result << std::endl; + std::cout << "Failed with result: " << result << std::endl; return 1; } From b4ce67d4077114c10eac25745969171cc043ae4a Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 12 Feb 2021 20:27:09 +0300 Subject: [PATCH 26/32] noinline function tests initial commit --- .../Inputs/noinline_func_args_size_common.hpp | 104 ++++++++++++++++++ .../noinline_func_args_char_int_size_192.cpp | 22 ++++ .../noinline_func_args_char_int_size_256.cpp | 22 ++++ .../noinline_func_args_char_int_size_512.cpp | 22 ++++ .../noinline_func_args_char_int_size_64.cpp | 22 ++++ .../noinline_func_args_char_int_size_96.cpp | 22 ++++ .../noinline_func_args_int_size_192.cpp | 22 ++++ .../noinline_func_args_int_size_256.cpp | 22 ++++ .../noinline_func_args_int_size_512.cpp | 22 ++++ .../noinline_func_args_int_size_64.cpp | 22 ++++ .../noinline_func_args_int_size_96.cpp | 22 ++++ SYCL/ESIMD/noinline_func_call_from_func.cpp | 74 +++++++++++++ SYCL/ESIMD/noinline_func_call_recursive.cpp | 82 ++++++++++++++ 13 files changed, 480 insertions(+) create mode 100644 SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp create mode 100644 SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp create mode 100644 SYCL/ESIMD/noinline_func_call_from_func.cpp create mode 100644 SYCL/ESIMD/noinline_func_call_recursive.cpp diff --git a/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp b/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp new file mode 100644 index 0000000000..55fa3f6054 --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp @@ -0,0 +1,104 @@ +//==--- noinline_func_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); + } + + simd vc = add(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_func_args_size/noinline_func_args_char_int_size_192.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp new file mode 100644 index 0000000000..589965729a --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp @@ -0,0 +1,22 @@ +//=- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp new file mode 100644 index 0000000000..321e84603e --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp @@ -0,0 +1,22 @@ +//=- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp new file mode 100644 index 0000000000..e3b5bd54c4 --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp @@ -0,0 +1,22 @@ +//=- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp new file mode 100644 index 0000000000..3789ea6d65 --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp @@ -0,0 +1,22 @@ +//==- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp new file mode 100644 index 0000000000..6b97d3170b --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp @@ -0,0 +1,22 @@ +//==- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp new file mode 100644 index 0000000000..b84d166812 --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp @@ -0,0 +1,22 @@ +//=- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp new file mode 100644 index 0000000000..3e20e484a2 --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp @@ -0,0 +1,22 @@ +//=- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp new file mode 100644 index 0000000000..e7c570e455 --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp @@ -0,0 +1,22 @@ +//=- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp new file mode 100644 index 0000000000..dc2fb61360 --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp @@ -0,0 +1,22 @@ +//==- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp new file mode 100644 index 0000000000..d76b3aba25 --- /dev/null +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp @@ -0,0 +1,22 @@ +//==- noinline_func_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 %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_func_args_size_common.hpp" diff --git a/SYCL/ESIMD/noinline_func_call_from_func.cpp b/SYCL/ESIMD/noinline_func_call_from_func.cpp new file mode 100644 index 0000000000..a52b84026f --- /dev/null +++ b/SYCL/ESIMD/noinline_func_call_from_func.cpp @@ -0,0 +1,74 @@ +//==---- noinline_func_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 +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out +// +// 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) { + try { + 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; + } + } catch (std::exception const &e) { + std::cout << "General exception caught: " << e.what() << std::endl; + return -1; + } + + std::cout << "passed" << std::endl; + return 0; +} diff --git a/SYCL/ESIMD/noinline_func_call_recursive.cpp b/SYCL/ESIMD/noinline_func_call_recursive.cpp new file mode 100644 index 0000000000..395de579f4 --- /dev/null +++ b/SYCL/ESIMD/noinline_func_call_recursive.cpp @@ -0,0 +1,82 @@ +//==---- noinline_func_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 +// UNSUPPORTED: cuda +// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out +// +// 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; +} From 9933741ecfd00c1f3f36bd92e1f685b3fdd382e1 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 12 Feb 2021 20:46:02 +0300 Subject: [PATCH 27/32] clang-format patch --- .../Inputs/noinline_func_args_size_common.hpp | 5 +- SYCL/ESIMD/noinline_func_call_from_func.cpp | 61 +++++++++---------- 2 files changed, 30 insertions(+), 36 deletions(-) diff --git a/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp b/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp index 55fa3f6054..f2c66458c8 100644 --- a/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp +++ b/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp @@ -6,9 +6,8 @@ // //===----------------------------------------------------------------------===// // -// The test checks that ESIMD kernels support call of noinline function from main -// function with different total arguments size and retval size. -// Cases: +// 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) diff --git a/SYCL/ESIMD/noinline_func_call_from_func.cpp b/SYCL/ESIMD/noinline_func_call_from_func.cpp index a52b84026f..366a6d8223 100644 --- a/SYCL/ESIMD/noinline_func_call_from_func.cpp +++ b/SYCL/ESIMD/noinline_func_call_from_func.cpp @@ -33,40 +33,35 @@ template ESIMD_NOINLINE void test(AccTy acc, int A, int B) { } 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 { - 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; - } - } catch (std::exception const &e) { - std::cout << "General exception caught: " << e.what() << std::endl; - return -1; + 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; From 19079ad51326c1ea5dd5bacbb06e552db5a75ab4 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 12 Feb 2021 20:56:08 +0300 Subject: [PATCH 28/32] clang-format patch 2 --- SYCL/ESIMD/noinline_func_call_from_func.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/SYCL/ESIMD/noinline_func_call_from_func.cpp b/SYCL/ESIMD/noinline_func_call_from_func.cpp index 366a6d8223..5cb770f112 100644 --- a/SYCL/ESIMD/noinline_func_call_from_func.cpp +++ b/SYCL/ESIMD/noinline_func_call_from_func.cpp @@ -48,12 +48,12 @@ int main(int argc, char **argv) { buffer buf(output, range<1>(1)); q.submit([&](handler &cgh) { - auto acc = buf.get_access(cgh); + auto acc = buf.get_access(cgh); - cgh.parallel_for( - sycl::range<1>{1}, - [=](id<1> i) SYCL_ESIMD_KERNEL { test(acc, in1, in2); }); - }); + 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(); From 01b57035ac0429d56f5a721e0048911192457f31 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 12 Feb 2021 20:59:11 +0300 Subject: [PATCH 29/32] remove erroneously added tests --- .../Inputs/fp_args_size_common.hpp | 106 ------------------ .../fp_args_char_int_size_192.cpp | 22 ---- .../fp_args_char_int_size_256.cpp | 22 ---- .../fp_args_char_int_size_512.cpp | 22 ---- .../fp_args_size/fp_args_char_int_size_64.cpp | 22 ---- .../fp_args_size/fp_args_char_int_size_96.cpp | 22 ---- .../fp_args_size/fp_args_int_size_192.cpp | 22 ---- .../fp_args_size/fp_args_int_size_256.cpp | 22 ---- .../fp_args_size/fp_args_int_size_512.cpp | 22 ---- .../fp_args_size/fp_args_int_size_64.cpp | 22 ---- .../fp_args_size/fp_args_int_size_96.cpp | 22 ---- SYCL/ESIMD/fp_call_from_func.cpp | 71 ------------ SYCL/ESIMD/fp_call_recursive.cpp | 84 -------------- SYCL/ESIMD/fp_in_phi.cpp | 91 --------------- SYCL/ESIMD/fp_in_select.cpp | 78 ------------- 15 files changed, 650 deletions(-) delete mode 100644 SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp delete mode 100644 SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp delete mode 100644 SYCL/ESIMD/fp_call_from_func.cpp delete mode 100644 SYCL/ESIMD/fp_call_recursive.cpp delete mode 100644 SYCL/ESIMD/fp_in_phi.cpp delete mode 100644 SYCL/ESIMD/fp_in_select.cpp diff --git a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp b/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp deleted file mode 100644 index a3cd76b192..0000000000 --- a/SYCL/ESIMD/fp_args_size/Inputs/fp_args_size_common.hpp +++ /dev/null @@ -1,106 +0,0 @@ -//==------- fp_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 use of function pointers 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 foo = &add, simd, - simd>; - auto vc = foo(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/fp_args_size/fp_args_char_int_size_192.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp deleted file mode 100644 index 39f38df192..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_192.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp deleted file mode 100644 index cfe71a18bb..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp deleted file mode 100644 index 2fee0d351c..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_512.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp deleted file mode 100644 index b33e13e475..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp b/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp deleted file mode 100644 index 02f4fecb28..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_char_int_size_96.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp deleted file mode 100644 index 73c7c2f83b..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_192.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp deleted file mode 100644 index c098139a48..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp deleted file mode 100644 index 67c4e63c35..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_512.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp deleted file mode 100644 index 6ed389d0fd..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp b/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp deleted file mode 100644 index 164025654c..0000000000 --- a/SYCL/ESIMD/fp_args_size/fp_args_int_size_96.cpp +++ /dev/null @@ -1,22 +0,0 @@ -//==------- fp_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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl -I%S/.. %s -o %t.out -// RUN: %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/fp_args_size_common.hpp" diff --git a/SYCL/ESIMD/fp_call_from_func.cpp b/SYCL/ESIMD/fp_call_from_func.cpp deleted file mode 100644 index c6b4c800c1..0000000000 --- a/SYCL/ESIMD/fp_call_from_func.cpp +++ /dev/null @@ -1,71 +0,0 @@ -//==--------------- fp_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 timeouts on Windows Level Zero -// UNSUPPORTED: windows && level_zero -// RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// UNSUPPORTED: cuda -// -// The test checks that ESIMD kernels support use of function pointers 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 foo = &add; - auto res = foo(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; - } - - return 0; -} diff --git a/SYCL/ESIMD/fp_call_recursive.cpp b/SYCL/ESIMD/fp_call_recursive.cpp deleted file mode 100644 index 3d4cc36f56..0000000000 --- a/SYCL/ESIMD/fp_call_recursive.cpp +++ /dev/null @@ -1,84 +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 -// -//===----------------------------------------------------------------------===// -// REQUIRES: gpu -// UNSUPPORTED: cuda -// RUN: %clangxx-esimd -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// -// The test checks that ESIMD kernels support use of function pointers -// recursively. - -#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; - - 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::INTEL::gpu; - - auto foo = &add; - auto res = foo(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 with result: " << result << std::endl; - return 1; - } - - return 0; -} diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp deleted file mode 100644 index ff8f7df991..0000000000 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ /dev/null @@ -1,91 +0,0 @@ -//==--------------- fp_in_phi.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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// -// The test checks that ESIMD kernels correctly handle function pointers as -// arguments of LLVM's PHI function. - -#include "esimd_test_utils.hpp" - -#include -#include - -#include -#include - -class KernelID; - -ESIMD_NOINLINE int f1(int x) { return x + 1; } - -ESIMD_NOINLINE int f2(int x) { return x + 2; } - -ESIMD_NOINLINE int f3(int x) { return x + 3; } - -bool test(queue q, bool flag) { - int result = 0; - int *output = &result; - - std::vector Y = {0, 1}; - - int in1 = 233; - int in2 = 1; - - try { - buffer o_buf(output, range<1>(1)); - buffer y_buf(Y.data(), Y.size()); - - q.submit([&](handler &cgh) { - auto o_acc = o_buf.get_access(cgh); - auto y_acc = y_buf.get_access(cgh); - - cgh.parallel_for(sycl::range<1>{1}, - [=](id<1> i) SYCL_ESIMD_KERNEL { - using namespace sycl::INTEL::gpu; - using f = int (*)(int); - - f a[] = {f1, f2}; - if (flag) { - a[0] = f3; - scalar_store(y_acc, 0, 2); - } - - auto res = a[0](in1) + a[1](in2); - - scalar_store(o_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 + (flag ? 3 : 1) + in2 + 2; - - if (result != etalon) { - std::cout << "Failed with result: " << result << std::endl; - return false; - } - - return true; -} - -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"; - - bool passed = true; - passed &= test(q, true); - passed &= test(q, false); - - return passed ? 0 : 1; -} diff --git a/SYCL/ESIMD/fp_in_select.cpp b/SYCL/ESIMD/fp_in_select.cpp deleted file mode 100644 index c63a2ff31f..0000000000 --- a/SYCL/ESIMD/fp_in_select.cpp +++ /dev/null @@ -1,78 +0,0 @@ -//==--------------- fp_in_select.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 -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// -// The test checks that ESIMD kernels correctly handle function pointers as -// arguments of LLVM's select function. - -#include "esimd_test_utils.hpp" - -#include -#include - -#include - -class KernelID; - -ESIMD_NOINLINE int add(int a, int b) { return a + b; } - -ESIMD_NOINLINE int sub(int a, int b) { return a - b; } - -bool test(queue q, bool flag) { - int result = 0; - int *output = &result; - - int in1 = 233; - int in2 = 100; - - 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 foo = flag ? &add : ⊂ - auto res = foo(in1, in2); - - scalar_store(acc, 0, res); - }); - }); - } catch (cl::sycl::exception const &e) { - std::cout << "SYCL exception caught: " << e.what() << std::endl; - return false; - } - - int etalon = flag ? in1 + in2 : in1 - in2; - - if (result != etalon) { - std::cout << "Failed with result: " << result << std::endl; - return false; - } - - return true; -} - -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"; - - bool passed = true; - passed &= test(q, true); - passed &= test(q, false); - - return passed ? 0 : 1; -} From 54ff602d241dd9be8c80c9db14514828dc5ccc94 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 12 Feb 2021 23:11:21 +0300 Subject: [PATCH 30/32] fix args size tests --- .../Inputs/noinline_func_args_size_common.hpp | 3 ++- .../noinline_func_args_char_int_size_192.cpp | 2 +- .../noinline_func_args_char_int_size_256.cpp | 2 +- .../noinline_func_args_char_int_size_512.cpp | 2 +- .../noinline_func_args_char_int_size_64.cpp | 2 +- .../noinline_func_args_char_int_size_96.cpp | 2 +- .../noinline_func_args_int_size_192.cpp | 2 +- .../noinline_func_args_int_size_256.cpp | 2 +- .../noinline_func_args_int_size_512.cpp | 2 +- .../noinline_func_args_size/noinline_func_args_int_size_64.cpp | 2 +- .../noinline_func_args_size/noinline_func_args_int_size_96.cpp | 2 +- 11 files changed, 12 insertions(+), 11 deletions(-) diff --git a/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp b/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp index f2c66458c8..a2ca8b4f4c 100644 --- a/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp +++ b/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp @@ -68,7 +68,8 @@ int main(void) { vb.select(j * VL) = block_load(B + j * VL); } - simd vc = add(va, vb); + auto vc = add, simd, + simd>(va, vb); for (int j = 0; j < ROWS; j++) block_store(C + j * VL, vc.select(j * VL)); diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp index 589965729a..d0601b49c1 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp index 321e84603e..b98d80ec68 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp index e3b5bd54c4..56cdb789de 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp index 3789ea6d65..9e0dafa56a 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp index 6b97d3170b..27037e700b 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp index b84d166812..86e5181a27 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp index 3e20e484a2..cf0c8a6855 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp index e7c570e455..77cf2cee8d 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp index dc2fb61360..eb7eb881cd 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp index d76b3aba25..80487596ed 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp +++ b/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp @@ -7,7 +7,7 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda -// RUN: %clangxx-esimd -fsycl %s -o %t.out +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out #include From 4c5b252e2f0604cf8e9e317017e87e19dfcee878 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Mon, 15 Feb 2021 17:22:44 +0300 Subject: [PATCH 31/32] rename and set UNSUPPORTED status --- .../Inputs/noinline_args_size_common.hpp} | 2 +- .../noinline_args_char_int_size_192.cpp} | 2 +- .../noinline_args_char_int_size_256.cpp} | 2 +- .../noinline_args_char_int_size_512.cpp} | 2 +- .../noinline_args_char_int_size_64.cpp} | 2 +- .../noinline_args_char_int_size_96.cpp} | 2 +- .../noinline_args_int_size_192.cpp} | 2 +- .../noinline_args_int_size_256.cpp} | 2 +- .../noinline_args_int_size_512.cpp} | 2 +- .../noinline_args_int_size_64.cpp} | 2 +- .../noinline_args_int_size_96.cpp} | 2 +- ..._func_call_from_func.cpp => noinline_call_from_func.cpp} | 6 ++++-- ..._func_call_recursive.cpp => noinline_call_recursive.cpp} | 4 ++-- 13 files changed, 17 insertions(+), 15 deletions(-) rename SYCL/ESIMD/{noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp => noinline_args_size/Inputs/noinline_args_size_common.hpp} (97%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_char_int_size_192.cpp => noinline_args_size/noinline_args_char_int_size_192.cpp} (89%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_char_int_size_256.cpp => noinline_args_size/noinline_args_char_int_size_256.cpp} (89%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_char_int_size_512.cpp => noinline_args_size/noinline_args_char_int_size_512.cpp} (89%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_char_int_size_64.cpp => noinline_args_size/noinline_args_char_int_size_64.cpp} (89%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_char_int_size_96.cpp => noinline_args_size/noinline_args_char_int_size_96.cpp} (89%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_int_size_192.cpp => noinline_args_size/noinline_args_int_size_192.cpp} (89%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_int_size_256.cpp => noinline_args_size/noinline_args_int_size_256.cpp} (89%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_int_size_512.cpp => noinline_args_size/noinline_args_int_size_512.cpp} (89%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_int_size_64.cpp => noinline_args_size/noinline_args_int_size_64.cpp} (89%) rename SYCL/ESIMD/{noinline_func_args_size/noinline_func_args_int_size_96.cpp => noinline_args_size/noinline_args_int_size_96.cpp} (89%) rename SYCL/ESIMD/{noinline_func_call_from_func.cpp => noinline_call_from_func.cpp} (91%) rename SYCL/ESIMD/{noinline_func_call_recursive.cpp => noinline_call_recursive.cpp} (96%) diff --git a/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp b/SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp similarity index 97% rename from SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp rename to SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp index a2ca8b4f4c..8a908497c0 100644 --- a/SYCL/ESIMD/noinline_func_args_size/Inputs/noinline_func_args_size_common.hpp +++ b/SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp @@ -1,4 +1,4 @@ -//==--- noinline_func_args_size_common.hpp - DPC++ ESIMD on-device test ---==// +//===------ 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_192.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_192.cpp index d0601b49c1..a741015d4d 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_192.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_192.cpp @@ -1,4 +1,4 @@ -//=- noinline_func_args_char_int_size_192.cpp - DPC++ ESIMD on-device test -=// +//===-- 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_256.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_256.cpp index b98d80ec68..28a8bd73d6 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_256.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_256.cpp @@ -1,4 +1,4 @@ -//=- noinline_func_args_char_int_size_256.cpp - DPC++ ESIMD on-device test -=// +//===-- 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_512.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_512.cpp index 56cdb789de..57bbe80330 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_512.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_512.cpp @@ -1,4 +1,4 @@ -//=- noinline_func_args_char_int_size_512.cpp - DPC++ ESIMD on-device test -=// +//===-- 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_64.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_64.cpp index 9e0dafa56a..afbcedf7c1 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_64.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_64.cpp @@ -1,4 +1,4 @@ -//==- noinline_func_args_char_int_size_64.cpp - DPC++ ESIMD on-device test ==// +//===--- 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_96.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_96.cpp index 27037e700b..36b042e8e3 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_char_int_size_96.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_96.cpp @@ -1,4 +1,4 @@ -//==- noinline_func_args_char_int_size_96.cpp - DPC++ ESIMD on-device test ==// +//===--- 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp index 86e5181a27..80098cdd69 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_192.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp @@ -1,4 +1,4 @@ -//=- noinline_func_args_char_int_size_192.cpp - DPC++ ESIMD on-device test -=// +//===-- 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp index cf0c8a6855..df65069baa 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_256.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp @@ -1,4 +1,4 @@ -//=- noinline_func_args_char_int_size_256.cpp - DPC++ ESIMD on-device test -=// +//===-- 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp index 77cf2cee8d..24b68a2dc4 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_512.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp @@ -1,4 +1,4 @@ -//=- noinline_func_args_char_int_size_512.cpp - DPC++ ESIMD on-device test -=// +//===-- 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp index eb7eb881cd..cfc771b0cd 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_64.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp @@ -1,4 +1,4 @@ -//==- noinline_func_args_char_int_size_64.cpp - DPC++ ESIMD on-device test ==// +//===--- 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. diff --git a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp similarity index 89% rename from SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp rename to SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp index 80487596ed..20a34586db 100644 --- a/SYCL/ESIMD/noinline_func_args_size/noinline_func_args_int_size_96.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp @@ -1,4 +1,4 @@ -//==- noinline_func_args_char_int_size_96.cpp - DPC++ ESIMD on-device test ==// +//===--- 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. diff --git a/SYCL/ESIMD/noinline_func_call_from_func.cpp b/SYCL/ESIMD/noinline_call_from_func.cpp similarity index 91% rename from SYCL/ESIMD/noinline_func_call_from_func.cpp rename to SYCL/ESIMD/noinline_call_from_func.cpp index 5cb770f112..6484bfee6d 100644 --- a/SYCL/ESIMD/noinline_func_call_from_func.cpp +++ b/SYCL/ESIMD/noinline_call_from_func.cpp @@ -1,4 +1,4 @@ -//==---- noinline_func_call_from_func.cpp - DPC++ ESIMD on-device test ----==// +//===------ 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. @@ -6,9 +6,11 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// UNSUPPORTED: cuda +// 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. diff --git a/SYCL/ESIMD/noinline_func_call_recursive.cpp b/SYCL/ESIMD/noinline_call_recursive.cpp similarity index 96% rename from SYCL/ESIMD/noinline_func_call_recursive.cpp rename to SYCL/ESIMD/noinline_call_recursive.cpp index 395de579f4..60d879ee88 100644 --- a/SYCL/ESIMD/noinline_func_call_recursive.cpp +++ b/SYCL/ESIMD/noinline_call_recursive.cpp @@ -1,4 +1,4 @@ -//==---- noinline_func_call_recursive.cpp - DPC++ ESIMD on-device test ----==// +//===------ 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. @@ -6,9 +6,9 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// UNSUPPORTED: cuda // 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. From 190a4b67f7e290fb276181a8edcf57c040586f51 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Mon, 15 Feb 2021 18:46:55 +0300 Subject: [PATCH 32/32] typo fix --- .../noinline_args_size/noinline_args_char_int_size_192.cpp | 2 +- .../noinline_args_size/noinline_args_char_int_size_256.cpp | 2 +- .../noinline_args_size/noinline_args_char_int_size_512.cpp | 2 +- .../ESIMD/noinline_args_size/noinline_args_char_int_size_64.cpp | 2 +- .../ESIMD/noinline_args_size/noinline_args_char_int_size_96.cpp | 2 +- SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp | 2 +- SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp | 2 +- SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp | 2 +- SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp | 2 +- SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp | 2 +- 10 files changed, 10 insertions(+), 10 deletions(-) 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 index a741015d4d..e152604bb0 100644 --- 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 @@ -19,4 +19,4 @@ using a_data_t = int8_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#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 index 28a8bd73d6..ff06f8d4fb 100644 --- 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 @@ -19,4 +19,4 @@ using a_data_t = int8_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#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 index 57bbe80330..1894de3c5a 100644 --- 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 @@ -19,4 +19,4 @@ using a_data_t = int8_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#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 index afbcedf7c1..5785b3e48b 100644 --- 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 @@ -19,4 +19,4 @@ using a_data_t = int8_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#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 index 36b042e8e3..cbd084e192 100644 --- 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 @@ -19,4 +19,4 @@ using a_data_t = int8_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#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 index 80098cdd69..2cb5ca0361 100644 --- a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp @@ -19,4 +19,4 @@ using a_data_t = int32_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#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 index df65069baa..ade1d5d91b 100644 --- a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp @@ -19,4 +19,4 @@ using a_data_t = int32_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#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 index 24b68a2dc4..22e3b14fe8 100644 --- a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp @@ -19,4 +19,4 @@ using a_data_t = int32_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#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 index cfc771b0cd..69fcf80022 100644 --- a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp @@ -19,4 +19,4 @@ using a_data_t = int32_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#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 index 20a34586db..7c9c198cea 100644 --- a/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp +++ b/SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp @@ -19,4 +19,4 @@ using a_data_t = int32_t; using b_data_t = int32_t; using c_data_t = int32_t; -#include "Inputs/noinline_func_args_size_common.hpp" +#include "Inputs/noinline_args_size_common.hpp"