[DO NOT MERGE] CI Testing

sycl/test/group-algorithm/broadcast.cpp

@@ -15,14 +15,11 @@
 using namespace sycl;
 using namespace sycl::ONEAPI;
 
-template <typename InputContainer, typename OutputContainer>
-class broadcast_kernel;
-
-template <typename InputContainer, typename OutputContainer>
+template <typename kernel_name, typename InputContainer,
+          typename OutputContainer>
 void test(queue q, InputContainer input, OutputContainer output) {
   typedef typename InputContainer::value_type InputT;
   typedef typename OutputContainer::value_type OutputT;
-  typedef class broadcast_kernel<InputContainer, OutputContainer> kernel_name;
   size_t N = input.size();
   size_t G = 4;
   range<2> R(G, G);
@@ -63,7 +60,7 @@ int main() {
     std::array<int, 3> output;
     std::iota(input.begin(), input.end(), 1);
     std::fill(output.begin(), output.end(), false);
-    test(q, input, output);
+    test<class KernelName_EFL>(q, input, output);
   }
 
   // Test pointer type
@@ -74,7 +71,7 @@ int main() {
       input[i] = static_cast<int *>(0x0) + i;
     }
     std::fill(output.begin(), output.end(), static_cast<int *>(0x0));
-    test(q, input, output);
+    test<class KernelName_NrqELzFQToOSPsRNMi>(q, input, output);
   }
 
   // Test user-defined type
@@ -88,7 +85,7 @@ int main() {
           std::complex<float>(0, 1) + (float)i * std::complex<float>(2, 2);
     }
     std::fill(output.begin(), output.end(), std::complex<float>(0, 0));
-    test(q, input, output);
+    test<class KernelName_rCblcml>(q, input, output);
   }
   {
     std::array<std::complex<double>, N> input;
@@ -98,7 +95,7 @@ int main() {
           std::complex<double>(0, 1) + (double)i * std::complex<double>(2, 2);
     }
     std::fill(output.begin(), output.end(), std::complex<float>(0, 0));
-    test(q, input, output);
+    test<class KernelName_NCWhjnQ>(q, input, output);
   }
   std::cout << "Test passed." << std::endl;
 }

sycl/test/group-algorithm/exclusive_scan.cpp

@@ -24,7 +24,7 @@
 using namespace sycl;
 using namespace sycl::ONEAPI;
 
-template <class BinaryOperation, int TestNumber>
+template <class SpecializationKernelName, int TestNumber>
 class exclusive_scan_kernel;
 
 // std::exclusive_scan isn't implemented yet, so use serial implementation
@@ -44,17 +44,17 @@ OutputIterator exclusive_scan(InputIterator first, InputIterator last,
 }
 } // namespace emu
 
-template <typename InputContainer, typename OutputContainer,
-          class BinaryOperation>
+template <typename SpecializationKernelName, typename InputContainer,
+          typename OutputContainer, class BinaryOperation>
 void test(queue q, InputContainer input, OutputContainer output,
           BinaryOperation binary_op,
           typename OutputContainer::value_type identity) {
   typedef typename InputContainer::value_type InputT;
   typedef typename OutputContainer::value_type OutputT;
-  typedef class exclusive_scan_kernel<BinaryOperation, 0> kernel_name0;
-  typedef class exclusive_scan_kernel<BinaryOperation, 1> kernel_name1;
-  typedef class exclusive_scan_kernel<BinaryOperation, 2> kernel_name2;
-  typedef class exclusive_scan_kernel<BinaryOperation, 3> kernel_name3;
+  typedef class exclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
+  typedef class exclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
+  typedef class exclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
+  typedef class exclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
   OutputT init = 42;
   size_t N = input.size();
   size_t G = 16;
@@ -159,19 +159,24 @@ int main() {
   std::iota(input.begin(), input.end(), 0);
   std::fill(output.begin(), output.end(), 0);
 
-  test(q, input, output, plus<>(), 0);
-  test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
+  test<class KernelNameMinimumV>(q, input, output, minimum<>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumV>(q, input, output, maximum<>(),
+                                 std::numeric_limits<int>::lowest());
 
-  test(q, input, output, plus<int>(), 0);
-  test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
+  test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
+                                 std::numeric_limits<int>::lowest());
 
 #ifdef SPIRV_1_3
-  test(q, input, output, multiplies<int>(), 1);
-  test(q, input, output, bit_or<int>(), 0);
-  test(q, input, output, bit_xor<int>(), 0);
-  test(q, input, output, bit_and<int>(), ~0);
+  test<class KernelName_VzAPutpBRRJrQPB>(q, input, output, multiplies<int>(),
+                                         1);
+  test<class KernelName_UXdGbr>(q, input, output, bit_or<int>(), 0);
+  test<class KernelName_saYaodNyJknrPW>(q, input, output, bit_xor<int>(), 0);
+  test<class KernelName_GPcuAlvAOjrDyP>(q, input, output, bit_and<int>(), ~0);
 #endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;

sycl/test/group-algorithm/inclusive_scan.cpp

@@ -24,7 +24,7 @@
 using namespace sycl;
 using namespace sycl::ONEAPI;
 
-template <class BinaryOperation, int TestNumber>
+template <class SpecializationKernelName, int TestNumber>
 class inclusive_scan_kernel;
 
 // std::inclusive_scan isn't implemented yet, so use serial implementation
@@ -44,17 +44,17 @@ OutputIterator inclusive_scan(InputIterator first, InputIterator last,
 }
 } // namespace emu
 
-template <typename InputContainer, typename OutputContainer,
-          class BinaryOperation>
+template <typename SpecializationKernelName, typename InputContainer,
+          typename OutputContainer, class BinaryOperation>
 void test(queue q, InputContainer input, OutputContainer output,
           BinaryOperation binary_op,
           typename OutputContainer::value_type identity) {
   typedef typename InputContainer::value_type InputT;
   typedef typename OutputContainer::value_type OutputT;
-  typedef class inclusive_scan_kernel<BinaryOperation, 0> kernel_name0;
-  typedef class inclusive_scan_kernel<BinaryOperation, 1> kernel_name1;
-  typedef class inclusive_scan_kernel<BinaryOperation, 2> kernel_name2;
-  typedef class inclusive_scan_kernel<BinaryOperation, 3> kernel_name3;
+  typedef class inclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
+  typedef class inclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
+  typedef class inclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
+  typedef class inclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
   OutputT init = 42;
   size_t N = input.size();
   size_t G = 16;
@@ -159,19 +159,25 @@ int main() {
   std::iota(input.begin(), input.end(), 0);
   std::fill(output.begin(), output.end(), 0);
 
-  test(q, input, output, plus<>(), 0);
-  test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
+  test<class KernelNameMinimumV>(q, input, output, minimum<>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumV>(q, input, output, maximum<>(),
+                                 std::numeric_limits<int>::lowest());
 
-  test(q, input, output, plus<int>(), 0);
-  test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
+  test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
+                                 std::numeric_limits<int>::lowest());
 
 #ifdef SPIRV_1_3
-  test(q, input, output, multiplies<int>(), 1);
-  test(q, input, output, bit_or<int>(), 0);
-  test(q, input, output, bit_xor<int>(), 0);
-  test(q, input, output, bit_and<int>(), ~0);
+  test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(q, input, output,
+                                              multiplies<int>(), 1);
+  test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output, bit_or<int>(), 0);
+  test<class KernelName_yXIZfjwjxQGiPeQAnc>(q, input, output, bit_xor<int>(),
+                                            0);
+  test<class KernelName_xGnAnMYHvqekCk>(q, input, output, bit_and<int>(), ~0);
 #endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;

sycl/test/group-algorithm/reduce.cpp

@@ -23,17 +23,13 @@
 using namespace sycl;
 using namespace sycl::ONEAPI;
 
-template <class BinaryOperation>
-class reduce_kernel;
-
-template <typename InputContainer, typename OutputContainer,
-          class BinaryOperation>
+template <typename SpecializationKernelName, typename InputContainer,
+          typename OutputContainer, class BinaryOperation>
 void test(queue q, InputContainer input, OutputContainer output,
           BinaryOperation binary_op,
           typename OutputContainer::value_type identity) {
   typedef typename InputContainer::value_type InputT;
   typedef typename OutputContainer::value_type OutputT;
-  typedef class reduce_kernel<BinaryOperation> kernel_name;
   OutputT init = 42;
   size_t N = input.size();
   size_t G = 16;
@@ -44,15 +40,17 @@ void test(queue q, InputContainer input, OutputContainer output,
     q.submit([&](handler &cgh) {
       auto in = in_buf.template get_access<access::mode::read>(cgh);
       auto out = out_buf.template get_access<access::mode::discard_write>(cgh);
-      cgh.parallel_for<kernel_name>(nd_range<1>(G, G), [=](nd_item<1> it) {
-        group<1> g = it.get_group();
-        int lid = it.get_local_id(0);
-        out[0] = reduce(g, in[lid], binary_op);
-        out[1] = reduce(g, in[lid], init, binary_op);
-        out[2] = reduce(g, in.get_pointer(), in.get_pointer() + N, binary_op);
-        out[3] =
-            reduce(g, in.get_pointer(), in.get_pointer() + N, init, binary_op);
-      });
+      cgh.parallel_for<SpecializationKernelName>(
+          nd_range<1>(G, G), [=](nd_item<1> it) {
+            group<1> g = it.get_group();
+            int lid = it.get_local_id(0);
+            out[0] = reduce(g, in[lid], binary_op);
+            out[1] = reduce(g, in[lid], init, binary_op);
+            out[2] =
+                reduce(g, in.get_pointer(), in.get_pointer() + N, binary_op);
+            out[3] = reduce(g, in.get_pointer(), in.get_pointer() + N, init,
+                            binary_op);
+          });
     });
   }
   // std::reduce is not implemented yet, so use std::accumulate instead
@@ -97,19 +95,24 @@ int main() {
   std::iota(input.begin(), input.end(), 0);
   std::fill(output.begin(), output.end(), 0);
 
-  test(q, input, output, plus<>(), 0);
-  test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
+  test<class KernelNameMinimumV>(q, input, output, minimum<>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumV>(q, input, output, maximum<>(),
+                                 std::numeric_limits<int>::lowest());
 
-  test(q, input, output, plus<int>(), 0);
-  test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
+  test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
+                                 std::numeric_limits<int>::lowest());
 
 #ifdef SPIRV_1_3
-  test(q, input, output, multiplies<int>(), 1);
-  test(q, input, output, bit_or<int>(), 0);
-  test(q, input, output, bit_xor<int>(), 0);
-  test(q, input, output, bit_and<int>(), ~0);
+  test<class KernelName_WonwuUVPUPOTKRKIBtT>(q, input, output,
+                                             multiplies<int>(), 1);
+  test<class KernelName_qYBaJDZTMGkdIwD>(q, input, output, bit_or<int>(), 0);
+  test<class KernelName_eLSFt>(q, input, output, bit_xor<int>(), 0);
+  test<class KernelName_uFhJnxSVhNAiFPTG>(q, input, output, bit_and<int>(), ~0);
 #endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;