Skip to content
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

[SYCL] Change some check_device_code tests to use SYCL_EXTERNAL #13899

Merged
merged 2 commits into from
May 27, 2024
Merged

[SYCL] Change some check_device_code tests to use SYCL_EXTERNAL #13899

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
c0809ed
Changed some tests in check_device_code to use SYCL_EXTERNAL instead …
ianayl May 23, 2024
0e25e0e
Applied clang-format to files
ianayl May 24, 2024
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
41 changes: 17 additions & 24 deletions sycl/test/check_device_code/atomic_fence.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,27 +2,20 @@

#include <sycl/sycl.hpp>

int main() {
sycl::queue Q;

Q.single_task([] {
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 896) #{{.*}}
sycl::atomic_fence(sycl::memory_order::relaxed,
sycl::memory_scope::work_group);
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 898) #{{.*}}
sycl::atomic_fence(sycl::memory_order::acquire,
sycl::memory_scope::work_group);
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 900) #{{.*}}
sycl::atomic_fence(sycl::memory_order::release,
sycl::memory_scope::work_group);
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 904) #{{.*}}
sycl::atomic_fence(sycl::memory_order::acq_rel,
sycl::memory_scope::work_group);
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 912) #{{.*}}
sycl::atomic_fence(sycl::memory_order::seq_cst,
sycl::memory_scope::work_group);
});
Q.wait();

return 0;
}
SYCL_EXTERNAL void atomic_fence() {
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 896) #{{.*}}
sycl::atomic_fence(sycl::memory_order::relaxed,
sycl::memory_scope::work_group);
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 898) #{{.*}}
sycl::atomic_fence(sycl::memory_order::acquire,
sycl::memory_scope::work_group);
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 900) #{{.*}}
sycl::atomic_fence(sycl::memory_order::release,
sycl::memory_scope::work_group);
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 904) #{{.*}}
sycl::atomic_fence(sycl::memory_order::acq_rel,
sycl::memory_scope::work_group);
// CHECK: tail call spir_func void @_Z21__spirv_MemoryBarrierjj(i32 noundef 2, i32 noundef 912) #{{.*}}
sycl::atomic_fence(sycl::memory_order::seq_cst,
sycl::memory_scope::work_group);
}
12 changes: 4 additions & 8 deletions sycl/test/check_device_code/device_global_ptr_use.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,11 +10,7 @@ using namespace sycl::ext::oneapi::experimental;

const device_global<int> DeviceGlobalVar;

int main() {
queue Q;
Q.single_task([]() {
// CHECK: load {{.*}} @_ZL15DeviceGlobalVar
volatile int ReadVal = DeviceGlobalVar;
});
return 0;
}
SYCL_EXTERNAL void global_ptr_use() {
// CHECK: load {{.*}} @_ZL15DeviceGlobalVar
volatile int ReadVal = DeviceGlobalVar;
}
15 changes: 4 additions & 11 deletions sycl/test/check_device_code/fpga_datapath_constructor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,14 +28,7 @@ class foo {

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

int main() {
queue Q;
int f = 5;

Q.single_task([=]() {
intel::fpga_datapath<foo> mem{42};

volatile int ReadVal = mem.get().secret;
});
return 0;
}
SYCL_EXTERNAL void fetch_secret() {
intel::fpga_datapath<foo> mem{42};
volatile int ReadVal = mem.get().secret;
}
7 changes: 2 additions & 5 deletions sycl/test/check_device_code/fpga_datapath_global.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,12 +13,9 @@ const intel::fpga_datapath<int[10]> empty{};
// CHECK: %[[datapath:.*]] = type { [10 x i32] }
// CHECK: {{.*}}empty = internal addrspace(1) constant %[[datapath]] zeroinitializer, align 4, !spirv.Decorations ![[empty_md:[0-9]*]]

int main() {
queue Q;
SYCL_EXTERNAL void fpga_datapath_global() {
Copy link
Contributor

Choose a reason for hiding this comment

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

JFYI: I think f should have been made an argument for the change to be a no-op (it was previously a kernel argument because of the lambda capture), but it likely doesn't matter here.

int f = 5;

Q.single_task([=]() { volatile int ReadVal = empty[f]; });
return 0;
volatile int ReadVal = empty[f];
Copy link
Contributor

Choose a reason for hiding this comment

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

I think the purpose of this volatile int was to prevent the optimizer from removing everything. In the new approach the idiomatic way would be to just return that int value, e.g.

SYCL_EXTERNAL auto /* or int, maybe */ fpga_datapath_global(int f) { return emtpy[f]; }

}

// CHECK: ![[empty_md]] = !{![[register:[0-9]*]]}
Expand Down
15 changes: 5 additions & 10 deletions sycl/test/check_device_code/fpga_datapath_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,15 +10,10 @@ namespace intel = sycl::ext::intel::experimental; // for fpga_datapath

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

int main() {
queue Q;
SYCL_EXTERNAL void fpga_datapath_local() {
int f = 5;

Q.single_task([=]() {
intel::fpga_datapath<int[10]> empty;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[RegisterINTEL]]
// CHECK-NOT: call void @llvm.memset
volatile int ReadVal = empty[f];
});
return 0;
intel::fpga_datapath<int[10]> empty;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[RegisterINTEL]]
// CHECK-NOT: call void @llvm.memset
volatile int ReadVal = empty[f];
}
13 changes: 3 additions & 10 deletions sycl/test/check_device_code/fpga_mem_constructor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,14 +29,7 @@ class foo {

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

int main() {
queue Q;
int f = 5;

Q.single_task([=]() {
intel::fpga_mem<foo> mem{42};

volatile int ReadVal = mem.get().secret;
});
return 0;
SYCL_EXTERNAL void fpga_mem_constructor() {
intel::fpga_mem<foo> mem{42};
volatile int ReadVal = mem.get().secret;
}
16 changes: 5 additions & 11 deletions sycl/test/check_device_code/fpga_mem_global.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,18 +65,12 @@ const intel::fpga_mem<int[10],
// CHECK: {{.*}}copies = internal addrspace(1) constant {{.*}} zeroinitializer, align 4, !spirv.Decorations ![[copies_md:[0-9]*]]
// CHECK: {{.*}}replicates = internal addrspace(1) constant {{.*}} zeroinitializer, align 4, !spirv.Decorations ![[replicates_md:[0-9]*]]

int main() {
queue Q;
SYCL_EXTERNAL void fpga_mem_global() {
int f = 5;

Q.single_task([=]() {
volatile int ReadVal = empty[f] + min_ram[f] + max_fmax[f] +
double_pumped[f] + single_pumped[f] + mlab[f] +
simple_dual_port[f] + true_dual_port[f] +
block_ram[f] + banks[f] + stride[f] + word[f] +
copies[f] + replicates[f];
});
return 0;
volatile int ReadVal =
empty[f] + min_ram[f] + max_fmax[f] + double_pumped[f] +
single_pumped[f] + mlab[f] + simple_dual_port[f] + true_dual_port[f] +
block_ram[f] + banks[f] + stride[f] + word[f] + copies[f] + replicates[f];
}

// CHECK: ![[empty_md]] = !{![[mem_default:[0-9]*]]}
Expand Down
142 changes: 67 additions & 75 deletions sycl/test/check_device_code/fpga_mem_local.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,81 +24,73 @@ namespace oneapi = sycl::ext::oneapi::experimental; // for properties
// CHECK: [[MaxPrivateCopiesINTEL:@.*]] = private unnamed_addr addrspace(1) constant [27 x i8] c"{5826:\22DEFAULT\22}{5829:\223\22}\00"
// CHECK: [[MaxReplicatesINTEL:@.*]] = private unnamed_addr addrspace(1) constant [27 x i8] c"{5826:\22DEFAULT\22}{5832:\225\22}\00"

int main() {
queue Q;
SYCL_EXTERNAL void fpga_mem_local() {
int f = 5;
intel::fpga_mem<int[10]> empty;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::ram_stitching_min_ram))>
min_ram;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_FALSE]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::ram_stitching_max_fmax))>
max_fmax;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_TRUE]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::clock_2x_true))>
double_pumped;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[DoublepumpINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::clock_2x_false))>
single_pumped;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SinglepumpINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::resource_mlab))>
mlab;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_mlab]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(
intel::bi_directional_ports_false))>
simple_dual_port;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SimpleDualPortINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(
intel::bi_directional_ports_true))>
true_dual_port;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[TrueDualPortINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::resource_block_ram))>
block_ram;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_block_ram]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::num_banks<4>))>
banks;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[NumbanksINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::stride_size<2>))>
stride;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[StridesizeINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::word_size<8>))>
word;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[WordsizeINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::max_private_copies<3>))>
copies;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxPrivateCopiesINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::num_replicates<5>))>
replicates;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxReplicatesINTEL]]
// CHECK-NOT: call void @llvm.memset

Q.single_task([=]() {
intel::fpga_mem<int[10]> empty;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::ram_stitching_min_ram))>
min_ram;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_FALSE]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::ram_stitching_max_fmax))>
max_fmax;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_TRUE]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::clock_2x_true))>
double_pumped;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[DoublepumpINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::clock_2x_false))>
single_pumped;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SinglepumpINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::resource_mlab))>
mlab;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_mlab]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(
intel::bi_directional_ports_false))>
simple_dual_port;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SimpleDualPortINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(
intel::bi_directional_ports_true))>
true_dual_port;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[TrueDualPortINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::resource_block_ram))>
block_ram;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_block_ram]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::num_banks<4>))>
banks;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[NumbanksINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::stride_size<2>))>
stride;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[StridesizeINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10], decltype(oneapi::properties(intel::word_size<8>))>
word;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[WordsizeINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::max_private_copies<3>))>
copies;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxPrivateCopiesINTEL]]
// CHECK-NOT: call void @llvm.memset
intel::fpga_mem<int[10],
decltype(oneapi::properties(intel::num_replicates<5>))>
replicates;
// CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxReplicatesINTEL]]
// CHECK-NOT: call void @llvm.memset

volatile int ReadVal = empty[f] + min_ram[f] + max_fmax[f] +
double_pumped[f] + single_pumped[f] + mlab[f] +
simple_dual_port[f] + true_dual_port[f] +
block_ram[f] + banks[f] + stride[f] + word[f] +
copies[f] + replicates[f];
});
return 0;
volatile int ReadVal =
empty[f] + min_ram[f] + max_fmax[f] + double_pumped[f] +
single_pumped[f] + mlab[f] + simple_dual_port[f] + true_dual_port[f] +
block_ram[f] + banks[f] + stride[f] + word[f] + copies[f] + replicates[f];
}
23 changes: 7 additions & 16 deletions sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,20 +21,11 @@ int arrayAdd(int *data1, int *data2, int N) {
return ret;
}

int main() {
sycl::queue myQueue;

myQueue.submit([&](sycl::handler &cgh) {
cgh.single_task([=]() {
int d1[kSize], d2[kSize];
task_sequence<arrayAdd,
decltype(properties{pipelined<0>, stall_enable_clusters,
invocation_capacity<1>,
response_capacity<1>})>
arrayAddTask;
arrayAddTask.async(d1, d2, kSize);
});
});
myQueue.wait();
return 0;
SYCL_EXTERNAL void task_sequence_no_explicit_get() {
int d1[kSize], d2[kSize];
task_sequence<arrayAdd, decltype(properties{
pipelined<0>, stall_enable_clusters,
invocation_capacity<1>, response_capacity<1>})>
arrayAddTask;
arrayAddTask.async(d1, d2, kSize);
}
24 changes: 7 additions & 17 deletions sycl/test/check_device_code/usm_pointers.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,20 +30,10 @@

using namespace sycl;

int main() {
sycl::queue queue;
{
queue.submit([&](sycl::handler &cgh) {
cgh.single_task<class check_adress_space>([=]() {
void *Ptr = nullptr;
ext::intel::device_ptr<void> DevPtr(Ptr);
ext::intel::host_ptr<void> HostPtr(Ptr);
global_ptr<void> GlobPtr = global_ptr<void>(DevPtr);
GlobPtr = global_ptr<void>(HostPtr);
});
});
queue.wait();
}

return 0;
}
SYCL_EXTERNAL void usm_pointers() {
void *Ptr = nullptr;
ext::intel::device_ptr<void> DevPtr(Ptr);
ext::intel::host_ptr<void> HostPtr(Ptr);
global_ptr<void> GlobPtr = global_ptr<void>(DevPtr);
GlobPtr = global_ptr<void>(HostPtr);
}
Loading