diff --git a/SYCL/ESIMD/BitonicSortK.cpp b/SYCL/ESIMD/BitonicSortK.cpp index 0bbfe53416..c2d5ab0195 100644 --- a/SYCL/ESIMD/BitonicSortK.cpp +++ b/SYCL/ESIMD/BitonicSortK.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/accessor.cpp b/SYCL/ESIMD/accessor.cpp index 0ee005b834..60d72edfbd 100644 --- a/SYCL/ESIMD/accessor.cpp +++ b/SYCL/ESIMD/accessor.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl -D_CRT_SECURE_NO_WARNINGS=1 %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/accessor_gather_scatter.cpp b/SYCL/ESIMD/accessor_gather_scatter.cpp index 7c1e53be9a..825af28770 100644 --- a/SYCL/ESIMD/accessor_gather_scatter.cpp +++ b/SYCL/ESIMD/accessor_gather_scatter.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated memory intrinsic +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/accessor_load_store.cpp b/SYCL/ESIMD/accessor_load_store.cpp index df218b872e..2895a0080f 100644 --- a/SYCL/ESIMD/accessor_load_store.cpp +++ b/SYCL/ESIMD/accessor_load_store.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated memory intrinsic +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/ballot.cpp b/SYCL/ESIMD/api/ballot.cpp index 470fa36880..e7d89927a8 100755 --- a/SYCL/ESIMD/api/ballot.cpp +++ b/SYCL/ESIMD/api/ballot.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated memory intrinsic +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp b/SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp index 014bf3beec..cab8b8cfb5 100644 --- a/SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp +++ b/SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/api/esimd_bit_ops.cpp b/SYCL/ESIMD/api/esimd_bit_ops.cpp index 4d36e7bb33..62f94f40a4 100644 --- a/SYCL/ESIMD/api/esimd_bit_ops.cpp +++ b/SYCL/ESIMD/api/esimd_bit_ops.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'single_task()' method +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/replicate_smoke.cpp b/SYCL/ESIMD/api/replicate_smoke.cpp index 5511037de8..5b587711ab 100644 --- a/SYCL/ESIMD/api/replicate_smoke.cpp +++ b/SYCL/ESIMD/api/replicate_smoke.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/saturation_smoke.cpp b/SYCL/ESIMD/api/saturation_smoke.cpp index 8459f203bd..d3352e090e 100644 --- a/SYCL/ESIMD/api/saturation_smoke.cpp +++ b/SYCL/ESIMD/api/saturation_smoke.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/simd_any_all.cpp b/SYCL/ESIMD/api/simd_any_all.cpp index 96073f5f2d..1161c18df7 100644 --- a/SYCL/ESIMD/api/simd_any_all.cpp +++ b/SYCL/ESIMD/api/simd_any_all.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/simd_binop_integer_promotion.cpp b/SYCL/ESIMD/api/simd_binop_integer_promotion.cpp index f94917d0af..d4c5825e7e 100644 --- a/SYCL/ESIMD/api/simd_binop_integer_promotion.cpp +++ b/SYCL/ESIMD/api/simd_binop_integer_promotion.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated memory intrinsic +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/simd_copy_to_from.cpp b/SYCL/ESIMD/api/simd_copy_to_from.cpp index ae8d73dbb6..ae41b3da71 100644 --- a/SYCL/ESIMD/api/simd_copy_to_from.cpp +++ b/SYCL/ESIMD/api/simd_copy_to_from.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/api/simd_memory_access.cpp b/SYCL/ESIMD/api/simd_memory_access.cpp index f283e6a082..1b2ffea904 100644 --- a/SYCL/ESIMD/api/simd_memory_access.cpp +++ b/SYCL/ESIMD/api/simd_memory_access.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/simd_negation_operator.cpp b/SYCL/ESIMD/api/simd_negation_operator.cpp index 72c1438b12..65d34b9bb0 100644 --- a/SYCL/ESIMD/api/simd_negation_operator.cpp +++ b/SYCL/ESIMD/api/simd_negation_operator.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/simd_subscript_operator.cpp b/SYCL/ESIMD/api/simd_subscript_operator.cpp index 7251423674..af742e35a2 100644 --- a/SYCL/ESIMD/api/simd_subscript_operator.cpp +++ b/SYCL/ESIMD/api/simd_subscript_operator.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated memory intrinsic +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/simd_view_copy_move_assign.cpp b/SYCL/ESIMD/api/simd_view_copy_move_assign.cpp index d23bf4142a..5204927637 100644 --- a/SYCL/ESIMD/api/simd_view_copy_move_assign.cpp +++ b/SYCL/ESIMD/api/simd_view_copy_move_assign.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/api/simd_view_negation_operator.cpp b/SYCL/ESIMD/api/simd_view_negation_operator.cpp index a8405192e1..48a75318f9 100644 --- a/SYCL/ESIMD/api/simd_view_negation_operator.cpp +++ b/SYCL/ESIMD/api/simd_view_negation_operator.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/simd_view_subscript_operator.cpp b/SYCL/ESIMD/api/simd_view_subscript_operator.cpp index 6de0a5f355..fcbbedd49f 100644 --- a/SYCL/ESIMD/api/simd_view_subscript_operator.cpp +++ b/SYCL/ESIMD/api/simd_view_subscript_operator.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/slm_gather_scatter.cpp b/SYCL/ESIMD/api/slm_gather_scatter.cpp index 42e5321be7..b279daed9f 100644 --- a/SYCL/ESIMD/api/slm_gather_scatter.cpp +++ b/SYCL/ESIMD/api/slm_gather_scatter.cpp @@ -1,5 +1,7 @@ // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp b/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp index 1334664c50..234d7412fe 100644 --- a/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp +++ b/SYCL/ESIMD/api/slm_gather_scatter_heavy.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated memory intrinsic +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/slm_gather_scatter_rgba.cpp b/SYCL/ESIMD/api/slm_gather_scatter_rgba.cpp index 6ec3de3de9..f0baf6f829 100644 --- a/SYCL/ESIMD/api/slm_gather_scatter_rgba.cpp +++ b/SYCL/ESIMD/api/slm_gather_scatter_rgba.cpp @@ -1,5 +1,7 @@ // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/ESIMD/api/unary_ops_heavy.cpp b/SYCL/ESIMD/api/unary_ops_heavy.cpp index d50887423d..e4c911fe20 100644 --- a/SYCL/ESIMD/api/unary_ops_heavy.cpp +++ b/SYCL/ESIMD/api/unary_ops_heavy.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/esimd_check_vc_codegen.cpp b/SYCL/ESIMD/esimd_check_vc_codegen.cpp new file mode 100644 index 0000000000..1c677f3366 --- /dev/null +++ b/SYCL/ESIMD/esimd_check_vc_codegen.cpp @@ -0,0 +1,104 @@ +//==-------- esimd_check_vc_codegen.cpp - DPC++ ESIMD on-device test -------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu +// UNSUPPORTED: cuda || hip +// esimd_emulator does not support online-compiler that invokes 'piProgramBuild' +// UNSUPPORTED: esimd_emulator +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: env SYCL_PI_TRACE=-1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER + +#include "esimd_test_utils.hpp" + +#include +#include +#include + +using namespace cl::sycl; + +int main(void) { + constexpr unsigned Size = 1024 * 128; + constexpr unsigned VL = 16; + + float *A = new float[Size]; + float *B = new float[Size]; + float *C = new float[Size]; + + for (unsigned i = 0; i < Size; ++i) { + A[i] = B[i] = i; + C[i] = 0.0f; + } + + try { + buffer bufa(A, range<1>(Size)); + buffer bufb(B, range<1>(Size)); + buffer bufc(C, range<1>(Size)); + + // We need that many workgroups + range<1> GlobalRange{Size / VL}; + + // We need that many threads in each group + range<1> LocalRange{1}; + + queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()); + + 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); + auto PB = bufb.get_access(cgh); + auto PC = bufc.get_access(cgh); + cgh.parallel_for( + GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL { + using namespace sycl::ext::intel::experimental::esimd; + unsigned int offset = i * VL * sizeof(float); + simd va; + va.copy_from(PA, offset); + simd vb; + vb.copy_from(PB, offset); + simd vc = va + vb; + vc.copy_to(PC, offset); + }); + }); + e.wait(); + } catch (sycl::exception const &e) { + std::cout << "SYCL exception caught: " << e.what() << '\n'; + + delete[] A; + delete[] B; + delete[] C; + return 1; + } + + int err_cnt = 0; + + for (unsigned i = 0; i < Size; ++i) { + if (A[i] + B[i] != C[i]) { + if (++err_cnt < 10) { + std::cout << "failed at index " << i << ", " << C[i] << " != " << A[i] + << " + " << B[i] << "\n"; + } + } + } + if (err_cnt > 0) { + std::cout << " pass rate: " + << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% (" + << (Size - err_cnt) << "/" << Size << ")\n"; + } + + delete[] A; + delete[] B; + delete[] C; + + std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n"); + return err_cnt > 0 ? 1 : 0; +} + +// CHECK: ---> piProgramBuild( +// CHECK: : {{.*}}-vc-codegen +// CHECK: ) ---> pi_result : PI_SUCCESS diff --git a/SYCL/ESIMD/ext_math.cpp b/SYCL/ESIMD/ext_math.cpp index c7c40c250e..3ffe931e6d 100644 --- a/SYCL/ESIMD/ext_math.cpp +++ b/SYCL/ESIMD/ext_math.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/fp_in_phi.cpp b/SYCL/ESIMD/fp_in_phi.cpp index 839d1139d7..3390b1ecab 100644 --- a/SYCL/ESIMD/fp_in_phi.cpp +++ b/SYCL/ESIMD/fp_in_phi.cpp @@ -14,6 +14,8 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip // UNSUPPORTED: ze_debug-1,ze_debug4 +// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled +// XFAIL: esimd_emulator // // The test checks that ESIMD kernels correctly handle function pointers as // arguments of LLVM's PHI function. diff --git a/SYCL/ESIMD/histogram.cpp b/SYCL/ESIMD/histogram.cpp index 9b062fbe80..30eb2c51a7 100644 --- a/SYCL/ESIMD/histogram.cpp +++ b/SYCL/ESIMD/histogram.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated __esimd_media_ld +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/histogram_256_slm.cpp b/SYCL/ESIMD/histogram_256_slm.cpp index 5a072c2f4f..844df448f4 100644 --- a/SYCL/ESIMD/histogram_256_slm.cpp +++ b/SYCL/ESIMD/histogram_256_slm.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/histogram_256_slm_spec_2020.cpp b/SYCL/ESIMD/histogram_256_slm_spec_2020.cpp index 3b33141d69..cc0ca87778 100644 --- a/SYCL/ESIMD/histogram_256_slm_spec_2020.cpp +++ b/SYCL/ESIMD/histogram_256_slm_spec_2020.cpp @@ -1,6 +1,8 @@ // TODO enable on Windows // REQUIRES: linux && gpu // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out 16 diff --git a/SYCL/ESIMD/histogram_2d.cpp b/SYCL/ESIMD/histogram_2d.cpp index 4fd0cd063a..b492f7df13 100644 --- a/SYCL/ESIMD/histogram_2d.cpp +++ b/SYCL/ESIMD/histogram_2d.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated __esimd_media_ld +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/histogram_raw_send.cpp b/SYCL/ESIMD/histogram_raw_send.cpp index 2389179109..9cf141841f 100644 --- a/SYCL/ESIMD/histogram_raw_send.cpp +++ b/SYCL/ESIMD/histogram_raw_send.cpp @@ -9,6 +9,8 @@ // REQUIRES: gpu // UNSUPPORTED: gpu-intel-dg1,cuda,hip // UNSUPPORTED: ze_debug-1,ze_debug4 +// TODO: esimd_emulator fails due to outdated __esimd_media_ld +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/linear/linear.cpp b/SYCL/ESIMD/linear/linear.cpp index 6b96a9970f..b975dee272 100644 --- a/SYCL/ESIMD/linear/linear.cpp +++ b/SYCL/ESIMD/linear/linear.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated __esimd_media_ld +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -I%S/.. -o %t.out // RUN: %HOST_RUN_PLACEHOLDER %t.out %S/linear_in.bmp %S/linear_gold_hw.bmp // RUN: %GPU_RUN_PLACEHOLDER %t.out %S/linear_in.bmp %S/linear_gold_hw.bmp diff --git a/SYCL/ESIMD/mandelbrot/mandelbrot.cpp b/SYCL/ESIMD/mandelbrot/mandelbrot.cpp index 119421fb6c..5eb339f47d 100644 --- a/SYCL/ESIMD/mandelbrot/mandelbrot.cpp +++ b/SYCL/ESIMD/mandelbrot/mandelbrot.cpp @@ -8,6 +8,8 @@ // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated __esimd_media_st +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -I%S/.. -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out %T/output.ppm %S/golden_hw.ppm diff --git a/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp b/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp index c668bd8274..c56186da24 100644 --- a/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp +++ b/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp @@ -9,6 +9,8 @@ // TODO enable on Windows // REQUIRES: linux && gpu // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -I%S/.. -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out %T/output_spec.ppm %S/golden_hw.ppm 512 -2.09798 -1.19798 0.004 4.0 diff --git a/SYCL/ESIMD/matrix_transpose.cpp b/SYCL/ESIMD/matrix_transpose.cpp index 889d13146c..e42aba9ec7 100644 --- a/SYCL/ESIMD/matrix_transpose.cpp +++ b/SYCL/ESIMD/matrix_transpose.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/matrix_transpose2.cpp b/SYCL/ESIMD/matrix_transpose2.cpp index ae642c3d2f..f7a4560d11 100644 --- a/SYCL/ESIMD/matrix_transpose2.cpp +++ b/SYCL/ESIMD/matrix_transpose2.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated __esimd_media_ld +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/matrix_transpose_glb.cpp b/SYCL/ESIMD/matrix_transpose_glb.cpp index 708098605f..84b681ee5e 100644 --- a/SYCL/ESIMD/matrix_transpose_glb.cpp +++ b/SYCL/ESIMD/matrix_transpose_glb.cpp @@ -20,6 +20,8 @@ using namespace cl::sycl; using namespace std; using namespace sycl::ext::intel::experimental::esimd; +const unsigned int ESIMD_EMULATOR_SIZE_LIMIT = 1U << 10; + void initMatrix(int *M, unsigned N) { assert(N >= 8 && (((N - 1) & N) == 0) && "only power of 2 (>= 16) is supported"); @@ -246,6 +248,16 @@ bool runTest(unsigned MZ, unsigned block_size, unsigned num_iters, cerr << "\nTranspose square matrix of size " << MZ << "\n"; // printMatrix("Initial matrix:", M, MZ); + if ((q.get_backend() == cl::sycl::backend::ext_intel_esimd_emulator) && + (MZ > ESIMD_EMULATOR_SIZE_LIMIT)) { + cerr << "Matrix Size larger than " << ESIMD_EMULATOR_SIZE_LIMIT + << " is skipped" + << "\n"; + cerr << "for esimd_emulator backend due to timeout" + << "\n"; + return true; + } + // Each C-for-Metal thread works on one or two blocks of size 8 x 8. int thread_width = MZ / block_size; int thread_height = MZ / block_size; diff --git a/SYCL/ESIMD/printf.cpp b/SYCL/ESIMD/printf.cpp index cb7d6e1cf0..de45c7ce77 100644 --- a/SYCL/ESIMD/printf.cpp +++ b/SYCL/ESIMD/printf.cpp @@ -9,6 +9,8 @@ // REQUIRES: gpu // UNSUPPORTED: cuda || hip // CUDA and HIP don't support printf. +// TODO: esimd_emulator fails due to unimplemented 'single_task()' method +// XFAIL: esimd_emulator // // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER diff --git a/SYCL/ESIMD/regression/big_const_initializer.cpp b/SYCL/ESIMD/regression/big_const_initializer.cpp index a4f1300398..f0e800ffdc 100644 --- a/SYCL/ESIMD/regression/big_const_initializer.cpp +++ b/SYCL/ESIMD/regression/big_const_initializer.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/regression/complex-lib-lin.cpp b/SYCL/ESIMD/regression/complex-lib-lin.cpp index 15e8013ba9..9c4a97b931 100644 --- a/SYCL/ESIMD/regression/complex-lib-lin.cpp +++ b/SYCL/ESIMD/regression/complex-lib-lin.cpp @@ -7,6 +7,8 @@ // // REQUIRES: linux,gpu // UNSUPPORTED: cuda || hip +// TODO/DEBUG Segmentation fault occurs with esimd_emulator backend +// XFAIL: esimd_emulator // // RUN: %clangxx -fsycl -fPIC -O3 %S/Inputs/complex-lib-sycl.cpp -c -o %t-lib-sycl.o // RUN: %clangxx -fsycl -fPIC -O3 %S/Inputs/complex-lib-esimd.cpp -c -o %t-lib-esimd.o diff --git a/SYCL/ESIMD/regression/globals.cpp b/SYCL/ESIMD/regression/globals.cpp index 2345027f8f..8a21764695 100644 --- a/SYCL/ESIMD/regression/globals.cpp +++ b/SYCL/ESIMD/regression/globals.cpp @@ -1,5 +1,7 @@ // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented sub-group support +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/regression/operator_decrement.cpp b/SYCL/ESIMD/regression/operator_decrement.cpp index 379af94a43..f9d96b6368 100644 --- a/SYCL/ESIMD/regression/operator_decrement.cpp +++ b/SYCL/ESIMD/regression/operator_decrement.cpp @@ -9,6 +9,8 @@ // // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/regression/unused_load.cpp b/SYCL/ESIMD/regression/unused_load.cpp index 3d9ae61e46..d4a626a6c1 100644 --- a/SYCL/ESIMD/regression/unused_load.cpp +++ b/SYCL/ESIMD/regression/unused_load.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/slm_barrier.cpp b/SYCL/ESIMD/slm_barrier.cpp index 2e58afc433..2276267204 100644 --- a/SYCL/ESIMD/slm_barrier.cpp +++ b/SYCL/ESIMD/slm_barrier.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_scatter_scaled4 +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/slm_split_barrier.cpp b/SYCL/ESIMD/slm_split_barrier.cpp index 9b180bca68..bb708f39d5 100644 --- a/SYCL/ESIMD/slm_split_barrier.cpp +++ b/SYCL/ESIMD/slm_split_barrier.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_scatter4_scaled +// XFAIL: esimd_emulator #include "esimd_test_utils.hpp" diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp index 7476d496aa..dec592f5cb 100644 --- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp index 3d7b6137c7..a04db2a283 100644 --- a/SYCL/ESIMD/spec_const/spec_const_char.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp index 146a170f37..658164f702 100644 --- a/SYCL/ESIMD/spec_const/spec_const_double.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp index 04e0209862..c53bbffff1 100644 --- a/SYCL/ESIMD/spec_const/spec_const_float.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp index 13b3b4227f..2458898c33 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_int64.cpp b/SYCL/ESIMD/spec_const/spec_const_int64.cpp index f67085eaf7..fb29bb3085 100644 --- a/SYCL/ESIMD/spec_const/spec_const_int64.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_int64.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp index 3c296dab27..754c0058de 100644 --- a/SYCL/ESIMD/spec_const/spec_const_short.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp index cd934e2ba0..9db9f82dc4 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp index a0f93df45b..32a26ed02f 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_uint64.cpp b/SYCL/ESIMD/spec_const/spec_const_uint64.cpp index aa6a2814bf..31fdeb66ea 100644 --- a/SYCL/ESIMD/spec_const/spec_const_uint64.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_uint64.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp index f0558d7b74..7349f5e67a 100644 --- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp +++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp @@ -9,6 +9,8 @@ // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO online_compiler check fails for esimd_emulator +// XFAIL: esimd_emulator #include diff --git a/SYCL/ESIMD/sycl_esimd_mix.cpp b/SYCL/ESIMD/sycl_esimd_mix.cpp index b3052f374f..4676bac97f 100644 --- a/SYCL/ESIMD/sycl_esimd_mix.cpp +++ b/SYCL/ESIMD/sycl_esimd_mix.cpp @@ -10,6 +10,9 @@ // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO/FIXME: esimd_emulator does not support online compilation that +// invokes 'piProgramBuild'/'piKernelCreate' +// UNSUPPORTED: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: env SYCL_PI_TRACE=-1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK,CHECK-NO-VAR // RUN: env SYCL_PROGRAM_COMPILE_OPTIONS="-g" SYCL_PI_TRACE=-1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK,CHECK-WITH-VAR diff --git a/SYCL/ESIMD/vadd_1d.cpp b/SYCL/ESIMD/vadd_1d.cpp index 78bf4e04a5..c3cfa236db 100644 --- a/SYCL/ESIMD/vadd_1d.cpp +++ b/SYCL/ESIMD/vadd_1d.cpp @@ -7,8 +7,10 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out -// RUN: env SYCL_PI_TRACE=-1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER +// RUN: %GPU_RUN_PLACEHOLDER %t.out #include "esimd_test_utils.hpp" @@ -96,7 +98,3 @@ int main(void) { std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n"); return err_cnt > 0 ? 1 : 0; } - -// CHECK: ---> piProgramBuild( -// CHECK: : {{.*}}-vc-codegen -// CHECK: ) ---> pi_result : PI_SUCCESS diff --git a/SYCL/ESIMD/vadd_2d.cpp b/SYCL/ESIMD/vadd_2d.cpp index 3037281833..8edfd32da8 100644 --- a/SYCL/ESIMD/vadd_2d.cpp +++ b/SYCL/ESIMD/vadd_2d.cpp @@ -7,6 +7,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to outdated __esimd_media_ld +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/vadd_2d_acc.cpp b/SYCL/ESIMD/vadd_2d_acc.cpp index 7533573e7f..bdb3313d23 100644 --- a/SYCL/ESIMD/vadd_2d_acc.cpp +++ b/SYCL/ESIMD/vadd_2d_acc.cpp @@ -10,6 +10,8 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented __esimd_oword_ld_unaligned +// XFAIL: esimd_emulator // The test checks that 2D workitem addressing works correctly with SIMD // kernels. diff --git a/SYCL/ESIMD/vadd_half.cpp b/SYCL/ESIMD/vadd_half.cpp index b92df7f748..2d5a2612a2 100644 --- a/SYCL/ESIMD/vadd_half.cpp +++ b/SYCL/ESIMD/vadd_half.cpp @@ -8,6 +8,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: cuda || hip +// TODO: esimd_emulator fails due to unimplemented 'half' type +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out diff --git a/SYCL/ESIMD/vadd_raw_send.cpp b/SYCL/ESIMD/vadd_raw_send.cpp index f74bc29a7e..1ca8ca2505 100644 --- a/SYCL/ESIMD/vadd_raw_send.cpp +++ b/SYCL/ESIMD/vadd_raw_send.cpp @@ -8,6 +8,8 @@ //===----------------------------------------------------------------------===// // REQUIRES: gpu // UNSUPPORTED: gpu-intel-dg1,cuda,hip +// TODO: esimd_emulator fails due to unimplemented 'raw_send' intrinsic +// XFAIL: esimd_emulator // RUN: %clangxx -fsycl %s -o %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out