Merge pull request #1157 from bader/assert

Hoist assert from the device code.

Author: tomdeakin

tests/common/semantics_reference.h

@@ -466,9 +466,10 @@ void check_kernel(InitFunc init_func, const std::string& type_name,
   std::size_t result_count = test_traits::result_count +
                              test_copy::result_count + test_move::result_count;
 
-  std::vector<int> results(result_count, false);
+  std::size_t result_size = result_count + 1;
+  std::vector<int> results(result_size, false);
   {
-    sycl::buffer<int> buffer(results.data(), sycl::range<1>{result_count});
+    sycl::buffer<int> buffer(results.data(), sycl::range<1>{result_size});
 
     queue.submit([&](sycl::handler& cgh) {
       auto accessor = buffer.template get_access<sycl::access_mode::write>(cgh);
@@ -484,8 +485,8 @@ void check_kernel(InitFunc init_func, const std::string& type_name,
             ptr += test_copy::result_count;
             test_move::run<storage>(ptr, t);  // last since invalidates instance
             ptr += test_move::result_count;
-            assert(static_cast<std::ptrdiff_t>(result_count) ==
-                   ptr - accessor.begin());
+            *ptr = static_cast<std::ptrdiff_t>(result_count) ==
+                   ptr - accessor.begin();
           });
     });
   }
@@ -497,7 +498,7 @@ void check_kernel(InitFunc init_func, const std::string& type_name,
   ptr += test_copy::result_count;
   test_move::evaluate(ptr);
   ptr += test_move::result_count;
-  assert(static_cast<std::ptrdiff_t>(result_count) == ptr - results.data());
+  assert(*ptr);
 }
 
 }  // namespace common_reference_semantics

tests/group_functions/group_scan.h

@@ -38,8 +38,6 @@ auto joint_inclusive_scan_helper(Group group, T* v_begin, T* v_end,
     return sycl::joint_inclusive_scan(group, v_begin, v_end, r_i_begin, op,
                                       I(init));
   }
-  assert((std::is_same_v<I, U> &&
-          "Without init value I and U should be the same type."));
   return (U*)sycl::joint_inclusive_scan(group, v_begin, v_end, (I*)r_i_begin,
                                         op);
 }
@@ -51,8 +49,6 @@ auto joint_exclusive_scan_helper(Group group, T* v_begin, T* v_end,
     return sycl::joint_exclusive_scan(group, v_begin, v_end, r_e_begin, I(init),
                                       op);
   }
-  assert((std::is_same_v<I, U> &&
-          "Without init value I and U should be the same type."));
   return (U*)sycl::joint_exclusive_scan(group, v_begin, v_end, (I*)r_e_begin,
                                         op);
 }
@@ -149,6 +145,9 @@ void check_scan(sycl::queue& queue, size_t size,
     sycl::buffer<bool, 1> end_sycl = host_data.create_end_buffer();
     sycl::buffer<bool, 1> ret_type_sycl = host_data.create_ret_type_buffer();
 
+    assert((with_init || std::is_same_v<I, U>) &&
+           "Without init value I and U must be the same type.");
+
     queue
         .submit([&](sycl::handler& cgh) {
           sycl::accessor<T, 1> ref_input_acc(ref_input_sycl, cgh);
@@ -315,8 +314,6 @@ auto inclusive_scan_over_group_helper(Group group, U x, OpT op,
   if (with_init) {
     return sycl::inclusive_scan_over_group(group, x, op, T(init));
   }
-  assert((std::is_same_v<T, U> &&
-          "Without init value T and U should be the same type."));
   return sycl::inclusive_scan_over_group(group, T(x), op);
 }
 
@@ -326,8 +323,6 @@ auto exclusive_scan_over_group_helper(Group group, U x, OpT op,
   if (with_init) {
     return sycl::exclusive_scan_over_group(group, x, T(init), op);
   }
-  assert((std::is_same_v<T, U> &&
-          "Without init value T and U should be the same type."));
   return sycl::exclusive_scan_over_group(group, T(x), op);
 }
 
@@ -457,6 +452,9 @@ void check_scan_over_group(sycl::queue& queue, sycl::range<D> range, OpT op,
     auto local_id_sycl = host_data.create_local_id_buffer();
     auto sub_group_id_sycl = host_data.create_sub_group_id_buffer();
 
+    assert((with_init || std::is_same_v<T, U>) &&
+           "Without init value T and U must be the same type.");
+
     queue
         .submit([&](sycl::handler& cgh) {
           sycl::accessor<U, 1, sycl::access_mode::read> ref_input_acc(