Skip to content

Commit 309b167

Browse files
ianaylianayl
authored
[SYCL] Change some check_device_code tests to use SYCL_EXTERNAL (#13899)
Changed some of the simpler test cases to use SYCL_EXTERNAL functions instead of submitting to a queue.
1 parent 601f121 commit 309b167

10 files changed

+121
-187
lines changed

‎sycl/test/check_device_code/atomic_fence.cpp

Lines changed: 17 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -2,27 +2,20 @@
22

33
#include <sycl/sycl.hpp>
44

5-
int main() {
6-
sycl::queue Q;
7-
8-
Q.single_task([] {
9-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 896) #{{.*}}
10-
sycl::atomic_fence(sycl::memory_order::relaxed,
11-
sycl::memory_scope::work_group);
12-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 898) #{{.*}}
13-
sycl::atomic_fence(sycl::memory_order::acquire,
14-
sycl::memory_scope::work_group);
15-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 900) #{{.*}}
16-
sycl::atomic_fence(sycl::memory_order::release,
17-
sycl::memory_scope::work_group);
18-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 904) #{{.*}}
19-
sycl::atomic_fence(sycl::memory_order::acq_rel,
20-
sycl::memory_scope::work_group);
21-
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 912) #{{.*}}
22-
sycl::atomic_fence(sycl::memory_order::seq_cst,
23-
sycl::memory_scope::work_group);
24-
});
25-
Q.wait();
26-
27-
return 0;
28-
}
5+
SYCL_EXTERNAL void atomic_fence() {
6+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 896) #{{.*}}
7+
sycl::atomic_fence(sycl::memory_order::relaxed,
8+
sycl::memory_scope::work_group);
9+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 898) #{{.*}}
10+
sycl::atomic_fence(sycl::memory_order::acquire,
11+
sycl::memory_scope::work_group);
12+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 900) #{{.*}}
13+
sycl::atomic_fence(sycl::memory_order::release,
14+
sycl::memory_scope::work_group);
15+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 904) #{{.*}}
16+
sycl::atomic_fence(sycl::memory_order::acq_rel,
17+
sycl::memory_scope::work_group);
18+
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 912) #{{.*}}
19+
sycl::atomic_fence(sycl::memory_order::seq_cst,
20+
sycl::memory_scope::work_group);
21+
}

‎sycl/test/check_device_code/device_global_ptr_use.cpp

Lines changed: 4 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -10,11 +10,7 @@ using namespace sycl::ext::oneapi::experimental;
1010

1111
const device_global<int> DeviceGlobalVar;
1212

13-
int main() {
14-
queue Q;
15-
Q.single_task([]() {
16-
// CHECK: load {{.*}} @_ZL15DeviceGlobalVar
17-
volatile int ReadVal = DeviceGlobalVar;
18-
});
19-
return 0;
20-
}
13+
SYCL_EXTERNAL void global_ptr_use() {
14+
// CHECK: load {{.*}} @_ZL15DeviceGlobalVar
15+
volatile int ReadVal = DeviceGlobalVar;
16+
}

‎sycl/test/check_device_code/fpga_datapath_constructor.cpp

Lines changed: 4 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -28,14 +28,7 @@ class foo {
2828

2929
// CHECK: call {{.*}}sqrt
3030

31-
int main() {
32-
queue Q;
33-
int f = 5;
34-
35-
Q.single_task([=]() {
36-
intel::fpga_datapath<foo> mem{42};
37-
38-
volatile int ReadVal = mem.get().secret;
39-
});
40-
return 0;
41-
}
31+
SYCL_EXTERNAL void fetch_secret() {
32+
intel::fpga_datapath<foo> mem{42};
33+
volatile int ReadVal = mem.get().secret;
34+
}

‎sycl/test/check_device_code/fpga_datapath_global.cpp

Lines changed: 2 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -13,12 +13,9 @@ const intel::fpga_datapath<int[10]> empty{};
1313
// CHECK: %[[datapath:.*]] = type { [10 x i32] }
1414
// CHECK: {{.*}}empty = internal addrspace(1) constant %[[datapath]] zeroinitializer, align 4, !spirv.Decorations ![[empty_md:[0-9]*]]
1515

16-
int main() {
17-
queue Q;
16+
SYCL_EXTERNAL void fpga_datapath_global() {
1817
int f = 5;
19-
20-
Q.single_task([=]() { volatile int ReadVal = empty[f]; });
21-
return 0;
18+
volatile int ReadVal = empty[f];
2219
}
2320

2421
// CHECK: ![[empty_md]] = !{![[register:[0-9]*]]}

‎sycl/test/check_device_code/fpga_datapath_local.cpp

Lines changed: 5 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -10,15 +10,10 @@ namespace intel = sycl::ext::intel::experimental; // for fpga_datapath
1010

1111
// CHECK: [[RegisterINTEL:@.*]] = private unnamed_addr addrspace(1) constant [7 x i8] c"{5825}\00"
1212

13-
int main() {
14-
queue Q;
13+
SYCL_EXTERNAL void fpga_datapath_local() {
1514
int f = 5;
16-
17-
Q.single_task([=]() {
18-
intel::fpga_datapath<int[10]> empty;
19-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[RegisterINTEL]]
20-
// CHECK-NOT: call void @llvm.memset
21-
volatile int ReadVal = empty[f];
22-
});
23-
return 0;
15+
intel::fpga_datapath<int[10]> empty;
16+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[RegisterINTEL]]
17+
// CHECK-NOT: call void @llvm.memset
18+
volatile int ReadVal = empty[f];
2419
}

‎sycl/test/check_device_code/fpga_mem_constructor.cpp

Lines changed: 3 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -29,14 +29,7 @@ class foo {
2929

3030
// CHECK: call {{.*}}sqrt
3131

32-
int main() {
33-
queue Q;
34-
int f = 5;
35-
36-
Q.single_task([=]() {
37-
intel::fpga_mem<foo> mem{42};
38-
39-
volatile int ReadVal = mem.get().secret;
40-
});
41-
return 0;
32+
SYCL_EXTERNAL void fpga_mem_constructor() {
33+
intel::fpga_mem<foo> mem{42};
34+
volatile int ReadVal = mem.get().secret;
4235
}

‎sycl/test/check_device_code/fpga_mem_global.cpp

Lines changed: 5 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -65,18 +65,12 @@ const intel::fpga_mem<int[10],
6565
// CHECK: {{.*}}copies = internal addrspace(1) constant {{.*}} zeroinitializer, align 4, !spirv.Decorations ![[copies_md:[0-9]*]]
6666
// CHECK: {{.*}}replicates = internal addrspace(1) constant {{.*}} zeroinitializer, align 4, !spirv.Decorations ![[replicates_md:[0-9]*]]
6767

68-
int main() {
69-
queue Q;
68+
SYCL_EXTERNAL void fpga_mem_global() {
7069
int f = 5;
71-
72-
Q.single_task([=]() {
73-
volatile int ReadVal = empty[f] + min_ram[f] + max_fmax[f] +
74-
double_pumped[f] + single_pumped[f] + mlab[f] +
75-
simple_dual_port[f] + true_dual_port[f] +
76-
block_ram[f] + banks[f] + stride[f] + word[f] +
77-
copies[f] + replicates[f];
78-
});
79-
return 0;
70+
volatile int ReadVal =
71+
empty[f] + min_ram[f] + max_fmax[f] + double_pumped[f] +
72+
single_pumped[f] + mlab[f] + simple_dual_port[f] + true_dual_port[f] +
73+
block_ram[f] + banks[f] + stride[f] + word[f] + copies[f] + replicates[f];
8074
}
8175

8276
// CHECK: ![[empty_md]] = !{![[mem_default:[0-9]*]]}

‎sycl/test/check_device_code/fpga_mem_local.cpp

Lines changed: 67 additions & 75 deletions
Original file line numberDiff line numberDiff line change
@@ -24,81 +24,73 @@ namespace oneapi = sycl::ext::oneapi::experimental; // for properties
2424
// CHECK: [[MaxPrivateCopiesINTEL:@.*]] = private unnamed_addr addrspace(1) constant [27 x i8] c"{5826:\22DEFAULT\22}{5829:\223\22}\00"
2525
// CHECK: [[MaxReplicatesINTEL:@.*]] = private unnamed_addr addrspace(1) constant [27 x i8] c"{5826:\22DEFAULT\22}{5832:\225\22}\00"
2626

27-
int main() {
28-
queue Q;
27+
SYCL_EXTERNAL void fpga_mem_local() {
2928
int f = 5;
29+
intel::fpga_mem<int[10]> empty;
30+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL]]
31+
// CHECK-NOT: call void @llvm.memset
32+
intel::fpga_mem<int[10],
33+
decltype(oneapi::properties(intel::ram_stitching_min_ram))>
34+
min_ram;
35+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_FALSE]]
36+
// CHECK-NOT: call void @llvm.memset
37+
intel::fpga_mem<int[10],
38+
decltype(oneapi::properties(intel::ram_stitching_max_fmax))>
39+
max_fmax;
40+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_TRUE]]
41+
// CHECK-NOT: call void @llvm.memset
42+
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::clock_2x_true))>
43+
double_pumped;
44+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[DoublepumpINTEL]]
45+
// CHECK-NOT: call void @llvm.memset
46+
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::clock_2x_false))>
47+
single_pumped;
48+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SinglepumpINTEL]]
49+
// CHECK-NOT: call void @llvm.memset
50+
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::resource_mlab))>
51+
mlab;
52+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_mlab]]
53+
// CHECK-NOT: call void @llvm.memset
54+
intel::fpga_mem<int[10], decltype(oneapi::properties(
55+
intel::bi_directional_ports_false))>
56+
simple_dual_port;
57+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SimpleDualPortINTEL]]
58+
// CHECK-NOT: call void @llvm.memset
59+
intel::fpga_mem<int[10], decltype(oneapi::properties(
60+
intel::bi_directional_ports_true))>
61+
true_dual_port;
62+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[TrueDualPortINTEL]]
63+
// CHECK-NOT: call void @llvm.memset
64+
intel::fpga_mem<int[10],
65+
decltype(oneapi::properties(intel::resource_block_ram))>
66+
block_ram;
67+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_block_ram]]
68+
// CHECK-NOT: call void @llvm.memset
69+
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::num_banks<4>))>
70+
banks;
71+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[NumbanksINTEL]]
72+
// CHECK-NOT: call void @llvm.memset
73+
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::stride_size<2>))>
74+
stride;
75+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[StridesizeINTEL]]
76+
// CHECK-NOT: call void @llvm.memset
77+
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::word_size<8>))>
78+
word;
79+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[WordsizeINTEL]]
80+
// CHECK-NOT: call void @llvm.memset
81+
intel::fpga_mem<int[10],
82+
decltype(oneapi::properties(intel::max_private_copies<3>))>
83+
copies;
84+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxPrivateCopiesINTEL]]
85+
// CHECK-NOT: call void @llvm.memset
86+
intel::fpga_mem<int[10],
87+
decltype(oneapi::properties(intel::num_replicates<5>))>
88+
replicates;
89+
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxReplicatesINTEL]]
90+
// CHECK-NOT: call void @llvm.memset
3091

31-
Q.single_task([=]() {
32-
intel::fpga_mem<int[10]> empty;
33-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL]]
34-
// CHECK-NOT: call void @llvm.memset
35-
intel::fpga_mem<int[10],
36-
decltype(oneapi::properties(intel::ram_stitching_min_ram))>
37-
min_ram;
38-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_FALSE]]
39-
// CHECK-NOT: call void @llvm.memset
40-
intel::fpga_mem<int[10],
41-
decltype(oneapi::properties(intel::ram_stitching_max_fmax))>
42-
max_fmax;
43-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_TRUE]]
44-
// CHECK-NOT: call void @llvm.memset
45-
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::clock_2x_true))>
46-
double_pumped;
47-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[DoublepumpINTEL]]
48-
// CHECK-NOT: call void @llvm.memset
49-
intel::fpga_mem<int[10],
50-
decltype(oneapi::properties(intel::clock_2x_false))>
51-
single_pumped;
52-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SinglepumpINTEL]]
53-
// CHECK-NOT: call void @llvm.memset
54-
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::resource_mlab))>
55-
mlab;
56-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_mlab]]
57-
// CHECK-NOT: call void @llvm.memset
58-
intel::fpga_mem<int[10], decltype(oneapi::properties(
59-
intel::bi_directional_ports_false))>
60-
simple_dual_port;
61-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SimpleDualPortINTEL]]
62-
// CHECK-NOT: call void @llvm.memset
63-
intel::fpga_mem<int[10], decltype(oneapi::properties(
64-
intel::bi_directional_ports_true))>
65-
true_dual_port;
66-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[TrueDualPortINTEL]]
67-
// CHECK-NOT: call void @llvm.memset
68-
intel::fpga_mem<int[10],
69-
decltype(oneapi::properties(intel::resource_block_ram))>
70-
block_ram;
71-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_block_ram]]
72-
// CHECK-NOT: call void @llvm.memset
73-
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::num_banks<4>))>
74-
banks;
75-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[NumbanksINTEL]]
76-
// CHECK-NOT: call void @llvm.memset
77-
intel::fpga_mem<int[10],
78-
decltype(oneapi::properties(intel::stride_size<2>))>
79-
stride;
80-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[StridesizeINTEL]]
81-
// CHECK-NOT: call void @llvm.memset
82-
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::word_size<8>))>
83-
word;
84-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[WordsizeINTEL]]
85-
// CHECK-NOT: call void @llvm.memset
86-
intel::fpga_mem<int[10],
87-
decltype(oneapi::properties(intel::max_private_copies<3>))>
88-
copies;
89-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxPrivateCopiesINTEL]]
90-
// CHECK-NOT: call void @llvm.memset
91-
intel::fpga_mem<int[10],
92-
decltype(oneapi::properties(intel::num_replicates<5>))>
93-
replicates;
94-
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxReplicatesINTEL]]
95-
// CHECK-NOT: call void @llvm.memset
96-
97-
volatile int ReadVal = empty[f] + min_ram[f] + max_fmax[f] +
98-
double_pumped[f] + single_pumped[f] + mlab[f] +
99-
simple_dual_port[f] + true_dual_port[f] +
100-
block_ram[f] + banks[f] + stride[f] + word[f] +
101-
copies[f] + replicates[f];
102-
});
103-
return 0;
92+
volatile int ReadVal =
93+
empty[f] + min_ram[f] + max_fmax[f] + double_pumped[f] +
94+
single_pumped[f] + mlab[f] + simple_dual_port[f] + true_dual_port[f] +
95+
block_ram[f] + banks[f] + stride[f] + word[f] + copies[f] + replicates[f];
10496
}

‎sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp

Lines changed: 7 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -21,20 +21,11 @@ int arrayAdd(int *data1, int *data2, int N) {
2121
return ret;
2222
}
2323

24-
int main() {
25-
sycl::queue myQueue;
26-
27-
myQueue.submit([&](sycl::handler &cgh) {
28-
cgh.single_task([=]() {
29-
int d1[kSize], d2[kSize];
30-
task_sequence<arrayAdd,
31-
decltype(properties{pipelined<0>, stall_enable_clusters,
32-
invocation_capacity<1>,
33-
response_capacity<1>})>
34-
arrayAddTask;
35-
arrayAddTask.async(d1, d2, kSize);
36-
});
37-
});
38-
myQueue.wait();
39-
return 0;
24+
SYCL_EXTERNAL void task_sequence_no_explicit_get() {
25+
int d1[kSize], d2[kSize];
26+
task_sequence<arrayAdd, decltype(properties{
27+
pipelined<0>, stall_enable_clusters,
28+
invocation_capacity<1>, response_capacity<1>})>
29+
arrayAddTask;
30+
arrayAddTask.async(d1, d2, kSize);
4031
}

‎sycl/test/check_device_code/usm_pointers.cpp

Lines changed: 7 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -30,20 +30,10 @@
3030

3131
using namespace sycl;
3232

33-
int main() {
34-
sycl::queue queue;
35-
{
36-
queue.submit([&](sycl::handler &cgh) {
37-
cgh.single_task<class check_adress_space>([=]() {
38-
void *Ptr = nullptr;
39-
ext::intel::device_ptr<void> DevPtr(Ptr);
40-
ext::intel::host_ptr<void> HostPtr(Ptr);
41-
global_ptr<void> GlobPtr = global_ptr<void>(DevPtr);
42-
GlobPtr = global_ptr<void>(HostPtr);
43-
});
44-
});
45-
queue.wait();
46-
}
47-
48-
return 0;
49-
}
33+
SYCL_EXTERNAL void usm_pointers() {
34+
void *Ptr = nullptr;
35+
ext::intel::device_ptr<void> DevPtr(Ptr);
36+
ext::intel::host_ptr<void> HostPtr(Ptr);
37+
global_ptr<void> GlobPtr = global_ptr<void>(DevPtr);
38+
GlobPtr = global_ptr<void>(HostPtr);
39+
}

0 commit comments

Comments
 (0)