Skip to content

[SYCL] fix e2e atomic ref tests #23559

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
sarnex merged 12 commits into from
Oct 21, 2026
+280 −138
Merged

[SYCL] fix e2e atomic ref tests #23559

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
350fd88
[SYCL][E2E] do not use conditional build and run in tests
dklochkov-emb Oct 20, 2026
564ae56
[SYCL][E2E] fix spirv tranclator tripple issue
dklochkov-emb Oct 20, 2026
f681676
Merge remote-tracking branch 'upstream/sycl' into sycl-atomic-ref-tes…
dklochkov-emb Oct 20, 2026
22a930e
[SYCL][E2E] exclude 16 ibt atomic testing when nvidia target
dklochkov-emb Oct 20, 2026
0f27b89
[SYCL][E2E] update if statement in test run
dklochkov-emb Oct 21, 2026
bed0d54
[SYCL][E2E] fix formatting
dklochkov-emb Oct 21, 2026
6901f14
Merge branch 'sycl' into sycl-atomic-ref-test-fail
dklochkov-emb Oct 21, 2026
75b5a7d
[SYCL][E2E] fix spirv translator issue
dklochkov-emb Oct 21, 2026
5a91a10
[SYCL][E2E] return requires into test
dklochkov-emb Oct 21, 2026
f5a0c33
[SYCL][E2E] guard bfloat tests
dklochkov-emb Oct 21, 2026
a7e9b21
[SYCL][E2E] add unsupported targets
dklochkov-emb Oct 21, 2026
9269024
Merge branch 'sycl' into sycl-atomic-ref-test-fail
dklochkov-emb Oct 21, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/add.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/add_generic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/add_generic_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/add_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/and.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
21 changes: 12 additions & 9 deletions sycl/test-e2e/AtomicRef/and_generic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,9 +1,12 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out

#include "and.h"

int main() { and_test_all<access::address_space::generic_space>(); }
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out

#include "and.h"

int main() { and_test_all<access::address_space::generic_space>(); }
25 changes: 14 additions & 11 deletions sycl/test-e2e/AtomicRef/and_generic_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,11 +1,14 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out

#define TEST_GENERIC_IN_LOCAL 1

#include "and.h"

int main() { and_test_all<access::address_space::generic_space>(); }
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out

#define TEST_GENERIC_IN_LOCAL 1

#include "and.h"

int main() { and_test_all<access::address_space::generic_space>(); }
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/and_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri

@againull againull Oct 21, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@dklochkov-emb Is there a particular reason we're limiting all atomic tests to CRI only?

@dklochkov-emb dklochkov-emb Oct 21, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@againull yes, it is not confirmed that other drivers support atomics.
Moreover, all tests were disabled before.

@dklochkov-emb dklochkov-emb Oct 21, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@againull in case of the build line
%if arch-intel_gpu_cri %{ %{build}..
fusion CI somehow skipped it but in case of run %if arch-intel_gpu_cri %{ %{run} %t.out %} somehow it tried to run resulting error that binary was not found

@againull againull Oct 21, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@againull yes, it is not confirmed that other drivers support atomics.

It may apply to the new 16-bit atomics, but not to atomics generally: 32-bit atomics are required by SYCL, and the existing tests were already running and passing on other platforms, including PVC.(https://github.com/intel/llvm/actions/runs/35870937631/job/107222525795).

For the 16-bit tests, we should at least use REQUIRES: aspect-ext_oneapi_atomic16 instead of hardcoding arch-intel_gpu_cri. Ideally, the 16-bit coverage should be in separate test wrappers: combining it with the existing 32-bit tests under a file-wide requirement causes us to lose 32-bit coverage on devices without atomic16 support.

Moreover, all tests were disabled before.

Could you elaborate on this? In this older PVC run (https://github.com/intel/llvm/actions/runs/35870937631/job/107222525795), the atomic tests were executed and passed. Were you referring specifically to the 16-bit cases?

@dklochkov-emb dklochkov-emb Oct 21, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ok, I will fix that.

Were you referring specifically to the 16-bit cases?

The fact is that all macro were disabled and they cut almost all tests except..1 data type. In some tests that was float:
https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/AtomicRef/add.h#L540
in some case that was int:
https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/AtomicRef/and.h#L233
I can guess that only reason is that atomic operations with these data types are supported by default and Iadd 32-bit is not supported in evert GPU

I will enable tests by default as it was done earlier.

// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
14 changes: 12 additions & 2 deletions sycl/test-e2e/AtomicRef/atomic_ref.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,9 +1,11 @@
// REQUIRES: arch-intel_gpu_cri
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -Xclang -freg-struct-return -Xspirv-translator=spir64 --spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out
// RUN: %{build} -o %t.out
// RUN: %{run} %t.out

// UNSUPPORTED: target-nvidia, target-amd, spirv-backend
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support

#include <sycl/atomic_ref.hpp>
Expand Down Expand Up @@ -147,22 +149,30 @@ int main() {
int ret = test_atomic_sub<sycl::half>();
ret |= test_atomic_sub<unsigned short>();
ret |= test_atomic_sub<short>();
#ifdef FULL_ATOMIC16_COVERAGE
ret |= test_atomic_sub<sycl::ext::oneapi::bfloat16>();
#endif

ret |= test_atomic_add<sycl::half>();
ret |= test_atomic_add<unsigned short>();
ret |= test_atomic_add<short>();
#ifdef FULL_ATOMIC16_COVERAGE
ret |= test_atomic_add<sycl::ext::oneapi::bfloat16>();
#endif

ret |= test_atomic_min<sycl::half>();
ret |= test_atomic_min<unsigned short>();
ret |= test_atomic_min<short>();
#ifdef FULL_ATOMIC16_COVERAGE
ret |= test_atomic_min<sycl::ext::oneapi::bfloat16>();
#endif

ret |= test_atomic_max<sycl::half>();
ret |= test_atomic_max<unsigned short>();
ret |= test_atomic_max<short>();
#ifdef FULL_ATOMIC16_COVERAGE
ret |= test_atomic_max<sycl::ext::oneapi::bfloat16>();
#endif

ret |= test_atomic_or<unsigned short>();
ret |= test_atomic_or<short>();
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/compare_exchange.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/compare_exchange_generic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/compare_exchange_generic_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/compare_exchange_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/exchange.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/exchange_generic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/exchange_generic_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/exchange_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/load.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/load_generic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/load_generic_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/load_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/max.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
7 changes: 5 additions & 2 deletions sycl/test-e2e/AtomicRef/max_generic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,8 @@
// RUN: %if arch-intel_gpu_cri %{ %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.out %}
// RUN: %if arch-intel_gpu_cri %{ %{run} %t.out %}
// REQUIRES: arch-intel_gpu_cri
// UNSUPPORTED: any-target-is-nvidia, any-target-is-amd, spirv-backend
// UNSUPPORTED-INTENDED: only supported by backends with atomic16 support
// RUN: %{build} -DFULL_ATOMIC16_COVERAGE -DFULL_ATOMIC32_COVERAGE -DFULL_ATOMIC64_COVERAGE -fsycl-device-code-split=per_kernel -Xspirv-translator=spir64 -spirv-ext=+SPV_KHR_bfloat16,+SPV_INTEL_16bit_atomics -o %t.cri.out
// RUN: %{run} %t.cri.out

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out
Expand Down
Loading
Loading