diff --git a/offload/test/api/omp_get_mapped_ptr.c b/offload/test/api/omp_get_mapped_ptr.c index 4436c02259b6..a8e11f912d66 100644 --- a/offload/test/api/omp_get_mapped_ptr.c +++ b/offload/test/api/omp_get_mapped_ptr.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-and-run-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/api/omp_get_num_procs.c b/offload/test/api/omp_get_num_procs.c index 24222d32983b..e958d0cf3492 100644 --- a/offload/test/api/omp_get_num_procs.c +++ b/offload/test/api/omp_get_num_procs.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include diff --git a/offload/test/mapping/array_section_implicit_capture.c b/offload/test/mapping/array_section_implicit_capture.c index fae18112763e..1042bd23bd5e 100644 --- a/offload/test/mapping/array_section_implicit_capture.c +++ b/offload/test/mapping/array_section_implicit_capture.c @@ -1,7 +1,6 @@ // RUN: %libomptarget-compile-generic // RUN: %libomptarget-run-generic 2>&1 \ // RUN: | %fcheck-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/chained_containing_structs_1.cc b/offload/test/mapping/chained_containing_structs_1.cc index e98d7fe1f975..c9a09bb657be 100644 --- a/offload/test/mapping/chained_containing_structs_1.cc +++ b/offload/test/mapping/chained_containing_structs_1.cc @@ -1,5 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/chained_containing_structs_2.cc b/offload/test/mapping/chained_containing_structs_2.cc index ed959dd6c495..272b0ebaa7c9 100644 --- a/offload/test/mapping/chained_containing_structs_2.cc +++ b/offload/test/mapping/chained_containing_structs_2.cc @@ -1,5 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/chained_containing_structs_3.cc b/offload/test/mapping/chained_containing_structs_3.cc index 95692e67a83b..763c9760918e 100644 --- a/offload/test/mapping/chained_containing_structs_3.cc +++ b/offload/test/mapping/chained_containing_structs_3.cc @@ -1,5 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/data_member_ref.cpp b/offload/test/mapping/data_member_ref.cpp index 2b5f978b6040..7947a62c169f 100644 --- a/offload/test/mapping/data_member_ref.cpp +++ b/offload/test/mapping/data_member_ref.cpp @@ -1,5 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu #include diff --git a/offload/test/mapping/declare_mapper_nested_default_mappers.cpp b/offload/test/mapping/declare_mapper_nested_default_mappers.cpp index 2d69f9573ad2..45fd042aedb0 100644 --- a/offload/test/mapping/declare_mapper_nested_default_mappers.cpp +++ b/offload/test/mapping/declare_mapper_nested_default_mappers.cpp @@ -1,5 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/declare_mapper_nested_mappers.cpp b/offload/test/mapping/declare_mapper_nested_mappers.cpp index 5d18739bf22e..a59ed6980ec4 100644 --- a/offload/test/mapping/declare_mapper_nested_mappers.cpp +++ b/offload/test/mapping/declare_mapper_nested_mappers.cpp @@ -1,5 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/declare_mapper_target_checks.cpp b/offload/test/mapping/declare_mapper_target_checks.cpp index 562e283d81bc..019ad5dec61d 100644 --- a/offload/test/mapping/declare_mapper_target_checks.cpp +++ b/offload/test/mapping/declare_mapper_target_checks.cpp @@ -1,5 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/duplicate_mappings_1.cpp b/offload/test/mapping/duplicate_mappings_1.cpp index 731bff9035d5..5bed3f161b4f 100644 --- a/offload/test/mapping/duplicate_mappings_1.cpp +++ b/offload/test/mapping/duplicate_mappings_1.cpp @@ -1,6 +1,5 @@ // clang-format off // RUN: %libomptarget-compilexx-generic -Wno-openmp-mapping && %libomptarget-run-generic -// XFAIL: intelgpu // clang-format on diff --git a/offload/test/mapping/duplicate_mappings_2.cpp b/offload/test/mapping/duplicate_mappings_2.cpp index 6ca7d3b55dfb..ca8112ee7254 100644 --- a/offload/test/mapping/duplicate_mappings_2.cpp +++ b/offload/test/mapping/duplicate_mappings_2.cpp @@ -1,6 +1,5 @@ // clang-format off // RUN: %libomptarget-compilexx-generic -Wno-openmp-mapping && %libomptarget-run-generic -// XFAIL: intelgpu #include diff --git a/offload/test/mapping/implicit_device_ptr.c b/offload/test/mapping/implicit_device_ptr.c index 6e32ece82a1c..baa75d21686a 100644 --- a/offload/test/mapping/implicit_device_ptr.c +++ b/offload/test/mapping/implicit_device_ptr.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/lambda_by_value.cpp b/offload/test/mapping/lambda_by_value.cpp index b4e3f16c295a..4c0278d40592 100644 --- a/offload/test/mapping/lambda_by_value.cpp +++ b/offload/test/mapping/lambda_by_value.cpp @@ -1,6 +1,5 @@ // RUN: %libomptarget-compileopt-generic -fno-exceptions // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/low_alignment.c b/offload/test/mapping/low_alignment.c index 71e90b3695a3..615a5a9c3111 100644 --- a/offload/test/mapping/low_alignment.c +++ b/offload/test/mapping/low_alignment.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/map_ordering_tgt_alloc_tofrom.c b/offload/test/mapping/map_ordering_tgt_alloc_tofrom.c index af7463066a80..806771e84916 100644 --- a/offload/test/mapping/map_ordering_tgt_alloc_tofrom.c +++ b/offload/test/mapping/map_ordering_tgt_alloc_tofrom.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include diff --git a/offload/test/mapping/map_ordering_tgt_data_alloc_from.c b/offload/test/mapping/map_ordering_tgt_data_alloc_from.c index eaeb38df9c65..b793551ee6ee 100644 --- a/offload/test/mapping/map_ordering_tgt_data_alloc_from.c +++ b/offload/test/mapping/map_ordering_tgt_data_alloc_from.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include diff --git a/offload/test/mapping/map_ordering_tgt_data_alloc_to_from.c b/offload/test/mapping/map_ordering_tgt_data_alloc_to_from.c index 2e501ecaa8bd..c636cccdb9fa 100644 --- a/offload/test/mapping/map_ordering_tgt_data_alloc_to_from.c +++ b/offload/test/mapping/map_ordering_tgt_data_alloc_to_from.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include diff --git a/offload/test/mapping/map_ordering_tgt_data_alloc_tofrom.c b/offload/test/mapping/map_ordering_tgt_data_alloc_tofrom.c index 883a59d0d408..0e44fc1531bd 100644 --- a/offload/test/mapping/map_ordering_tgt_data_alloc_tofrom.c +++ b/offload/test/mapping/map_ordering_tgt_data_alloc_tofrom.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c index d685d99c350a..563cdfa3c3a6 100644 --- a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c +++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include diff --git a/offload/test/mapping/ompx_hold/struct.c b/offload/test/mapping/ompx_hold/struct.c index 0b898d5d3b56..da2b38028762 100644 --- a/offload/test/mapping/ompx_hold/struct.c +++ b/offload/test/mapping/ompx_hold/struct.c @@ -1,6 +1,5 @@ // RUN: %libomptarget-compile-generic -fopenmp-extensions // RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/pr38704.c b/offload/test/mapping/pr38704.c index 612c2683b37c..72b4a171ccea 100644 --- a/offload/test/mapping/pr38704.c +++ b/offload/test/mapping/pr38704.c @@ -3,7 +3,6 @@ // Clang 6.0 doesn't use the new map interface, undefined behavior when // the compiler emits "old" interface code for structures. // UNSUPPORTED: clang-6 -// XFAIL: intelgpu #include #include diff --git a/offload/test/mapping/ptr_and_obj_motion.c b/offload/test/mapping/ptr_and_obj_motion.c index 6a5f16eb8f30..a94c07aadc1b 100644 --- a/offload/test/mapping/ptr_and_obj_motion.c +++ b/offload/test/mapping/ptr_and_obj_motion.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include diff --git a/offload/test/mapping/target_data_array_extension_at_exit.c b/offload/test/mapping/target_data_array_extension_at_exit.c index cfb09a98e06f..e300c800b255 100644 --- a/offload/test/mapping/target_data_array_extension_at_exit.c +++ b/offload/test/mapping/target_data_array_extension_at_exit.c @@ -15,7 +15,6 @@ // RUN: -DEXTENDS=AFTER // RUN: %libomptarget-run-generic 2>&1 \ // RUN: | %fcheck-generic -// XFAIL: intelgpu // END. diff --git a/offload/test/mapping/target_map_for_member_data.cpp b/offload/test/mapping/target_map_for_member_data.cpp index 9f2adc29d667..fafff6b18a2b 100644 --- a/offload/test/mapping/target_map_for_member_data.cpp +++ b/offload/test/mapping/target_map_for_member_data.cpp @@ -1,7 +1,6 @@ // RUN: %libomptarget-compile-generic -fopenmp-version=51 // RUN: %libomptarget-run-generic 2>&1 \ // RUN: | %fcheck-generic -// XFAIL: intelgpu extern "C" int printf(const char *, ...); template class A { diff --git a/offload/test/mapping/use_device_addr/target_use_device_addr.c b/offload/test/mapping/use_device_addr/target_use_device_addr.c index 00a323f0499d..4a9dbe252f76 100644 --- a/offload/test/mapping/use_device_addr/target_use_device_addr.c +++ b/offload/test/mapping/use_device_addr/target_use_device_addr.c @@ -1,7 +1,6 @@ // RUN: %libomptarget-compile-generic -fopenmp-version=51 // RUN: %libomptarget-run-generic 2>&1 \ // RUN: | %fcheck-generic -// XFAIL: intelgpu #include int main() { diff --git a/offload/test/offloading/atomic-compare-signedness.c b/offload/test/offloading/atomic-compare-signedness.c index e1c8a2f846a8..7ce981f46383 100644 --- a/offload/test/offloading/atomic-compare-signedness.c +++ b/offload/test/offloading/atomic-compare-signedness.c @@ -7,7 +7,8 @@ // RUN: %libomptarget-run-generic | %fcheck-generic // RUN: %libomptarget-compileopt-generic -fopenmp-version=51 // RUN: %libomptarget-run-generic | %fcheck-generic -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu // High parallelism increases our chances of detecting a lack of atomicity. #define NUM_THREADS_TRY 256 diff --git a/offload/test/offloading/back2back_distribute.c b/offload/test/offloading/back2back_distribute.c index 4bd3d52a8c2f..d5ab0500fe80 100644 --- a/offload/test/offloading/back2back_distribute.c +++ b/offload/test/offloading/back2back_distribute.c @@ -1,6 +1,7 @@ // clang-format off // RUN: %libomptarget-compile-generic -O3 && %libomptarget-run-generic | %fcheck-generic -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu // clang-format on #include #include diff --git a/offload/test/offloading/bug47654.cpp b/offload/test/offloading/bug47654.cpp index 3203e9f6da48..eba7904e76ca 100644 --- a/offload/test/offloading/bug47654.cpp +++ b/offload/test/offloading/bug47654.cpp @@ -1,6 +1,5 @@ // RUN: %libomptarget-compilexx-run-and-check-generic // RUN: %libomptarget-compileoptxx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/bug49021.cpp b/offload/test/offloading/bug49021.cpp index 1bd35c66f5db..83fc66317e9b 100644 --- a/offload/test/offloading/bug49021.cpp +++ b/offload/test/offloading/bug49021.cpp @@ -3,7 +3,8 @@ // RUN: %libomptarget-compilexx-generic -O3 -ffast-math && %libomptarget-run-generic // RUN: %libomptarget-compileoptxx-generic -O3 && %libomptarget-run-generic // RUN: %libomptarget-compileoptxx-generic -O3 -ffast-math && %libomptarget-run-generic -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu // clang-format on #include diff --git a/offload/test/offloading/bug50022.cpp b/offload/test/offloading/bug50022.cpp index 1bd48be9f9d8..d2ae02b7054b 100644 --- a/offload/test/offloading/bug50022.cpp +++ b/offload/test/offloading/bug50022.cpp @@ -1,6 +1,5 @@ // RUN: %libomptarget-compilexx-and-run-generic // RUN: %libomptarget-compileoptxx-and-run-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/bug51781.c b/offload/test/offloading/bug51781.c index 3c7f2ce09022..b5841146788f 100644 --- a/offload/test/offloading/bug51781.c +++ b/offload/test/offloading/bug51781.c @@ -33,7 +33,8 @@ // RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=CUSTOM -input-file=%t.custom // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu // // CUSTOM: Rewriting generic-mode kernel with a customized state machine. diff --git a/offload/test/offloading/bug51982.c b/offload/test/offloading/bug51982.c index a6910b1fbb3f..ca32ff29437c 100644 --- a/offload/test/offloading/bug51982.c +++ b/offload/test/offloading/bug51982.c @@ -1,7 +1,8 @@ // RUN: %libomptarget-compile-generic -O2 && %libomptarget-run-generic // -O2 to run openmp-opt // RUN: %libomptarget-compileopt-generic -O2 && %libomptarget-run-generic -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu int main(void) { long int aa = 0; diff --git a/offload/test/offloading/bug53727.cpp b/offload/test/offloading/bug53727.cpp index a1a01f007850..8ce8b7e9b87c 100644 --- a/offload/test/offloading/bug53727.cpp +++ b/offload/test/offloading/bug53727.cpp @@ -1,6 +1,5 @@ // RUN: %libomptarget-compilexx-and-run-generic // RUN: %libomptarget-compileoptxx-and-run-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/bug64959.c b/offload/test/offloading/bug64959.c index eddc55325ffe..a31b8449d34f 100644 --- a/offload/test/offloading/bug64959.c +++ b/offload/test/offloading/bug64959.c @@ -6,6 +6,7 @@ // UNSUPPORTED: amdgcn-amd-amdhsa // UNSUPPORTED: nvptx64-nvidia-cuda // UNSUPPORTED: nvptx64-nvidia-cuda-LTO +// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/extern.c b/offload/test/offloading/extern.c index f594e377d8f2..4828855ee9c4 100644 --- a/offload/test/offloading/extern.c +++ b/offload/test/offloading/extern.c @@ -1,7 +1,6 @@ // clang-format off // RUN: %libomptarget-compile-generic -DVAR -c -o %t.o // RUN: %libomptarget-compile-generic %t.o && %libomptarget-run-generic | %fcheck-generic -// XFAIL: intelgpu // clang-format on #ifdef VAR int x = 1; diff --git a/offload/test/offloading/generic_multiple_parallel_regions.c b/offload/test/offloading/generic_multiple_parallel_regions.c index 80939d147ebc..cb7ea2d81b57 100644 --- a/offload/test/offloading/generic_multiple_parallel_regions.c +++ b/offload/test/offloading/generic_multiple_parallel_regions.c @@ -1,6 +1,7 @@ // RUN: %libomptarget-compile-run-and-check-generic // RUN: %libomptarget-compileopt-run-and-check-generic -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu #include #include diff --git a/offload/test/offloading/global_constructor.cpp b/offload/test/offloading/global_constructor.cpp index 0f869788e146..9673e2b6b622 100644 --- a/offload/test/offloading/global_constructor.cpp +++ b/offload/test/offloading/global_constructor.cpp @@ -1,6 +1,5 @@ // clang-format off // RUN: %libomptarget-compilexx-generic && %libomptarget-run-generic | %fcheck-generic -// XFAIL: intelgpu // clang-format on #include diff --git a/offload/test/offloading/host_as_target.c b/offload/test/offloading/host_as_target.c index 24e870b441ed..8ab991b72338 100644 --- a/offload/test/offloading/host_as_target.c +++ b/offload/test/offloading/host_as_target.c @@ -6,7 +6,6 @@ // - Works whether it's specified directly or as the default device. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/spmdization.c b/offload/test/offloading/spmdization.c index ca881a758024..ea76d9c1a31c 100644 --- a/offload/test/offloading/spmdization.c +++ b/offload/test/offloading/spmdization.c @@ -9,7 +9,8 @@ // clang-format on // REQUIRES: gpu -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu #include #include diff --git a/offload/test/offloading/static_linking.c b/offload/test/offloading/static_linking.c index 000ed11052d3..4487ab83dd5f 100644 --- a/offload/test/offloading/static_linking.c +++ b/offload/test/offloading/static_linking.c @@ -2,7 +2,6 @@ // RUN: %libomptarget-compile-generic -DLIBRARY -c -o %t.o // RUN: ar rcs %t.a %t.o // RUN: %libomptarget-compile-generic %t.a && %libomptarget-run-generic 2>&1 | %fcheck-generic -// XFAIL: intelgpu // clang-format on #ifdef LIBRARY diff --git a/offload/test/offloading/strided_multiple_update_from.c b/offload/test/offloading/strided_multiple_update_from.c index 4f2df8193373..a3e8d10863ae 100644 --- a/offload/test/offloading/strided_multiple_update_from.c +++ b/offload/test/offloading/strided_multiple_update_from.c @@ -3,7 +3,6 @@ // from the device to the host. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/strided_multiple_update_to.c b/offload/test/offloading/strided_multiple_update_to.c index bb16d7a09269..02b6210807d2 100644 --- a/offload/test/offloading/strided_multiple_update_to.c +++ b/offload/test/offloading/strided_multiple_update_to.c @@ -3,7 +3,6 @@ // from the host to the device. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/strided_partial_update_from.c b/offload/test/offloading/strided_partial_update_from.c index 4a2977a08f70..15d477f2b9b7 100644 --- a/offload/test/offloading/strided_partial_update_from.c +++ b/offload/test/offloading/strided_partial_update_from.c @@ -3,7 +3,6 @@ // across the array // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/strided_partial_update_to.c b/offload/test/offloading/strided_partial_update_to.c index f9c960fb9c96..1cbdc0917c31 100644 --- a/offload/test/offloading/strided_partial_update_to.c +++ b/offload/test/offloading/strided_partial_update_to.c @@ -3,7 +3,6 @@ // across the array // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/strided_ptr_multiple_update_from.c b/offload/test/offloading/strided_ptr_multiple_update_from.c index 33deb59534f6..b671a1e2b688 100644 --- a/offload/test/offloading/strided_ptr_multiple_update_from.c +++ b/offload/test/offloading/strided_ptr_multiple_update_from.c @@ -3,7 +3,6 @@ // from the device to the host using dynamically allocated memory. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include #include diff --git a/offload/test/offloading/strided_ptr_multiple_update_to.c b/offload/test/offloading/strided_ptr_multiple_update_to.c index be56f3dc5c24..7ae9189ead8e 100644 --- a/offload/test/offloading/strided_ptr_multiple_update_to.c +++ b/offload/test/offloading/strided_ptr_multiple_update_to.c @@ -3,7 +3,6 @@ // from the host to the device using dynamically allocated memory. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include #include diff --git a/offload/test/offloading/strided_ptr_partial_update_from.c b/offload/test/offloading/strided_ptr_partial_update_from.c index be235d3d9c53..d30a67139300 100644 --- a/offload/test/offloading/strided_ptr_partial_update_from.c +++ b/offload/test/offloading/strided_ptr_partial_update_from.c @@ -3,7 +3,6 @@ // across the array using dynamically allocated memory. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include #include diff --git a/offload/test/offloading/strided_ptr_partial_update_to.c b/offload/test/offloading/strided_ptr_partial_update_to.c index 662e1a2655ca..37c165351a76 100644 --- a/offload/test/offloading/strided_ptr_partial_update_to.c +++ b/offload/test/offloading/strided_ptr_partial_update_to.c @@ -3,7 +3,6 @@ // across the array using dynamically allocated memory. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include #include diff --git a/offload/test/offloading/strided_update_count_expression.c b/offload/test/offloading/strided_update_count_expression.c index 40e8b89de259..e18fc881a802 100644 --- a/offload/test/offloading/strided_update_count_expression.c +++ b/offload/test/offloading/strided_update_count_expression.c @@ -4,7 +4,6 @@ // where the count (len/2) is a variable expression, not a constant. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/strided_update_from.c b/offload/test/offloading/strided_update_from.c index 9910bed9c546..fe875b7fd55c 100644 --- a/offload/test/offloading/strided_update_from.c +++ b/offload/test/offloading/strided_update_from.c @@ -4,7 +4,6 @@ // other element (stride 2) from the device to the host // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/strided_update_multiple_arrays_count_expression.c b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c index fdc7653f36d9..a1472cacc4a3 100644 --- a/offload/test/offloading/strided_update_multiple_arrays_count_expression.c +++ b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c @@ -4,7 +4,6 @@ // same array with various count expressions. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c index 315bfef950ef..32859202a20f 100644 --- a/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c +++ b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // Tests multiple arrays with different variable strides in single update // clause. diff --git a/offload/test/offloading/strided_update_to.c b/offload/test/offloading/strided_update_to.c index eca20abed4b8..b90a240d1abe 100644 --- a/offload/test/offloading/strided_update_to.c +++ b/offload/test/offloading/strided_update_to.c @@ -4,7 +4,6 @@ // other element (stride 2) from the host to the device // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/strided_update_variable_count_and_stride.c b/offload/test/offloading/strided_update_variable_count_and_stride.c index 98f0b5b51e32..1e1e41653c2c 100644 --- a/offload/test/offloading/strided_update_variable_count_and_stride.c +++ b/offload/test/offloading/strided_update_variable_count_and_stride.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // Tests combining variable count expression AND variable stride in array // sections. diff --git a/offload/test/offloading/strided_update_variable_stride.c b/offload/test/offloading/strided_update_variable_stride.c index 31fff587eed0..7c8079efa56e 100644 --- a/offload/test/offloading/strided_update_variable_stride.c +++ b/offload/test/offloading/strided_update_variable_stride.c @@ -2,7 +2,6 @@ // Tests data[0:5:stride] where stride is a variable, making it non-contiguous. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/strided_update_variable_stride_misc.c b/offload/test/offloading/strided_update_variable_stride_misc.c index d5579fe48887..d7f5fe8dcc12 100644 --- a/offload/test/offloading/strided_update_variable_stride_misc.c +++ b/offload/test/offloading/strided_update_variable_stride_misc.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // Miscellaneous variable stride tests: stride=1, stride=array_size, stride from // array subscript. diff --git a/offload/test/offloading/target-tile.c b/offload/test/offloading/target-tile.c index e5413a9493f4..8460b43b6f9c 100644 --- a/offload/test/offloading/target-tile.c +++ b/offload/test/offloading/target-tile.c @@ -3,7 +3,6 @@ // RUN: %libomptarget-compile-generic -fopenmp-version=51 // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic -// XFAIL: intelgpu #include diff --git a/offload/test/offloading/target_constexpr_mapping.cpp b/offload/test/offloading/target_constexpr_mapping.cpp index 240005b48763..14cf92a7cc26 100644 --- a/offload/test/offloading/target_constexpr_mapping.cpp +++ b/offload/test/offloading/target_constexpr_mapping.cpp @@ -1,5 +1,4 @@ // RUN: %libomptarget-compileoptxx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/target_critical_region.cpp b/offload/test/offloading/target_critical_region.cpp index 836cce68e4fc..ddab2e754a66 100644 --- a/offload/test/offloading/target_critical_region.cpp +++ b/offload/test/offloading/target_critical_region.cpp @@ -4,8 +4,7 @@ // UNSUPPORTED: nvptx64-nvidia-cuda // UNSUPPORTED: nvptx64-nvidia-cuda-LTO // UNSUPPORTED: amdgcn-amd-amdhsa -// https://github.com/llvm/llvm-project/issues/182119 -// UNSUPPORTED: intelgpu +// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/target_depend_nowait.cpp b/offload/test/offloading/target_depend_nowait.cpp index 0a6b903db688..b0d5408770bc 100644 --- a/offload/test/offloading/target_depend_nowait.cpp +++ b/offload/test/offloading/target_depend_nowait.cpp @@ -1,5 +1,4 @@ // RUN: %libomptarget-compilexx-run-and-check-generic -// XFAIL: intelgpu #include #include diff --git a/offload/test/offloading/target_nowait_target.cpp b/offload/test/offloading/target_nowait_target.cpp index f6c2b56963cf..7c20e76ce138 100644 --- a/offload/test/offloading/target_nowait_target.cpp +++ b/offload/test/offloading/target_nowait_target.cpp @@ -1,6 +1,5 @@ // RUN: %libomptarget-compilexx-and-run-generic // RUN: %libomptarget-compileoptxx-and-run-generic -// XFAIL: intelgpu #include diff --git a/offload/test/offloading/target_update_from.c b/offload/test/offloading/target_update_from.c index 92697bc6b802..4f84968f6d9a 100644 --- a/offload/test/offloading/target_update_from.c +++ b/offload/test/offloading/target_update_from.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // This test checks that "update from" clause in OpenMP supports strided // sections. #pragma omp target update from(result[0:N/2:2]) updates every other // element from device diff --git a/offload/test/offloading/target_update_ptr_count_expression.c b/offload/test/offloading/target_update_ptr_count_expression.c index 97ee9fc5d019..c4b9fd566d40 100644 --- a/offload/test/offloading/target_update_ptr_count_expression.c +++ b/offload/test/offloading/target_update_ptr_count_expression.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // Tests non-contiguous array sections with expression-based count on // heap-allocated pointer arrays with both FROM and TO directives. diff --git a/offload/test/offloading/target_update_ptr_variable_count_and_stride.c b/offload/test/offloading/target_update_ptr_variable_count_and_stride.c index 3969cac05c5c..1a28595969c6 100644 --- a/offload/test/offloading/target_update_ptr_variable_count_and_stride.c +++ b/offload/test/offloading/target_update_ptr_variable_count_and_stride.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // Tests heap-allocated pointers with both variable count expression and // variable stride. diff --git a/offload/test/offloading/target_update_ptr_variable_stride.c b/offload/test/offloading/target_update_ptr_variable_stride.c index 8c0c62a2c8e8..bea396065b76 100644 --- a/offload/test/offloading/target_update_ptr_variable_stride.c +++ b/offload/test/offloading/target_update_ptr_variable_stride.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // Tests non-contiguous array sections with variable stride on heap-allocated // pointers. diff --git a/offload/test/offloading/target_update_strided_struct_count_expression.c b/offload/test/offloading/target_update_strided_struct_count_expression.c index 2dc3ea5d9b82..89b0e1b4c7ea 100644 --- a/offload/test/offloading/target_update_strided_struct_count_expression.c +++ b/offload/test/offloading/target_update_strided_struct_count_expression.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // Tests non-contiguous array sections with expression-based count on struct // member arrays with both FROM and TO directives. diff --git a/offload/test/offloading/target_update_strided_struct_from.c b/offload/test/offloading/target_update_strided_struct_from.c index 7a8805ab35e9..46448d10189f 100644 --- a/offload/test/offloading/target_update_strided_struct_from.c +++ b/offload/test/offloading/target_update_strided_struct_from.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // This test checks that "update from" with user-defined mapper supports strided // sections using fixed-size arrays in structs. diff --git a/offload/test/offloading/target_update_strided_struct_multiple_from.c b/offload/test/offloading/target_update_strided_struct_multiple_from.c index 16023be5970a..fa57d35b86a4 100644 --- a/offload/test/offloading/target_update_strided_struct_multiple_from.c +++ b/offload/test/offloading/target_update_strided_struct_multiple_from.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // This test checks that #pragma omp target update from(s1.data[0:6:2], // s2.data[0:4:3]) correctly updates strided sections covering the full arrays // from device to host. diff --git a/offload/test/offloading/target_update_strided_struct_multiple_to.c b/offload/test/offloading/target_update_strided_struct_multiple_to.c index 182b40cb06b3..69c401515358 100644 --- a/offload/test/offloading/target_update_strided_struct_multiple_to.c +++ b/offload/test/offloading/target_update_strided_struct_multiple_to.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // This test checks that #pragma omp target update to(s1.data[0:6:2], // s2.data[0:4:3]) correctly updates strided sections covering the full arrays // from host to device. diff --git a/offload/test/offloading/target_update_strided_struct_partial_from.c b/offload/test/offloading/target_update_strided_struct_partial_from.c index fb852c07f023..c0ffc72f900a 100644 --- a/offload/test/offloading/target_update_strided_struct_partial_from.c +++ b/offload/test/offloading/target_update_strided_struct_partial_from.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // This test checks that #pragma omp target update from(s.data[0:2:3]) correctly // updates every third element (stride 3) from the device to the host // using a struct with fixed-size array member. diff --git a/offload/test/offloading/target_update_strided_struct_partial_to.c b/offload/test/offloading/target_update_strided_struct_partial_to.c index 35c52ea2db20..018dcac7cdcf 100644 --- a/offload/test/offloading/target_update_strided_struct_partial_to.c +++ b/offload/test/offloading/target_update_strided_struct_partial_to.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // This test checks that #pragma omp target update to(s.data[0:2:3]) correctly // updates every third element (stride 3) from the host to the device // for struct member arrays. diff --git a/offload/test/offloading/target_update_strided_struct_to.c b/offload/test/offloading/target_update_strided_struct_to.c index 3c02232df280..90b3a5fe2698 100644 --- a/offload/test/offloading/target_update_strided_struct_to.c +++ b/offload/test/offloading/target_update_strided_struct_to.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // This test checks that "update to" with struct member arrays supports strided // sections using fixed-size arrays in structs. diff --git a/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c b/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c index 65af249e67ba..6daf10383e92 100644 --- a/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c +++ b/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // Tests struct member arrays with both variable count expression and variable // stride. diff --git a/offload/test/offloading/target_update_strided_struct_variable_stride.c b/offload/test/offloading/target_update_strided_struct_variable_stride.c index 54605cafac47..4cd9da629ca9 100644 --- a/offload/test/offloading/target_update_strided_struct_variable_stride.c +++ b/offload/test/offloading/target_update_strided_struct_variable_stride.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // Tests non-contiguous array sections with variable stride on struct member // arrays. diff --git a/offload/test/offloading/target_update_to.c b/offload/test/offloading/target_update_to.c index ca4984850a49..d45d3182a166 100644 --- a/offload/test/offloading/target_update_to.c +++ b/offload/test/offloading/target_update_to.c @@ -1,5 +1,4 @@ // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu // This test checks that "update to" clause in OpenMP supports strided sections. // #pragma omp target update to(result[0:8:2]) updates every other element // (stride 2) diff --git a/offload/test/offloading/thread_state_1.c b/offload/test/offloading/thread_state_1.c index 404ed6812928..b5f17f0c45dc 100644 --- a/offload/test/offloading/thread_state_1.c +++ b/offload/test/offloading/thread_state_1.c @@ -1,6 +1,7 @@ // RUN: %libomptarget-compile-run-and-check-generic // RUN: %libomptarget-compileopt-run-and-check-generic -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu #include #include diff --git a/offload/test/offloading/thread_state_2.c b/offload/test/offloading/thread_state_2.c index d40b987883c8..e709f5f33944 100644 --- a/offload/test/offloading/thread_state_2.c +++ b/offload/test/offloading/thread_state_2.c @@ -1,6 +1,7 @@ // This fails when optimized for now. // RUN: %libomptarget-compile-run-and-check-generic -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu // XUN: %libomptarget-compileopt-run-and-check-generic #include diff --git a/offload/test/sanitizer/kernel_crash.c b/offload/test/sanitizer/kernel_crash.c index 87133bfcfbd6..a0d1ecab9aff 100644 --- a/offload/test/sanitizer/kernel_crash.c +++ b/offload/test/sanitizer/kernel_crash.c @@ -12,7 +12,8 @@ // UNSUPPORTED: aarch64-unknown-linux-gnu // UNSUPPORTED: x86_64-unknown-linux-gnu // UNSUPPORTED: s390x-ibm-linux-gnu -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu #include diff --git a/offload/test/sanitizer/kernel_crash_async.c b/offload/test/sanitizer/kernel_crash_async.c index 48bd65162db8..73c34671acad 100644 --- a/offload/test/sanitizer/kernel_crash_async.c +++ b/offload/test/sanitizer/kernel_crash_async.c @@ -12,7 +12,8 @@ // UNSUPPORTED: aarch64-unknown-linux-gnu // UNSUPPORTED: x86_64-unknown-linux-gnu // UNSUPPORTED: s390x-ibm-linux-gnu -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu #include diff --git a/offload/test/sanitizer/kernel_crash_many.c b/offload/test/sanitizer/kernel_crash_many.c index 146b44586fa9..694b01792e8b 100644 --- a/offload/test/sanitizer/kernel_crash_many.c +++ b/offload/test/sanitizer/kernel_crash_many.c @@ -10,7 +10,8 @@ // UNSUPPORTED: aarch64-unknown-linux-gnu // UNSUPPORTED: x86_64-unknown-linux-gnu // UNSUPPORTED: s390x-ibm-linux-gnu -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu #include diff --git a/offload/test/sanitizer/kernel_trap.c b/offload/test/sanitizer/kernel_trap.c index 98d5faf241c8..3281bd2c8ca2 100644 --- a/offload/test/sanitizer/kernel_trap.c +++ b/offload/test/sanitizer/kernel_trap.c @@ -12,7 +12,8 @@ // UNSUPPORTED: aarch64-unknown-linux-gnu // UNSUPPORTED: x86_64-unknown-linux-gnu // UNSUPPORTED: s390x-ibm-linux-gnu -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu #include diff --git a/offload/test/sanitizer/kernel_trap.cpp b/offload/test/sanitizer/kernel_trap.cpp index a83e6cbaa78a..cb3a42a9daa4 100644 --- a/offload/test/sanitizer/kernel_trap.cpp +++ b/offload/test/sanitizer/kernel_trap.cpp @@ -13,7 +13,8 @@ // UNSUPPORTED: aarch64-unknown-linux-gnu // UNSUPPORTED: x86_64-unknown-linux-gnu // UNSUPPORTED: s390x-ibm-linux-gnu -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu struct S {}; diff --git a/offload/test/sanitizer/kernel_trap_async.c b/offload/test/sanitizer/kernel_trap_async.c index ae1bb8d322e4..466fee11f5d0 100644 --- a/offload/test/sanitizer/kernel_trap_async.c +++ b/offload/test/sanitizer/kernel_trap_async.c @@ -12,7 +12,8 @@ // UNSUPPORTED: aarch64-unknown-linux-gnu // UNSUPPORTED: x86_64-unknown-linux-gnu // UNSUPPORTED: s390x-ibm-linux-gnu -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu #include diff --git a/offload/test/sanitizer/kernel_trap_many.c b/offload/test/sanitizer/kernel_trap_many.c index c21cc9915969..72640728483f 100644 --- a/offload/test/sanitizer/kernel_trap_many.c +++ b/offload/test/sanitizer/kernel_trap_many.c @@ -10,7 +10,8 @@ // UNSUPPORTED: aarch64-unknown-linux-gnu // UNSUPPORTED: x86_64-unknown-linux-gnu // UNSUPPORTED: s390x-ibm-linux-gnu -// XFAIL: intelgpu +// https://github.com/llvm/llvm-project/issues/182119 +// UNSUPPORTED: intelgpu #include diff --git a/offload/test/unified_shared_memory/associate_ptr.c b/offload/test/unified_shared_memory/associate_ptr.c index 08975ebef002..0e9d25cd7354 100644 --- a/offload/test/unified_shared_memory/associate_ptr.c +++ b/offload/test/unified_shared_memory/associate_ptr.c @@ -2,7 +2,6 @@ // REQUIRES: unified_shared_memory // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9 -// XFAIL: intelgpu #include #include diff --git a/offload/test/unified_shared_memory/close_member.c b/offload/test/unified_shared_memory/close_member.c index d4743a587231..10ec42e90b02 100644 --- a/offload/test/unified_shared_memory/close_member.c +++ b/offload/test/unified_shared_memory/close_member.c @@ -2,7 +2,6 @@ // REQUIRES: unified_shared_memory // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9 -// XFAIL: intelgpu #include #include diff --git a/openmp/device/src/Synchronization.cpp b/openmp/device/src/Synchronization.cpp index 3554226d2ee7..93f10ebd3729 100644 --- a/openmp/device/src/Synchronization.cpp +++ b/openmp/device/src/Synchronization.cpp @@ -181,8 +181,27 @@ void setCriticalLock(omp_lock_t *Lock) { setLock(Lock); } ///} #if defined(__SPIRV__) -void namedBarrierInit() { __builtin_trap(); } // TODO -void namedBarrier() { __builtin_trap(); } // TODO + +[[clang::loader_uninitialized]] Local namedBarrierTracker; + +void namedBarrierInit() { + atomic::store(&namedBarrierTracker, 0u, atomic::seq_cst, atomic::workgroup); +} + +void namedBarrier() { + uint32_t NumThreads = omp_get_num_threads(); + uint32_t load = + atomic::add(&namedBarrierTracker, 1, atomic::seq_cst, atomic::workgroup); + + if (load >= NumThreads - 1) { + atomic::store(&namedBarrierTracker, 0u, atomic::seq_cst, atomic::workgroup); + } else { + do { + load = atomic::load(&namedBarrierTracker, atomic::seq_cst, + atomic::workgroup); + } while (load != 0); + } +} void unsetLock(omp_lock_t *Lock) { atomic::store((int32_t *)Lock, 0, atomic::seq_cst); @@ -194,13 +213,7 @@ void initLock(omp_lock_t *Lock) { unsetLock(Lock); } void destroyLock(omp_lock_t *Lock) { unsetLock(Lock); } void setLock(omp_lock_t *Lock) { int32_t *Lock_ptr = (int32_t *)Lock; - bool Acquired = false; - int32_t Expected; - while (!Acquired) { - Expected = 0; - if (Expected == atomic::load(Lock_ptr, atomic::seq_cst)) - Acquired = - atomic::cas(Lock_ptr, Expected, 1, atomic::seq_cst, atomic::seq_cst); + while (!atomic::cas(Lock_ptr, 0, 1, atomic::seq_cst, atomic::seq_cst)) { } }