Merge pull request #783 from steffenlarsen/steffen/fix_sub_group_mask…

…_marray

Fix sub_group_mask_extract_bits expectation and comparison

tests/extension/oneapi_sub_group_mask/sub_group_mask_extract_bits.cpp

@@ -49,16 +49,24 @@ void get_expected_bits(T &out, uint32_t mask_size, int pos) {
   }
 }
 
+// For marray the extracted bits start from the first element and moves position
+// into the consecutive elements.
+template <typename T, size_t N>
+void get_expected_bits(sycl::marray<T, N>& out, uint32_t mask_size, int pos) {
+  for (size_t i = 0; i < N; ++i)
+    get_expected_bits(out[i], mask_size, pos + i * sizeof(T) * CHAR_BIT);
+}
+
 template <typename T>
 struct check_result_extract_bits {
   bool operator()(const sycl::ext::oneapi::sub_group_mask sub_group_mask,
-                  const sycl::sub_group &) {
+                  const sycl::sub_group&) {
     for (size_t pos = 0; pos <= sub_group_mask.size(); pos++) {
       T bits;
       sub_group_mask.extract_bits(bits, sycl::id(pos));
-      T expected(0);
+      T expected = value_operations::init<T>(0);
       get_expected_bits(expected, sub_group_mask.size(), pos);
-      if (bits != expected) return false;
+      if (!value_operations::are_equal(bits, expected)) return false;
     }
     return true;
   }