EXSWHTEC-256 - Implement tests for grid_group APIs (#153)

- Migrate basic Cooperative Groups tests and integrate to catch
- Refactor basic Cooperative Groups tests
- Rename tiled partition related files and fix minor bug
- Add LaunchCooperativeKernal and LaunchCooperativeKernelMultiDevice tests
- Refactor hipCGThreadBlockTileType to use common function
- Fix updated file not added during merge
- Add coalesced_group type tests
- Add coalesced_group shuffle_up and shuffle_down tests
- Add coalesced_group shuffle tests - test fails
- Implement common code for cooperative group tests
- Fixed compilation errror in cooperative_groups_common.hh
- Implement busy wait device function
- Reimplement tests for grid_group APIs
- Add tests for grid_group member and non-member APIs
- Refactor existing test for grid_group sync testing
- Add thread and block dimensions generators
- Add check of grid and block dimensions
- Modify doxygen comments
- Move cpu_grid.h and supporting functions to catch/include
- Use warp_size from properties in grid/block dims generators
- Fix condition for warp size 32 on AMD
- Fix cpu_grid.h for warp function tests
- Add missing include into cpu_grid.h
- Code cleanup
- Fix doxygen comments
- Add missing include in utils header

catch/include/utils.hh

@@ -20,6 +20,7 @@ THE SOFTWARE.
 #pragma once
 
 #include <chrono>
+#include <optional>
 
 #include <hip_test_common.hh>
 #include <hip/hip_runtime_api.h>
@@ -54,6 +55,20 @@ void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_ele
   ArrayFindIfNot(array, array + num_elements, expected_value);
 }
 
+template <typename T, typename F>
+static inline void ArrayAllOf(const T* arr, uint32_t count, F value_gen) {
+  for (auto i = 0u; i < count; ++i) {
+    const std::optional<T> expected_val = value_gen(i);
+    if (!expected_val.has_value()) continue;
+    // Using require on every iteration leads to a noticeable performance loss on large arrays,
+    // even when the require passes.
+    if (arr[i] != expected_val.value()) {
+      INFO("Mismatch at index: " << i);
+      REQUIRE(arr[i] == expected_val.value());
+    }
+  }
+}
+
 template <typename T, typename F>
 void PitchedMemoryVerify(T* const ptr, const size_t pitch, const size_t width, const size_t height,
                          const size_t depth, F expected_value_generator) {