diff --git a/sycl/test/check_device_code/atomic_fence.cpp b/sycl/test/check_device_code/atomic_fence.cpp index 2d16cace866c2..d59f3bdc7af2f 100644 --- a/sycl/test/check_device_code/atomic_fence.cpp +++ b/sycl/test/check_device_code/atomic_fence.cpp @@ -2,27 +2,20 @@ #include -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); +} \ No newline at end of file diff --git a/sycl/test/check_device_code/device_global_ptr_use.cpp b/sycl/test/check_device_code/device_global_ptr_use.cpp index 23397a7b97036..8804272f3234b 100644 --- a/sycl/test/check_device_code/device_global_ptr_use.cpp +++ b/sycl/test/check_device_code/device_global_ptr_use.cpp @@ -10,11 +10,7 @@ using namespace sycl::ext::oneapi::experimental; const device_global 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; +} \ No newline at end of file diff --git a/sycl/test/check_device_code/fpga_datapath_constructor.cpp b/sycl/test/check_device_code/fpga_datapath_constructor.cpp index c03d4712427f8..4adb3a413fc50 100644 --- a/sycl/test/check_device_code/fpga_datapath_constructor.cpp +++ b/sycl/test/check_device_code/fpga_datapath_constructor.cpp @@ -28,14 +28,7 @@ class foo { // CHECK: call {{.*}}sqrt -int main() { - queue Q; - int f = 5; - - Q.single_task([=]() { - intel::fpga_datapath mem{42}; - - volatile int ReadVal = mem.get().secret; - }); - return 0; -} +SYCL_EXTERNAL void fetch_secret() { + intel::fpga_datapath mem{42}; + volatile int ReadVal = mem.get().secret; +} \ No newline at end of file diff --git a/sycl/test/check_device_code/fpga_datapath_global.cpp b/sycl/test/check_device_code/fpga_datapath_global.cpp index 43ebce7cc1f9f..07952b3e5be4e 100644 --- a/sycl/test/check_device_code/fpga_datapath_global.cpp +++ b/sycl/test/check_device_code/fpga_datapath_global.cpp @@ -13,12 +13,9 @@ const intel::fpga_datapath 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() { int f = 5; - - Q.single_task([=]() { volatile int ReadVal = empty[f]; }); - return 0; + volatile int ReadVal = empty[f]; } // CHECK: ![[empty_md]] = !{![[register:[0-9]*]]} diff --git a/sycl/test/check_device_code/fpga_datapath_local.cpp b/sycl/test/check_device_code/fpga_datapath_local.cpp index 901ef4441b7aa..8cf07c3da7af5 100644 --- a/sycl/test/check_device_code/fpga_datapath_local.cpp +++ b/sycl/test/check_device_code/fpga_datapath_local.cpp @@ -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 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 empty; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[RegisterINTEL]] + // CHECK-NOT: call void @llvm.memset + volatile int ReadVal = empty[f]; } diff --git a/sycl/test/check_device_code/fpga_mem_constructor.cpp b/sycl/test/check_device_code/fpga_mem_constructor.cpp index ececf587ece0f..2f9b5c81d7545 100644 --- a/sycl/test/check_device_code/fpga_mem_constructor.cpp +++ b/sycl/test/check_device_code/fpga_mem_constructor.cpp @@ -29,14 +29,7 @@ class foo { // CHECK: call {{.*}}sqrt -int main() { - queue Q; - int f = 5; - - Q.single_task([=]() { - intel::fpga_mem mem{42}; - - volatile int ReadVal = mem.get().secret; - }); - return 0; +SYCL_EXTERNAL void fpga_mem_constructor() { + intel::fpga_mem mem{42}; + volatile int ReadVal = mem.get().secret; } diff --git a/sycl/test/check_device_code/fpga_mem_global.cpp b/sycl/test/check_device_code/fpga_mem_global.cpp index 24c7c12815f1f..c30c14cb7ef9a 100644 --- a/sycl/test/check_device_code/fpga_mem_global.cpp +++ b/sycl/test/check_device_code/fpga_mem_global.cpp @@ -65,18 +65,12 @@ const intel::fpga_mem empty; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem + min_ram; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_FALSE]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem + max_fmax; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_TRUE]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem + double_pumped; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[DoublepumpINTEL]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem + single_pumped; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SinglepumpINTEL]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem + mlab; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_mlab]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem + simple_dual_port; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SimpleDualPortINTEL]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem + true_dual_port; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[TrueDualPortINTEL]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem + block_ram; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_block_ram]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem))> + banks; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[NumbanksINTEL]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem))> + stride; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[StridesizeINTEL]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem))> + word; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[WordsizeINTEL]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem))> + copies; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxPrivateCopiesINTEL]] + // CHECK-NOT: call void @llvm.memset + intel::fpga_mem))> + replicates; + // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxReplicatesINTEL]] + // CHECK-NOT: call void @llvm.memset - Q.single_task([=]() { - intel::fpga_mem empty; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem - min_ram; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_FALSE]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem - max_fmax; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[ForcePow2DepthINTEL_TRUE]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem - double_pumped; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[DoublepumpINTEL]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem - single_pumped; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SinglepumpINTEL]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem - mlab; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_mlab]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem - simple_dual_port; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[SimpleDualPortINTEL]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem - true_dual_port; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[TrueDualPortINTEL]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem - block_ram; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MemoryINTEL_block_ram]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem))> - banks; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[NumbanksINTEL]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem))> - stride; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[StridesizeINTEL]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem))> - word; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[WordsizeINTEL]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem))> - copies; - // CHECK: @llvm.ptr.annotation{{.*}}(ptr addrspace(4) {{.*}}, ptr addrspace(1) [[MaxPrivateCopiesINTEL]] - // CHECK-NOT: call void @llvm.memset - intel::fpga_mem))> - 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]; } diff --git a/sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp b/sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp index d5c4ba65a846e..bd208f9e0e0e1 100644 --- a/sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp +++ b/sycl/test/check_device_code/task_sequence_intel_no_explicit_get.cpp @@ -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, 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, stall_enable_clusters, + invocation_capacity<1>, response_capacity<1>})> + arrayAddTask; + arrayAddTask.async(d1, d2, kSize); } \ No newline at end of file diff --git a/sycl/test/check_device_code/usm_pointers.cpp b/sycl/test/check_device_code/usm_pointers.cpp index b8af6c4eefebc..28127a9187250 100644 --- a/sycl/test/check_device_code/usm_pointers.cpp +++ b/sycl/test/check_device_code/usm_pointers.cpp @@ -30,20 +30,10 @@ using namespace sycl; -int main() { - sycl::queue queue; - { - queue.submit([&](sycl::handler &cgh) { - cgh.single_task([=]() { - void *Ptr = nullptr; - ext::intel::device_ptr DevPtr(Ptr); - ext::intel::host_ptr HostPtr(Ptr); - global_ptr GlobPtr = global_ptr(DevPtr); - GlobPtr = global_ptr(HostPtr); - }); - }); - queue.wait(); - } - - return 0; -} +SYCL_EXTERNAL void usm_pointers() { + void *Ptr = nullptr; + ext::intel::device_ptr DevPtr(Ptr); + ext::intel::host_ptr HostPtr(Ptr); + global_ptr GlobPtr = global_ptr(DevPtr); + GlobPtr = global_ptr(HostPtr); +} \ No newline at end of file