From 5d55ee256063201d5703249a6e072e5f22069445 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 4 Dec 2020 01:03:44 +0300 Subject: [PATCH 01/39] 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/39] [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/39] [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/39] [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/39] [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/39] [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/39] [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/39] 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/39] 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/39] 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/39] 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/39] 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/39] [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/39] 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 544ac3f332d0780940c7e101131c51ed328437f0 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Thu, 4 Feb 2021 22:32:54 +0300 Subject: [PATCH 15/39] [SYCL][ESIMD] spec const tests for all basic types --- .../spec_const/Inputs/spec_const_common.hpp | 68 +++++++++++++++++++ SYCL/ESIMD/spec_const/spec_const_bool.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_char.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_double.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_float.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_int.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_long.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_short.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_uint.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 31 +++++++++ SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 31 +++++++++ 12 files changed, 409 insertions(+) create mode 100644 SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_bool.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_char.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_double.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_float.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_int.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_long.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_short.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_uchar.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_uint.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_ulong.cpp create mode 100644 SYCL/ESIMD/spec_const/spec_const_ushort.cpp diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp new file mode 100644 index 0000000000..a2c0e38776 --- /dev/null +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -0,0 +1,68 @@ +//==--------------- spec_const_common.h - 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 specialization constants for all +// basic types, particularly a specialization constant can be redifined and +// correct new value is used after redefinition. + +#include +#include + +using namespace sycl; + +class ConstID; +class TestKernel; + +int main(int argc, char **argv) { + queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()); + auto dev = q.get_device(); + auto ctx = q.get_context(); + + const int n_times = 2; + std::vector output(n_times); + std::vector etalon = {DEF_VAL, REDEF_VAL}; + + bool passed = true; + for (int i = 0; i < n_times; i++) { + sycl::program prg(q.get_context()); + + auto spec_const = prg.set_spec_constant((spec_const_t)DEF_VAL); + if (i % 2 != 0) + spec_const = prg.set_spec_constant((spec_const_t)REDEF_VAL); + + prg.build_with_kernel_type(); + + { + sycl::buffer buf(output.data(), output.size()); + + auto e = q.submit([&](sycl::handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task( + prg.get_kernel(), [=]() SYCL_ESIMD_KERNEL { + do_the_store(acc, i, spec_const.get()); + }); + }); + + e.wait(); + } + + if (output[i] != etalon[i]) { + passed = false; + std::cout << "comparison error -- case #" << i << " -- "; + std::cout << "output: " << output[i] << ", "; + std::cout << "etalon: " << etalon[i] << std::endl; + } + } + + if (passed) { + std::cout << "passed" << std::endl; + return 0; + } else { + std::cout << "FAILED" << std::endl; + return 1; + } +} diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp new file mode 100644 index 0000000000..833493265c --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_bool.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL true +#define REDEF_VAL false + +typedef bool spec_const_t; +typedef unsigned char container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + scalar_store(acc, i, val ? 1 : 0); +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp new file mode 100644 index 0000000000..6cc8c4472d --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_char.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL -22 +#define REDEF_VAL 33 + +typedef char spec_const_t; +typedef char container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + scalar_store(acc, i, val); +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp new file mode 100644 index 0000000000..4f2c03a50d --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_double.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL 9.1029384756e+11 +#define REDEF_VAL -1.4432211654e-10 + +typedef double spec_const_t; +typedef double container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + block_store(acc, i, simd{val}); // 2 doubles per 1 oword +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp new file mode 100644 index 0000000000..624d827d0f --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_float.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL -1.456789e-5 +#define REDEF_VAL 2.9865432e+5 + +typedef float spec_const_t; +typedef float container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + scalar_store(acc, i, val); +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp new file mode 100644 index 0000000000..331ae1415a --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_int.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL 100500 +#define REDEF_VAL -44556677 + +typedef int spec_const_t; +typedef int container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + scalar_store(acc, i, val); +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp new file mode 100644 index 0000000000..feeee13625 --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_long.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL -99776644220011 +#define REDEF_VAL 22001144668855 + +typedef long spec_const_t; +typedef long container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + block_store(acc, i, simd{val}); // 2 long int per 1 oword +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp new file mode 100644 index 0000000000..089267e7f6 --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_short.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL -30572 +#define REDEF_VAL 24794 + +typedef short spec_const_t; +typedef short container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + scalar_store(acc, i, val); +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp new file mode 100644 index 0000000000..7f54fbff5e --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_uchar.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL 128 +#define REDEF_VAL 33 + +typedef unsigned char spec_const_t; +typedef unsigned char container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + scalar_store(acc, i, val); +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp new file mode 100644 index 0000000000..672c7622e4 --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_uint.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL 0xdeadcafe +#define REDEF_VAL 0x4badbeaf + +typedef unsigned int spec_const_t; +typedef unsigned int container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + scalar_store(acc, i, val); +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp new file mode 100644 index 0000000000..bc4f7f6a5d --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_ulong.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL 0xdeaddeaf4badbeaf +#define REDEF_VAL 0x4cafebad00112233 + +typedef unsigned long spec_const_t; +typedef unsigned long container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + block_store(acc, i, simd{val}); // 2 long int per 1 oword +} + +#include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp new file mode 100644 index 0000000000..32282b51ab --- /dev/null +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -0,0 +1,31 @@ +//==--------------- spec_const_ushort.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 +// REQUIRES: linux && gpu +// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// UNSUPPORTED: cuda + +#define DEF_VAL 0xcafe +#define REDEF_VAL 0xdeaf + +typedef unsigned short spec_const_t; +typedef unsigned short container_t; + +#include "esimd_test_utils.hpp" + +#include +#include + +template +ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; + scalar_store(acc, i, val); +} + +#include "Inputs/spec_const_common.hpp" From d84e3b7ddd5719312102b0a45310b72caec5b7d8 Mon Sep 17 00:00:00 2001 From: Fedor Veselovskiy Date: Thu, 4 Feb 2021 23:04:24 +0300 Subject: [PATCH 16/39] Update pm_access_1.cpp --- SYCL/ESIMD/private_memory/pm_access_1.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/ESIMD/private_memory/pm_access_1.cpp b/SYCL/ESIMD/private_memory/pm_access_1.cpp index e35a8a7c71..0f8621009c 100644 --- a/SYCL/ESIMD/private_memory/pm_access_1.cpp +++ b/SYCL/ESIMD/private_memory/pm_access_1.cpp @@ -9,4 +9,4 @@ // REQUIRES: gpu // UNSUPPORTED: cuda // RUN: %clangxx-esimd -fsycl -I%S/.. %S/Inputs/pm_common.cpp -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 1 \ No newline at end of file +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 1 From a47856d702f46179ee551badd845425a3a1d4051 Mon Sep 17 00:00:00 2001 From: Fedor Veselovskiy Date: Thu, 4 Feb 2021 23:04:47 +0300 Subject: [PATCH 17/39] Update pm_access_2.cpp --- SYCL/ESIMD/private_memory/pm_access_2.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/ESIMD/private_memory/pm_access_2.cpp b/SYCL/ESIMD/private_memory/pm_access_2.cpp index efb3987459..1afa615d66 100644 --- a/SYCL/ESIMD/private_memory/pm_access_2.cpp +++ b/SYCL/ESIMD/private_memory/pm_access_2.cpp @@ -9,4 +9,4 @@ // REQUIRES: gpu // UNSUPPORTED: cuda // RUN: %clangxx-esimd -fsycl -I%S/.. %S/Inputs/pm_common.cpp -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 2 \ No newline at end of file +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 2 From 309968df48b855a61d844259312b3e232d4f1c2e Mon Sep 17 00:00:00 2001 From: Fedor Veselovskiy Date: Thu, 4 Feb 2021 23:05:01 +0300 Subject: [PATCH 18/39] Update pm_access_3.cpp --- SYCL/ESIMD/private_memory/pm_access_3.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/ESIMD/private_memory/pm_access_3.cpp b/SYCL/ESIMD/private_memory/pm_access_3.cpp index 20345ce74a..274d06a897 100644 --- a/SYCL/ESIMD/private_memory/pm_access_3.cpp +++ b/SYCL/ESIMD/private_memory/pm_access_3.cpp @@ -9,4 +9,4 @@ // REQUIRES: gpu // UNSUPPORTED: cuda // RUN: %clangxx-esimd -fsycl -I%S/.. %S/Inputs/pm_common.cpp -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 3 \ No newline at end of file +// RUN: %ESIMD_RUN_PLACEHOLDER %t.out 3 From 2061249e46dfc289a13e07585fe4df76a8f5e214 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Thu, 4 Feb 2021 23:20:40 +0300 Subject: [PATCH 19/39] clang-format patch --- SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index a2c0e38776..6f96834e7d 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -41,10 +41,10 @@ int main(int argc, char **argv) { auto e = q.submit([&](sycl::handler &cgh) { auto acc = buf.get_access(cgh); - cgh.single_task( - prg.get_kernel(), [=]() SYCL_ESIMD_KERNEL { - do_the_store(acc, i, spec_const.get()); - }); + cgh.single_task(prg.get_kernel(), + [=]() SYCL_ESIMD_KERNEL { + do_the_store(acc, i, spec_const.get()); + }); }); e.wait(); From e6cc0d3ea20388e58dc1a9e46b8ccb73eb7dea1e Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 5 Feb 2021 05:11:28 +0300 Subject: [PATCH 20/39] more generalization --- .../spec_const/Inputs/spec_const_common.hpp | 24 +++++++++++++++++-- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_char.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_double.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_float.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_int.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_long.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_short.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 12 +--------- SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 12 +--------- 12 files changed, 33 insertions(+), 123 deletions(-) diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index 6f96834e7d..b221129bf3 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -9,10 +9,30 @@ // basic types, particularly a specialization constant can be redifined and // correct new value is used after redefinition. +#include "esimd_test_utils.hpp" + +#include +#include + #include #include -using namespace sycl; +using namespace cl::sycl; + +template +ESIMD_INLINE void do_store(AccessorTy acc, int i, spec_const_t val) { + using namespace sycl::INTEL::gpu; +#if STORE == 0 + // bool + scalar_store(acc, i, val ? 1 : 0); +#elif STORE == 1 + // block + block_store(acc, i, simd{val}); +#elif STORE == 2 + // scatter + scalar_store(acc, i, val); +#endif +} class ConstID; class TestKernel; @@ -43,7 +63,7 @@ int main(int argc, char **argv) { auto acc = buf.get_access(cgh); cgh.single_task(prg.get_kernel(), [=]() SYCL_ESIMD_KERNEL { - do_the_store(acc, i, spec_const.get()); + do_store(acc, i, spec_const.get()); }); }); diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index 833493265c..d109a573a3 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -13,19 +13,9 @@ #define DEF_VAL true #define REDEF_VAL false +#define STORE 0 typedef bool spec_const_t; typedef unsigned char container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - scalar_store(acc, i, val ? 1 : 0); -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 6cc8c4472d..91205ec781 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -13,19 +13,9 @@ #define DEF_VAL -22 #define REDEF_VAL 33 +#define STORE 2 typedef char spec_const_t; typedef char container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - scalar_store(acc, i, val); -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index 4f2c03a50d..e9002d17d6 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -13,19 +13,9 @@ #define DEF_VAL 9.1029384756e+11 #define REDEF_VAL -1.4432211654e-10 +#define STORE 1 typedef double spec_const_t; typedef double container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - block_store(acc, i, simd{val}); // 2 doubles per 1 oword -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 624d827d0f..9eb0ef94c3 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -13,19 +13,9 @@ #define DEF_VAL -1.456789e-5 #define REDEF_VAL 2.9865432e+5 +#define STORE 2 typedef float spec_const_t; typedef float container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - scalar_store(acc, i, val); -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index 331ae1415a..26e296316d 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -13,19 +13,9 @@ #define DEF_VAL 100500 #define REDEF_VAL -44556677 +#define STORE 2 typedef int spec_const_t; typedef int container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - scalar_store(acc, i, val); -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index feeee13625..c48585e5a0 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -13,19 +13,9 @@ #define DEF_VAL -99776644220011 #define REDEF_VAL 22001144668855 +#define STORE 1 typedef long spec_const_t; typedef long container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - block_store(acc, i, simd{val}); // 2 long int per 1 oword -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index 089267e7f6..695d6da6a2 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -13,19 +13,9 @@ #define DEF_VAL -30572 #define REDEF_VAL 24794 +#define STORE 2 typedef short spec_const_t; typedef short container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - scalar_store(acc, i, val); -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index 7f54fbff5e..3b1c99b908 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -13,19 +13,9 @@ #define DEF_VAL 128 #define REDEF_VAL 33 +#define STORE 2 typedef unsigned char spec_const_t; typedef unsigned char container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - scalar_store(acc, i, val); -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index 672c7622e4..fecbcc8172 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -13,19 +13,9 @@ #define DEF_VAL 0xdeadcafe #define REDEF_VAL 0x4badbeaf +#define STORE 2 typedef unsigned int spec_const_t; typedef unsigned int container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - scalar_store(acc, i, val); -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index bc4f7f6a5d..6b59947bea 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -13,19 +13,9 @@ #define DEF_VAL 0xdeaddeaf4badbeaf #define REDEF_VAL 0x4cafebad00112233 +#define STORE 1 typedef unsigned long spec_const_t; typedef unsigned long container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - block_store(acc, i, simd{val}); // 2 long int per 1 oword -} - #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index 32282b51ab..3da1626495 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -13,19 +13,9 @@ #define DEF_VAL 0xcafe #define REDEF_VAL 0xdeaf +#define STORE 2 typedef unsigned short spec_const_t; typedef unsigned short container_t; -#include "esimd_test_utils.hpp" - -#include -#include - -template -ESIMD_INLINE void do_the_store(AccessorTy acc, int i, spec_const_t val) { - using namespace sycl::INTEL::gpu; - scalar_store(acc, i, val); -} - #include "Inputs/spec_const_common.hpp" From 371dd8441795713032bcd5fbba0dbe822a5db446 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 5 Feb 2021 05:19:48 +0300 Subject: [PATCH 21/39] typo fix --- SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index b221129bf3..f8994d2e59 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -29,7 +29,7 @@ ESIMD_INLINE void do_store(AccessorTy acc, int i, spec_const_t val) { // block block_store(acc, i, simd{val}); #elif STORE == 2 - // scatter + // scalar scalar_store(acc, i, val); #endif } From c0684b9b3b6b37940353f319c481abc77701b8a1 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 5 Feb 2021 05:29:40 +0300 Subject: [PATCH 22/39] clang-format patch --- SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index f8994d2e59..c448313975 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -61,10 +61,9 @@ int main(int argc, char **argv) { auto e = q.submit([&](sycl::handler &cgh) { auto acc = buf.get_access(cgh); - cgh.single_task(prg.get_kernel(), - [=]() SYCL_ESIMD_KERNEL { - do_store(acc, i, spec_const.get()); - }); + cgh.single_task( + prg.get_kernel(), + [=]() SYCL_ESIMD_KERNEL { do_store(acc, i, spec_const.get()); }); }); e.wait(); From 04a4069c221e59f5e539fda6b576a2c6ea2a3476 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Sat, 6 Feb 2021 16:54:24 +0300 Subject: [PATCH 23/39] cosmetic changes --- SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index c448313975..eb58b361d0 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -39,8 +39,9 @@ class TestKernel; int main(int argc, char **argv) { queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()); + auto dev = q.get_device(); - auto ctx = q.get_context(); + std::cout << "Running on " << dev.get_info() << "\n"; const int n_times = 2; std::vector output(n_times); @@ -80,8 +81,8 @@ int main(int argc, char **argv) { if (passed) { std::cout << "passed" << std::endl; return 0; - } else { - std::cout << "FAILED" << std::endl; - return 1; } + + std::cout << "FAILED" << std::endl; + return 1; } From 2772abba6db8dd49c354050121007cf1c8212274 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Mon, 8 Feb 2021 23:33:03 +0300 Subject: [PATCH 24/39] set expect fail --- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_char.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_double.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_float.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_int.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_long.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_short.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 3 ++- 11 files changed, 15 insertions(+), 11 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index d109a573a3..f014d86611 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -8,7 +8,7 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda #define DEF_VAL true diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 91205ec781..0971e54dbc 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -8,8 +8,9 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +// XFAIL: * #define DEF_VAL -22 #define REDEF_VAL 33 diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index e9002d17d6..d64bcf252a 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -8,7 +8,7 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda #define DEF_VAL 9.1029384756e+11 diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 9eb0ef94c3..08b8d0a576 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -8,7 +8,7 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda #define DEF_VAL -1.456789e-5 diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index 26e296316d..f773d6ae7f 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -8,7 +8,7 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda #define DEF_VAL 100500 diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index c48585e5a0..a09afefa85 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -8,7 +8,7 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda #define DEF_VAL -99776644220011 diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index 695d6da6a2..bc98878f28 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -8,8 +8,9 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +// XFAIL: * #define DEF_VAL -30572 #define REDEF_VAL 24794 diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index 3b1c99b908..63d3af33d0 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -8,8 +8,9 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +// XFAIL: * #define DEF_VAL 128 #define REDEF_VAL 33 diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index fecbcc8172..d0c1ef9393 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -8,7 +8,7 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda #define DEF_VAL 0xdeadcafe diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index 6b59947bea..7972e414dd 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -8,7 +8,7 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda #define DEF_VAL 0xdeaddeaf4badbeaf diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index 3da1626495..af1c95a562 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -8,8 +8,9 @@ // TODO enable on Windows // REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out -// RUN: %ESIMD_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +// XFAIL: * #define DEF_VAL 0xcafe #define REDEF_VAL 0xdeaf From 59cb60cb3ba2b9fc225763da2b27b9846eee5409 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 00:19:51 +0300 Subject: [PATCH 25/39] enable windows --- SYCL/ESIMD/spec_const/spec_const_char.cpp | 5 ++--- SYCL/ESIMD/spec_const/spec_const_double.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_float.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_int.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_long.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_short.cpp | 5 ++--- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 5 ++--- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 5 ++--- 10 files changed, 14 insertions(+), 24 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 0971e54dbc..587893500a 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -5,12 +5,11 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu +// XFAIL: * // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -// XFAIL: * #define DEF_VAL -22 #define REDEF_VAL 33 diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index d64bcf252a..109f5be5db 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 08b8d0a576..2705f66cb9 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index f773d6ae7f..3ac0a5ed2d 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index a09afefa85..55afd58e0b 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index bc98878f28..6c56fd7ca4 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -5,12 +5,11 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu +// XFAIL: * // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -// XFAIL: * #define DEF_VAL -30572 #define REDEF_VAL 24794 diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index 63d3af33d0..976c616225 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -5,12 +5,11 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu +// XFAIL: * // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -// XFAIL: * #define DEF_VAL 128 #define REDEF_VAL 33 diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index d0c1ef9393..54d86227c4 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index 7972e414dd..4d939c7967 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index af1c95a562..fcc5ead646 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -5,12 +5,11 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu +// XFAIL: * // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -// XFAIL: * #define DEF_VAL 0xcafe #define REDEF_VAL 0xdeaf From f477d473693c095dcacf7a65538676fbcab9f53d Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 01:00:50 +0300 Subject: [PATCH 26/39] disable windows and remove xfail --- SYCL/ESIMD/spec_const/spec_const_char.cpp | 4 ++-- SYCL/ESIMD/spec_const/spec_const_double.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_float.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_int.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_long.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_short.cpp | 4 ++-- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 4 ++-- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 4 ++-- 10 files changed, 20 insertions(+), 14 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 587893500a..7a53a765f0 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -5,8 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu -// XFAIL: * +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index 109f5be5db..d64bcf252a 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -5,7 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 2705f66cb9..08b8d0a576 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -5,7 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index 3ac0a5ed2d..f773d6ae7f 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -5,7 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index 55afd58e0b..a09afefa85 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -5,7 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index 6c56fd7ca4..6f3375bc7d 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -5,8 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu -// XFAIL: * +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index 976c616225..c4348aa042 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -5,8 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu -// XFAIL: * +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index 54d86227c4..d0c1ef9393 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -5,7 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index 4d939c7967..7972e414dd 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -5,7 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index fcc5ead646..136543d4bf 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -5,8 +5,8 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// REQUIRES: gpu -// XFAIL: * +// TODO enable on Windows +// REQUIRES: linux && gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda From cc71d2013d57cb4081854edcacf44b4742156b44 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 01:44:35 +0300 Subject: [PATCH 27/39] set xfail for level_zero --- SYCL/ESIMD/spec_const/spec_const_char.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_short.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 1 + 4 files changed, 4 insertions(+) diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 7a53a765f0..65bfe01368 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// // TODO enable on Windows // REQUIRES: linux && gpu +// XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index 6f3375bc7d..8fe3638fed 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// // TODO enable on Windows // REQUIRES: linux && gpu +// XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index c4348aa042..cf9e0ad536 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// // TODO enable on Windows // REQUIRES: linux && gpu +// XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index 136543d4bf..f41c8adadf 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// // TODO enable on Windows // REQUIRES: linux && gpu +// XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda From 3635f1a03a1b119fa30665911b79f0feef36d5c4 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 17:33:00 +0300 Subject: [PATCH 28/39] handle synchronous SYCL exceptions; remove unnecessary code; add comment --- SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp | 9 +++++---- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 2 ++ 2 files changed, 7 insertions(+), 4 deletions(-) diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index eb58b361d0..05201ece7a 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -57,17 +57,18 @@ int main(int argc, char **argv) { prg.build_with_kernel_type(); - { + try { sycl::buffer buf(output.data(), output.size()); - auto e = q.submit([&](sycl::handler &cgh) { + q.submit([&](sycl::handler &cgh) { auto acc = buf.get_access(cgh); cgh.single_task( prg.get_kernel(), [=]() SYCL_ESIMD_KERNEL { do_store(acc, i, spec_const.get()); }); }); - - e.wait(); + } catch (cl::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + return e.get_cl_code(); } if (output[i] != etalon[i]) { diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index f014d86611..90cd8a3926 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -15,6 +15,8 @@ #define REDEF_VAL false #define STORE 0 +// In this case container type is set to unsigned char to be able to use +// esimd memory interfaces to pollute container. typedef bool spec_const_t; typedef unsigned char container_t; From 6e482b21fdcd41d0ac388361ff045f9543198fb3 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 17:38:00 +0300 Subject: [PATCH 29/39] enable windows to run jenkins check(will be reverted) --- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_char.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_double.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_float.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_int.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_long.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_short.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 3 +-- SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 3 +-- 11 files changed, 11 insertions(+), 22 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index 90cd8a3926..34e2f84547 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 65bfe01368..3886dc4094 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index d64bcf252a..109f5be5db 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 08b8d0a576..2705f66cb9 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index f773d6ae7f..3ac0a5ed2d 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index a09afefa85..55afd58e0b 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index 8fe3638fed..2ab3846019 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index cf9e0ad536..de0be3c077 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index d0c1ef9393..54d86227c4 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index 7972e414dd..4d939c7967 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index f41c8adadf..0c896da9da 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -5,8 +5,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// TODO enable on Windows -// REQUIRES: linux && gpu +// REQUIRES: gpu // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out From 845bd2a7d83631d5e8a8b5685436f1005ec69e49 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Tue, 9 Feb 2021 18:56:11 +0300 Subject: [PATCH 30/39] set expect fail for Windows --- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_char.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_double.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_float.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_int.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_long.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_short.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 1 + SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 2 +- 11 files changed, 11 insertions(+), 4 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index 34e2f84547..5a1e325e2d 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 3886dc4094..597e587071 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: level_zero +// XFAIL: level_zero || windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index 109f5be5db..ce9198d2f0 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 2705f66cb9..4829975d1b 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index 3ac0a5ed2d..eb4001c1a5 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index 55afd58e0b..4e3b123b39 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index 2ab3846019..f3bd8d63c6 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: level_zero +// XFAIL: level_zero || windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index de0be3c077..fa8fc2f3ac 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: level_zero +// XFAIL: level_zero || windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index 54d86227c4..fc15fb5917 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index 4d939c7967..834229a614 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index 0c896da9da..ec281a1dd3 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: level_zero +// XFAIL: level_zero || windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda From 4bc6ca278b81abfc28e813eb90a44c40be1c3e1f Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Wed, 10 Feb 2021 00:44:25 +0300 Subject: [PATCH 31/39] comments and C++ re-style --- .../spec_const/Inputs/spec_const_common.hpp | 22 ++++++++++++++----- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 6 +++-- SYCL/ESIMD/spec_const/spec_const_char.cpp | 6 +++-- SYCL/ESIMD/spec_const/spec_const_double.cpp | 6 +++-- SYCL/ESIMD/spec_const/spec_const_float.cpp | 6 +++-- SYCL/ESIMD/spec_const/spec_const_int.cpp | 7 +++--- SYCL/ESIMD/spec_const/spec_const_long.cpp | 11 +++++----- SYCL/ESIMD/spec_const/spec_const_short.cpp | 6 +++-- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 6 +++-- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 6 +++-- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 11 +++++----- SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 6 +++-- 12 files changed, 65 insertions(+), 34 deletions(-) diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index 05201ece7a..ea418aac43 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -22,13 +22,16 @@ using namespace cl::sycl; template ESIMD_INLINE void do_store(AccessorTy acc, int i, spec_const_t val) { using namespace sycl::INTEL::gpu; -#if STORE == 0 + // scatter function, that is used in scalar_store, can only process types + // whose size is no more than 4 bytes. +#if (STORE == 0) // bool scalar_store(acc, i, val ? 1 : 0); -#elif STORE == 1 +#elif (STORE == 1) // block block_store(acc, i, simd{val}); -#elif STORE == 2 +#else + static_assert(STORE == 2, "Unspecified store"); // scalar scalar_store(acc, i, val); #endif @@ -43,14 +46,23 @@ int main(int argc, char **argv) { auto dev = q.get_device(); std::cout << "Running on " << dev.get_info() << "\n"; - const int n_times = 2; - std::vector output(n_times); std::vector etalon = {DEF_VAL, REDEF_VAL}; + const size_t n_times = etalon.size(); + std::vector output(n_times); bool passed = true; for (int i = 0; i < n_times; i++) { sycl::program prg(q.get_context()); + // Checking that already initialized constant can be overwritten. + // According to standards proposals: + // A cl::sycl::experimental::spec_constant object is considered + // initialized once the result of a cl::sycl::program::set_spec_constant + // is assigned to it. + // A specialization constant value can be overwritten if the program was + // not built before by recalling set_spec_constant with the same ID and + // the new value. Although the type T of the specialization constant must + // remain the same. auto spec_const = prg.set_spec_constant((spec_const_t)DEF_VAL); if (i % 2 != 0) spec_const = prg.set_spec_constant((spec_const_t)REDEF_VAL); diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index 5a1e325e2d..d3536df80b 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -11,13 +11,15 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +#include + #define DEF_VAL true #define REDEF_VAL false #define STORE 0 // In this case container type is set to unsigned char to be able to use // esimd memory interfaces to pollute container. -typedef bool spec_const_t; -typedef unsigned char container_t; +using spec_const_t = bool; +using container_t = uint8_t; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 597e587071..b3c849d633 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -11,11 +11,13 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +#include + #define DEF_VAL -22 #define REDEF_VAL 33 #define STORE 2 -typedef char spec_const_t; -typedef char container_t; +using spec_const_t = int8_t; +using container_t = int8_t; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index ce9198d2f0..6faae972d4 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -11,11 +11,13 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +#include + #define DEF_VAL 9.1029384756e+11 #define REDEF_VAL -1.4432211654e-10 #define STORE 1 -typedef double spec_const_t; -typedef double container_t; +using spec_const_t = double; +using container_t = double; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 4829975d1b..5643782e79 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -11,11 +11,13 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +#include + #define DEF_VAL -1.456789e-5 #define REDEF_VAL 2.9865432e+5 #define STORE 2 -typedef float spec_const_t; -typedef float container_t; +using spec_const_t = float; +using container_t = float; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index eb4001c1a5..f494fda7b8 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -6,16 +6,17 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +#include + #define DEF_VAL 100500 #define REDEF_VAL -44556677 #define STORE 2 -typedef int spec_const_t; -typedef int container_t; +using spec_const_t = int32_t; +using container_t = int32_t; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index 4e3b123b39..df3afeca0f 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -6,16 +6,17 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#define DEF_VAL -99776644220011 -#define REDEF_VAL 22001144668855 +#include + +#define DEF_VAL -99776644220011ll +#define REDEF_VAL 22001144668855ll #define STORE 1 -typedef long spec_const_t; -typedef long container_t; +using spec_const_t = int64_t; +using container_t = int64_t; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index f3bd8d63c6..923ae92724 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -11,11 +11,13 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +#include + #define DEF_VAL -30572 #define REDEF_VAL 24794 #define STORE 2 -typedef short spec_const_t; -typedef short container_t; +using spec_const_t = int16_t; +using container_t = int16_t; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index fa8fc2f3ac..182752fe11 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -11,11 +11,13 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +#include + #define DEF_VAL 128 #define REDEF_VAL 33 #define STORE 2 -typedef unsigned char spec_const_t; -typedef unsigned char container_t; +using spec_const_t = uint8_t; +using container_t = uint8_t; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index fc15fb5917..e4a5e2a1f3 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -11,11 +11,13 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +#include + #define DEF_VAL 0xdeadcafe #define REDEF_VAL 0x4badbeaf #define STORE 2 -typedef unsigned int spec_const_t; -typedef unsigned int container_t; +using spec_const_t = uint32_t; +using container_t = uint32_t; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index 834229a614..a27adb87a9 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -6,16 +6,17 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#define DEF_VAL 0xdeaddeaf4badbeaf -#define REDEF_VAL 0x4cafebad00112233 +#include + +#define DEF_VAL 0xdeaddeaf4badbeafull +#define REDEF_VAL 0x4cafebad00112233ull #define STORE 1 -typedef unsigned long spec_const_t; -typedef unsigned long container_t; +using spec_const_t = uint64_t; +using container_t = uint64_t; #include "Inputs/spec_const_common.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index ec281a1dd3..622007549d 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -11,11 +11,13 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda +#include + #define DEF_VAL 0xcafe #define REDEF_VAL 0xdeaf #define STORE 2 -typedef unsigned short spec_const_t; -typedef unsigned short container_t; +using spec_const_t = uint16_t; +using container_t = uint16_t; #include "Inputs/spec_const_common.hpp" From cf53dd9ee12c0c7a1e0cc9630db95c0c36c3d521 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Wed, 10 Feb 2021 01:05:38 +0300 Subject: [PATCH 32/39] clang-format patch --- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_char.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_double.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_float.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_int.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_long.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_short.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 2 +- 11 files changed, 11 insertions(+), 11 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index d3536df80b..8ed6fac93f 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -11,7 +11,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL true #define REDEF_VAL false diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index b3c849d633..3c7ac72186 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -11,7 +11,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL -22 #define REDEF_VAL 33 diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index 6faae972d4..239ad1e794 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -11,7 +11,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL 9.1029384756e+11 #define REDEF_VAL -1.4432211654e-10 diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 5643782e79..638a2d55ad 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -11,7 +11,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL -1.456789e-5 #define REDEF_VAL 2.9865432e+5 diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index f494fda7b8..3bb7c041a4 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -10,7 +10,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL 100500 #define REDEF_VAL -44556677 diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index df3afeca0f..54e0117827 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -10,7 +10,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL -99776644220011ll #define REDEF_VAL 22001144668855ll diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index 923ae92724..a9bd9b28c2 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -11,7 +11,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL -30572 #define REDEF_VAL 24794 diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index 182752fe11..4a85fbde88 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -11,7 +11,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL 128 #define REDEF_VAL 33 diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index e4a5e2a1f3..ec4e20fdfb 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -11,7 +11,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL 0xdeadcafe #define REDEF_VAL 0x4badbeaf diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index a27adb87a9..0cecc7b504 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -10,7 +10,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL 0xdeaddeaf4badbeafull #define REDEF_VAL 0x4cafebad00112233ull diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index 622007549d..ac28a0a43b 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -11,7 +11,7 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda -#include +#include #define DEF_VAL 0xcafe #define REDEF_VAL 0xdeaf From 5441b6456ae18e81feadb3ebb595c93c1fb35c38 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Wed, 10 Feb 2021 15:56:04 +0300 Subject: [PATCH 33/39] std exception handle --- SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index ea418aac43..aae2c18cfe 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -81,6 +81,9 @@ int main(int argc, char **argv) { } catch (cl::sycl::exception const &e) { std::cout << "SYCL exception caught: " << e.what() << '\n'; return e.get_cl_code(); + } catch (std::exception const &e) { + std::cout << "General exception caught: " << e.what() << '\n'; + return 2; } if (output[i] != etalon[i]) { From dc94ab4c9b90833822393c511d5eb0b62982f079 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Wed, 10 Feb 2021 16:56:47 +0300 Subject: [PATCH 34/39] move spec const init into try block --- .../spec_const/Inputs/spec_const_common.hpp | 37 +++++++++---------- 1 file changed, 17 insertions(+), 20 deletions(-) diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp index aae2c18cfe..6d13598d74 100644 --- a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp +++ b/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp @@ -52,24 +52,24 @@ int main(int argc, char **argv) { bool passed = true; for (int i = 0; i < n_times; i++) { - sycl::program prg(q.get_context()); - - // Checking that already initialized constant can be overwritten. - // According to standards proposals: - // A cl::sycl::experimental::spec_constant object is considered - // initialized once the result of a cl::sycl::program::set_spec_constant - // is assigned to it. - // A specialization constant value can be overwritten if the program was - // not built before by recalling set_spec_constant with the same ID and - // the new value. Although the type T of the specialization constant must - // remain the same. - auto spec_const = prg.set_spec_constant((spec_const_t)DEF_VAL); - if (i % 2 != 0) - spec_const = prg.set_spec_constant((spec_const_t)REDEF_VAL); - - prg.build_with_kernel_type(); - try { + sycl::program prg(q.get_context()); + + // Checking that already initialized constant can be overwritten. + // According to standards proposals: + // A cl::sycl::experimental::spec_constant object is considered + // initialized once the result of a cl::sycl::program::set_spec_constant + // is assigned to it. + // A specialization constant value can be overwritten if the program was + // not built before by recalling set_spec_constant with the same ID and + // the new value. Although the type T of the specialization constant + // must remain the same. + auto spec_const = prg.set_spec_constant((spec_const_t)DEF_VAL); + if (i % 2 != 0) + spec_const = prg.set_spec_constant((spec_const_t)REDEF_VAL); + + prg.build_with_kernel_type(); + sycl::buffer buf(output.data(), output.size()); q.submit([&](sycl::handler &cgh) { @@ -81,9 +81,6 @@ int main(int argc, char **argv) { } catch (cl::sycl::exception const &e) { std::cout << "SYCL exception caught: " << e.what() << '\n'; return e.get_cl_code(); - } catch (std::exception const &e) { - std::cout << "General exception caught: " << e.what() << '\n'; - return 2; } if (output[i] != etalon[i]) { From bdcdc3ce049775a403ee55be90d8ce2c9b931eff Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Wed, 10 Feb 2021 20:50:32 +0300 Subject: [PATCH 35/39] set expect fail and unsupported status --- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_char.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_double.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_float.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_int.cpp | 2 ++ SYCL/ESIMD/spec_const/spec_const_long.cpp | 2 ++ SYCL/ESIMD/spec_const/spec_const_short.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 3 ++- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 2 ++ SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 5 ++++- 11 files changed, 30 insertions(+), 8 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index 8ed6fac93f..756a505d5f 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -6,7 +6,8 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: windows +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 3c7ac72186..e704116198 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: level_zero || windows +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows +// Linux Level Zero fail: VMIT-8680 +// XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index 239ad1e794..ee37f22e69 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -6,7 +6,8 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: windows +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 638a2d55ad..4de4d6ed58 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -6,7 +6,8 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: windows +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index 3bb7c041a4..523c3d8c53 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -6,6 +6,8 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index 54e0117827..2f4a1718e6 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -6,6 +6,8 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index a9bd9b28c2..cad507b970 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: level_zero || windows +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows +// Linux Level Zero fail: VMIT-8680 +// XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index 4a85fbde88..de2dea3230 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: level_zero || windows +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows +// Linux Level Zero fail: VMIT-8680 +// XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index ec4e20fdfb..c275e6db70 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -6,7 +6,8 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: windows +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index 0cecc7b504..b33828f397 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -6,6 +6,8 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index ac28a0a43b..48aef3d67e 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// XFAIL: level_zero || windows +// Windows fail: XDEPS-1100 +// UNSUPPORTED: windows +// Linux Level Zero fail: VMIT-8680 +// XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda From 63237bdd384c9d9b0ecf59925550b240859380b2 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Thu, 11 Feb 2021 15:37:43 +0300 Subject: [PATCH 36/39] add unsupported and xfail description --- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_char.cpp | 7 +++++-- SYCL/ESIMD/spec_const/spec_const_double.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_float.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_int.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_long.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_short.cpp | 7 +++++-- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 7 +++++-- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 5 ++++- SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 7 +++++-- 11 files changed, 48 insertions(+), 15 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index 756a505d5f..77c2d4ddd6 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index e704116198..d64cf1a5d0 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -6,9 +6,12 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows -// Linux Level Zero fail: VMIT-8680 +// Linux Level Zero fail with assertion in SPIRV about specialization constant type size. // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index ee37f22e69..67fe4ed5ab 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 4de4d6ed58..f4a8a3e229 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index 523c3d8c53..9e12c40731 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index 2f4a1718e6..8d70aa0bf0 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index cad507b970..0b2f2b7b24 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -6,9 +6,12 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows -// Linux Level Zero fail: VMIT-8680 +// Linux Level Zero fail with assertion in SPIRV about specialization constant type size. // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index de2dea3230..915e6874c9 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -6,9 +6,12 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows -// Linux Level Zero fail: VMIT-8680 +// Linux Level Zero fail with assertion in SPIRV about specialization constant type size. // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index c275e6db70..e5627b671c 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index b33828f397..b6db5f77fc 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -6,7 +6,10 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index 48aef3d67e..745b6ff551 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -6,9 +6,12 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// Windows fail: XDEPS-1100 +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. +// This translator doesn't have the ability to overwrite the default specialization constant value. +// That is why the support in Windows driver is disabled at all. +// This feature will start working on Windows when the llvm version is switched to 9. // UNSUPPORTED: windows -// Linux Level Zero fail: VMIT-8680 +// Linux Level Zero fail with assertion in SPIRV about specialization constant type size. // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out From dec083693d19e539ee1388aee80494fa0bfc44bd Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Thu, 11 Feb 2021 15:48:05 +0300 Subject: [PATCH 37/39] cosmetic changes --- SYCL/ESIMD/spec_const/spec_const_bool.cpp | 9 +++++---- SYCL/ESIMD/spec_const/spec_const_char.cpp | 12 +++++++----- SYCL/ESIMD/spec_const/spec_const_double.cpp | 9 +++++---- SYCL/ESIMD/spec_const/spec_const_float.cpp | 9 +++++---- SYCL/ESIMD/spec_const/spec_const_int.cpp | 9 +++++---- SYCL/ESIMD/spec_const/spec_const_long.cpp | 9 +++++---- SYCL/ESIMD/spec_const/spec_const_short.cpp | 12 +++++++----- SYCL/ESIMD/spec_const/spec_const_uchar.cpp | 12 +++++++----- SYCL/ESIMD/spec_const/spec_const_uint.cpp | 9 +++++---- SYCL/ESIMD/spec_const/spec_const_ulong.cpp | 9 +++++---- SYCL/ESIMD/spec_const/spec_const_ushort.cpp | 12 +++++++----- 11 files changed, 63 insertions(+), 48 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index 77c2d4ddd6..44d558e425 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -6,10 +6,11 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index d64cf1a5d0..8828985c2c 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -6,12 +6,14 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows -// Linux Level Zero fail with assertion in SPIRV about specialization constant type size. +// Linux Level Zero fail with assertion in SPIRV about specialization constant +// type size. // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index 67fe4ed5ab..2e53ed0a9f 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -6,10 +6,11 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index f4a8a3e229..0217160659 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -6,10 +6,11 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index 9e12c40731..7882021da6 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -6,10 +6,11 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_long.cpp index 8d70aa0bf0..900d377283 100644 --- a/SYCL/ESIMD/spec_const/spec_const_long.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_long.cpp @@ -6,10 +6,11 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index 0b2f2b7b24..2099ce43ab 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -6,12 +6,14 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows -// Linux Level Zero fail with assertion in SPIRV about specialization constant type size. +// Linux Level Zero fail with assertion in SPIRV about specialization constant +// type size. // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index 915e6874c9..776840c96b 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -6,12 +6,14 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows -// Linux Level Zero fail with assertion in SPIRV about specialization constant type size. +// Linux Level Zero fail with assertion in SPIRV about specialization constant +// type size. // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index e5627b671c..5461b90a64 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -6,10 +6,11 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp index b6db5f77fc..ebb6b9622d 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ulong.cpp @@ -6,10 +6,11 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index 745b6ff551..8515b6cb29 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -6,12 +6,14 @@ // //===----------------------------------------------------------------------===// // REQUIRES: gpu -// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 based spirv translator. -// This translator doesn't have the ability to overwrite the default specialization constant value. -// That is why the support in Windows driver is disabled at all. -// This feature will start working on Windows when the llvm version is switched to 9. +// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7 +// based spirv translator. This translator doesn't have the ability to overwrite +// the default specialization constant value. That is why the support in Windows +// driver is disabled at all. This feature will start working on Windows when +// the llvm version is switched to 9. // UNSUPPORTED: windows -// Linux Level Zero fail with assertion in SPIRV about specialization constant type size. +// Linux Level Zero fail with assertion in SPIRV about specialization constant +// type size. // XFAIL: level_zero // RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out From 3ad7ffe52d237aca037666f8738aaceff39e228b Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 12 Feb 2021 13:58:51 +0300 Subject: [PATCH 38/39] rename to int64 --- .../spec_const/{spec_const_long.cpp => spec_const_int64.cpp} | 0 .../spec_const/{spec_const_ulong.cpp => spec_const_uint64.cpp} | 0 2 files changed, 0 insertions(+), 0 deletions(-) rename SYCL/ESIMD/spec_const/{spec_const_long.cpp => spec_const_int64.cpp} (100%) rename SYCL/ESIMD/spec_const/{spec_const_ulong.cpp => spec_const_uint64.cpp} (100%) diff --git a/SYCL/ESIMD/spec_const/spec_const_long.cpp b/SYCL/ESIMD/spec_const/spec_const_int64.cpp similarity index 100% rename from SYCL/ESIMD/spec_const/spec_const_long.cpp rename to SYCL/ESIMD/spec_const/spec_const_int64.cpp diff --git a/SYCL/ESIMD/spec_const/spec_const_ulong.cpp b/SYCL/ESIMD/spec_const/spec_const_uint64.cpp similarity index 100% rename from SYCL/ESIMD/spec_const/spec_const_ulong.cpp rename to SYCL/ESIMD/spec_const/spec_const_uint64.cpp From d27ab0d70c8bc5cd515a61bb7d522020a7ab62b9 Mon Sep 17 00:00:00 2001 From: Fedor Veselovsky Date: Fri, 12 Feb 2021 15:30:48 +0300 Subject: [PATCH 39/39] cosmetic fix --- SYCL/ESIMD/spec_const/spec_const_int64.cpp | 2 +- SYCL/ESIMD/spec_const/spec_const_uint64.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/SYCL/ESIMD/spec_const/spec_const_int64.cpp b/SYCL/ESIMD/spec_const/spec_const_int64.cpp index 900d377283..bebd2c947b 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int64.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int64.cpp @@ -1,4 +1,4 @@ -//==--------------- spec_const_long.cpp - DPC++ ESIMD on-device test -----===// +//==-------------- spec_const_int64.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/spec_const/spec_const_uint64.cpp b/SYCL/ESIMD/spec_const/spec_const_uint64.cpp index ebb6b9622d..368c8bb087 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint64.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint64.cpp @@ -1,4 +1,4 @@ -//==--------------- spec_const_ulong.cpp - DPC++ ESIMD on-device test ----===// +//==-------------- spec_const_uint64.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.