[SYCL] Change some check_device_code tests to use SYCL_EXTERNAL (inte…

…l#13899) Changed some of the simpler test cases to use SYCL_EXTERNAL functions instead of submitting to a queue.
hdelan · May 27, 2024 · 309b167 · 309b167
1 parent 601f121
commit 309b167
Show file tree

Hide file tree

Showing 10 changed files with 121 additions and 187 deletions.
diff --git a/sycl/test/check_device_code/atomic_fence.cpp b/sycl/test/check_device_code/atomic_fence.cpp
@@ -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);
+}
diff --git 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<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;
+}
diff --git 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<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;
+}
diff --git 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<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() {
   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
@@ -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];
 }
diff --git 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<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;
 }
diff --git 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<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]*]]}

diff --git a/sycl/test/check_device_code/fpga_mem_local.cpp b/sycl/test/check_device_code/fpga_mem_local.cpp
@@ -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];
 }
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
@@ -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);
 }
diff --git 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<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);
+}