diff --git a/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp new file mode 100644 index 0000000000..046a767dd1 --- /dev/null +++ b/SYCL/ESIMD/regression/sycl_esimd_mixed_unnamed.cpp @@ -0,0 +1,123 @@ +//==--- sycl_esimd_mixed_unnamed.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 +// +//===----------------------------------------------------------------------===// +// This is basic test for mixing unnamed SYCL and ESIMD kernels in the same +// source and in the same program . + +// REQUIRES: gpu +// UNSUPPORTED: cuda || hip +// UNSUPPORTED: esimd_emulator +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// XFAIL: esimd_emulator + +#include +#include +#include + +using namespace ::sycl; + +bool checkResult(const std::vector &A, int Inc) { + int err_cnt = 0; + unsigned Size = A.size(); + + for (unsigned i = 0; i < Size; ++i) { + if (A[i] != i + Inc) + if (++err_cnt < 10) + std::cerr << "failed at A[" << i << "]: " << A[i] << " != " << i + Inc + << "\n"; + } + + if (err_cnt > 0) { + std::cout << " pass rate: " + << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% (" + << (Size - err_cnt) << "/" << Size << ")\n"; + return false; + } + return true; +} + +int main(void) { + constexpr unsigned Size = 32; + constexpr unsigned VL = 16; + + std::vector A(Size); + + for (unsigned i = 0; i < Size; ++i) { + A[i] = i; + } + + try { + buffer bufa(A.data(), range<1>(Size)); + + // We need that many workgroups + ::sycl::range<1> GlobalRange{Size}; + // We need that many threads in each group + ::sycl::range<1> LocalRange{1}; + + queue q; + + auto dev = q.get_device(); + std::cout << "Running on " << dev.get_info() << "\n"; + + auto e = q.submit([&](handler &cgh) { + auto PA = bufa.get_access(cgh); + cgh.parallel_for(GlobalRange * LocalRange, + [=](id<1> i) { PA[i] = PA[i] + 1; }); + }); + e.wait(); + } catch (::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + return 2; + } + + if (checkResult(A, 1)) { + std::cout << "SYCL kernel passed\n"; + } else { + std::cout << "SYCL kernel failed\n"; + return 1; + } + + try { + buffer bufa(A.data(), range<1>(Size)); + + // We need that many workgroups + ::sycl::range<1> GlobalRange{Size / VL}; + // We need that many threads in each group + ::sycl::range<1> LocalRange{1}; + + queue q; + + auto dev = q.get_device(); + std::cout << "Running on " << dev.get_info() << "\n"; + + auto e = q.submit([&](handler &cgh) { + auto PA = bufa.get_access(cgh); + cgh.parallel_for(GlobalRange * LocalRange, + [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::ext::intel::esimd; + unsigned int offset = i * VL * sizeof(float); + simd va; + va.copy_from(PA, offset); + simd vc = va + 1; + vc.copy_to(PA, offset); + }); + }); + e.wait(); + } catch (::sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + return 2; + } + + if (checkResult(A, 2)) { + std::cout << "ESIMD kernel passed\n"; + } else { + std::cout << "ESIMD kernel failed\n"; + return 1; + } + return 0; +}