Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 68e4ef3

Browse files
raaiq1raaiq1
authored
Merge branch 'intel:intel' into ext_namespace
2 parents 8a3827a + d6e43f6 commit 68e4ef3

File tree

369 files changed

+4665
-6416
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

369 files changed

+4665
-6416
lines changed

‎SYCL/AOT/multiple-devices.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -15,26 +15,22 @@
1515

1616
// CPU, GPU, FPGA
1717
// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen,spir64_fpga -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_all_aot.out
18-
// RUN: %HOST_RUN_PLACEHOLDER %t_all_aot.out
1918
// RUN: %CPU_RUN_PLACEHOLDER %t_all_aot.out
2019
// RUN: %GPU_RUN_PLACEHOLDER %t_all_aot.out
2120
// RUN: %ACC_RUN_PLACEHOLDER %t_all_aot.out
2221

2322
// CPU, GPU
2423
// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_cpu_gpu.out
25-
// RUN: %HOST_RUN_PLACEHOLDER %t_cpu_gpu.out
2624
// RUN: %CPU_RUN_PLACEHOLDER %t_cpu_gpu.out
2725
// RUN: %GPU_RUN_PLACEHOLDER %t_cpu_gpu.out
2826

2927
// CPU, FPGA
3028
// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_fpga %t.o -o %t_cpu_fpga.out
31-
// RUN: %HOST_RUN_PLACEHOLDER %t_cpu_fpga.out
3229
// RUN: %CPU_RUN_PLACEHOLDER %t_cpu_fpga.out
3330
// RUN: %ACC_RUN_PLACEHOLDER %t_cpu_fpga.out
3431

3532
// GPU, FPGA
3633
// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen,spir64_fpga -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %t.o -o %t_gpu_fpga.out
37-
// RUN: %HOST_RUN_PLACEHOLDER %t_gpu_fpga.out
3834
// RUN: %GPU_RUN_PLACEHOLDER %t_gpu_fpga.out
3935
// RUN: %ACC_RUN_PLACEHOLDER %t_gpu_fpga.out
4036

‎SYCL/Assert/assert_in_simultaneously_multiple_tus.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,6 @@
1+
// FIXME flaky fail on CUDA and HIP
2+
// UNSUPPORTED: cuda || hip
3+
//
14
// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -I %S/Inputs %s %S/Inputs/kernels_in_file2.cpp -o %t.out %threads_lib
25
// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
36
// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt

‎SYCL/Assert/assert_in_simultaneously_multiple_tus_one_ndebug.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
1+
// FIXME flaky fail on CUDA
2+
// UNSUPPORTED: cuda
13
// RUN: %clangxx -DSYCL_FALLBACK_ASSERT=1 -fsycl -fsycl-targets=%sycl_triple -DDEFINE_NDEBUG_INFILE2 -I %S/Inputs %S/assert_in_simultaneously_multiple_tus.cpp %S/Inputs/kernels_in_file2.cpp -o %t.out %threads_lib
24
// RUN: %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
35
// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt

‎SYCL/AtomicRef/accessor.cpp

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %CPU_RUN_PLACEHOLDER %t.out
43
// RUN: %GPU_RUN_PLACEHOLDER %t.out
54

@@ -101,9 +100,6 @@ int main() {
101100
queue q;
102101
constexpr int N = 32;
103102
accessor_test<int>(q, N);
104-
// TODO: Enable local accessor test for host when barrier is supported
105-
if (!q.get_device().is_host()) {
106-
local_accessor_test<int>(q, N);
107-
}
103+
local_accessor_test<int>(q, N);
108104
std::cout << "Test passed." << std::endl;
109105
}

‎SYCL/AtomicRef/add.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,6 @@
22
// UNSUPPORTED: hip
33

44
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-device-code-split=per_kernel -fsycl-targets=%sycl_triple %s -o %t.out
5-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
65
// RUN: %GPU_RUN_PLACEHOLDER %t.out
76
// RUN: %CPU_RUN_PLACEHOLDER %t.out
87
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/add.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void add_fetch_local_test(queue q, size_t N) {
3030
auto sum = sum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/add_generic.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/add_generic_local.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,10 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65

76
// CUDA and HIP backends have had no support for the generic address space yet.
8-
// Host does not support barrier.
9-
// XFAIL: cuda || hip || host
7+
// XFAIL: cuda || hip
108

119
#define TEST_GENERIC_IN_LOCAL 1
1210

‎SYCL/AtomicRef/add_generic_local_native_fp.cpp

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,10 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65

76
// CUDA and HIP backends have had no support for the generic address space yet.
8-
// Host does not support barrier. HIP does not support native floating point
9-
// atomics
10-
// XFAIL: cuda, hip, host
7+
// XFAIL: cuda, hip
118

129
#define SYCL_USE_NATIVE_FP_ATOMICS
1310
#define FP_TESTS_ONLY

‎SYCL/AtomicRef/add_generic_native_fp.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/add_local.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,10 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65

7-
// Barrier is not supported on host. HIP does not support floating point
8-
// atomics.
9-
// XFAIL: host, hip
6+
// HIP does not support floating point atomics.
7+
// XFAIL: hip
108

119
#include "add.h"
1210

‎SYCL/AtomicRef/add_local_native_fp.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,10 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65

7-
// Barrier is not supported on host. HIP does not support floating point
8-
// atomics.
9-
// XFAIL: host, hip
6+
// HIP does not support floating point atomics.
7+
// XFAIL: hip
108

119
#define SYCL_USE_NATIVE_FP_ATOMICS
1210
#define FP_TESTS_ONLY

‎SYCL/AtomicRef/add_native_fp.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/and.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/and.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,8 +30,7 @@ void and_local_test(queue q) {
3030
auto cum = cum_buf.template get_access<access::mode::read_write>(cgh);
3131
auto out =
3232
output_buf.template get_access<access::mode::discard_write>(cgh);
33-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
34-
cgh);
33+
local_accessor<T, 1> loc(1, cgh);
3534

3635
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3736
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/and_generic.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/and_generic_local.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,10 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65

76
// CUDA and HIP backends have had no support for the generic address space yet.
8-
// Host does not support barrier.
9-
// XFAIL: cuda || hip || host
7+
// XFAIL: cuda || hip
108

119
#define TEST_GENERIC_IN_LOCAL 1
1210

‎SYCL/AtomicRef/and_local.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,8 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65

7-
// Barrier is not supported on host.
8-
// XFAIL: host
9-
106
#include "and.h"
117

128
int main() { and_test_all<access::address_space::local_space>(); }

‎SYCL/AtomicRef/assignment.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %CPU_RUN_PLACEHOLDER %t.out
43
// RUN: %GPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/assignment_atomic64.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %CPU_RUN_PLACEHOLDER %t.out
43
// RUN: %GPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/assignment_atomic64_generic.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %CPU_RUN_PLACEHOLDER %t.out
43
// RUN: %GPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/assignment_generic.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %CPU_RUN_PLACEHOLDER %t.out
43
// RUN: %GPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/atomic_memory_order.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %CPU_RUN_PLACEHOLDER %t.out
43
// RUN: %GPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/atomic_memory_order_acq_rel.cpp

Lines changed: 2 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1,15 +1,11 @@
11
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -O3 -o %t.out -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %CPU_RUN_PLACEHOLDER %t.out
43
// RUN: %GPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65
// L0, OpenCL, and HIP backends don't currently support
76
// info::device::atomic_memory_order_capabilities
87
// UNSUPPORTED: level_zero, opencl, hip
98

10-
// host does not support barrier
11-
// XFAIL: host
12-
139
// NOTE: Tests fetch_add for acquire and release memory ordering.
1410

1511
#include "atomic_memory_order.h"
@@ -76,8 +72,7 @@ template <memory_order order> void test_acquire_local() {
7672
q.submit([&](handler &cgh) {
7773
auto error =
7874
error_buf.template get_access<access::mode::read_write>(cgh);
79-
accessor<int, 1, access::mode::read_write, access::target::local> val(
80-
2, cgh);
75+
local_accessor<int, 1> val(2, cgh);
8176
cgh.parallel_for(
8277
nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
8378
size_t lid = it.get_local_id(0);
@@ -168,8 +163,7 @@ template <memory_order order> void test_release_local() {
168163
q.submit([&](handler &cgh) {
169164
auto error =
170165
error_buf.template get_access<access::mode::read_write>(cgh);
171-
accessor<int, 1, access::mode::read_write, access::target::local> val(
172-
2, cgh);
166+
local_accessor<int, 1> val(2, cgh);
173167
cgh.parallel_for(
174168
nd_range<1>(global_size, local_size), [=](nd_item<1> it) {
175169
size_t lid = it.get_local_id(0);

‎SYCL/AtomicRef/atomic_memory_order_seq_cst.cpp

Lines changed: 1 addition & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,15 +1,11 @@
11
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -O3 -o %t.out -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %CPU_RUN_PLACEHOLDER %t.out
43
// RUN: %GPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65
// L0, OpenCL, and HIP backends don't currently support
76
// info::device::atomic_memory_order_capabilities
87
// UNSUPPORTED: level_zero, opencl, hip
98

10-
// host does not support barrier
11-
// XFAIL: host
12-
139
#include "atomic_memory_order.h"
1410
#include <iostream>
1511
#include <numeric>
@@ -120,8 +116,7 @@ template <memory_order order> void test_local() {
120116

121117
q.submit([&](handler &cgh) {
122118
auto res = res_buf.template get_access<access::mode::discard_write>(cgh);
123-
accessor<int, 1, access::mode::read_write, access::target::local> val(2,
124-
cgh);
119+
local_accessor<int, 1> val(2, cgh);
125120
cgh.parallel_for(nd_range<1>(N_items, N_items), [=](nd_item<1> it) {
126121
val[0] = 0;
127122
it.barrier(access::fence_space::local_space);

‎SYCL/AtomicRef/compare_exchange.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/compare_exchange.h

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -32,8 +32,7 @@ void compare_exchange_local_test(queue q, size_t N) {
3232
cgh);
3333
auto out =
3434
output_buf.template get_access<access::mode::discard_write>(cgh);
35-
accessor<T, 1, access::mode::read_write, access::target::local> loc(1,
36-
cgh);
35+
local_accessor<T, 1> loc(1, cgh);
3736

3837
cgh.parallel_for(nd_range<1>(N, N), [=](nd_item<1> it) {
3938
int gid = it.get_global_id(0);

‎SYCL/AtomicRef/compare_exchange_generic.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out

‎SYCL/AtomicRef/compare_exchange_generic_local.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,10 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65

76
// CUDA and HIP backends have had no support for the generic address space yet.
8-
// Host does not support barrier.
9-
// XFAIL: cuda || hip || host
7+
// XFAIL: cuda || hip
108

119
#define TEST_GENERIC_IN_LOCAL 1
1210

‎SYCL/AtomicRef/compare_exchange_local.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,8 @@
11
// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
2-
// RUN: %HOST_RUN_PLACEHOLDER %t.out
32
// RUN: %GPU_RUN_PLACEHOLDER %t.out
43
// RUN: %CPU_RUN_PLACEHOLDER %t.out
54
// RUN: %ACC_RUN_PLACEHOLDER %t.out
65

7-
// Barrier is not supported on host.
8-
// XFAIL: host
9-
106
#include "compare_exchange.h"
117

128
int main() { compare_exchange_test_all<access::address_space::local_space>(); }

‎SYCL/AtomicRef/device_has_aspect_atomic64_cuda_and_hip.cpp

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
// REQUIRES: cuda || hip
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
3+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
4+
5+
// XFAIL: hip
6+
// Expected failure because hip does not have atomic64 check implementation
7+
8+
#include <CL/sycl.hpp>
9+
#include <iostream>
10+
11+
using namespace sycl;
12+
13+
int main() {
14+
queue Queue;
15+
device Dev = Queue.get_device();
16+
// cout in order to ensure that the query hasn't been optimized out
17+
std::cout << Dev.has(aspect::atomic64) << std::endl;
18+
return 0;
19+
}

‎SYCL/AtomicRef/device_has_aspect_atomic64_level_zero.cpp

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
// REQUIRES: level_zero, level_zero_dev_kit
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out %level_zero_options
3+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
4+
5+
#include <CL/sycl.hpp>
6+
#include <level_zero/ze_api.h>
7+
8+
using namespace sycl;
9+
10+
int main() {
11+
queue Queue;
12+
device Dev = Queue.get_device();
13+
bool Result;
14+
ze_device_module_properties_t Properties;
15+
zeDeviceGetModuleProperties(get_native<backend::ext_oneapi_level_zero>(Dev),
16+
&Properties);
17+
if (Properties.flags & ZE_DEVICE_MODULE_FLAG_INT64_ATOMICS)
18+
Result = true;
19+
else
20+
Result = false;
21+
assert(Dev.has(aspect::atomic64) == Result &&
22+
"The Result value differs from the implemented atomic64 check on "
23+
"the L0 backend.");
24+
return 0;
25+
}

0 commit comments

Comments
 (0)