SWDEV-256723 - Added group partitioning feature in ROCm CG.

Change-Id: Ie54046feef3baba857a7068972ec1fc0a60c2df9


[ROCm/clr commit: 0a0b026304]

projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups.h

@@ -39,15 +39,15 @@ namespace cooperative_groups {
 
 /** \brief The base type of all cooperative group types
  *
- *  \details Holds the key properties of a constructed cooperative group type
+ *  \details Holds the key properties of a constructed cooperative group types
  *           object, like the group type, its size, etc
  */
 class thread_group {
  protected:
-  uint32_t _type; // thread_group type
-  uint32_t _size; // total number of threads in the tread_group
-  uint64_t _mask; // Lanemask for coalesced and tiled partitioned group types,
-                  // LSB represents lane 0, and MSB represents lane 63
+  uint32_t _type;  // thread_group type
+  uint32_t _size;  // total number of threads in the tread_group
+  uint64_t _mask;  // Lanemask for coalesced and tiled partitioned group types,
+                   // LSB represents lane 0, and MSB represents lane 63
 
   // Construct a thread group, and set thread group type and other essential
   // thread group properties. This generic thread group is directly constructed
@@ -61,13 +61,21 @@ class thread_group {
     _mask = mask;
   }
 
+  struct _tiled_info {
+    bool is_tiled;
+    unsigned int size;
+  } tiled_info;
+
+  friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent,
+                                                       unsigned int tile_size);
+  friend class thread_block;
+
  public:
   // Total number of threads in the thread group, and this serves the purpose
   // for all derived cooperative group types since their `size` is directly
   // saved during the construction
-  __CG_QUALIFIER__ uint32_t size() const {
-    return _size;
-  }
+  __CG_QUALIFIER__ uint32_t size() const { return _size; }
+  __CG_QUALIFIER__ unsigned int cg_type() const { return _type; }
   // Rank of the calling thread within [0, size())
   __CG_QUALIFIER__ uint32_t thread_rank() const;
   // Is this cooperative group type valid?
@@ -90,28 +98,18 @@ class multi_grid_group : public thread_group {
  protected:
   // Construct mutli-grid thread group (through the API this_multi_grid())
   explicit __CG_QUALIFIER__ multi_grid_group(uint32_t size)
-      : thread_group(internal::cg_multi_grid, size) { }
+      : thread_group(internal::cg_multi_grid, size) {}
 
  public:
   // Number of invocations participating in this multi-grid group. In other
   // words, the number of GPUs
-	__CG_QUALIFIER__ uint32_t num_grids() {
-    return internal::multi_grid::num_grids();
-  }
+  __CG_QUALIFIER__ uint32_t num_grids() { return internal::multi_grid::num_grids(); }
   // Rank of this invocation. In other words, an ID number within the range
   // [0, num_grids()) of the GPU, this kernel is running on
-	__CG_QUALIFIER__ uint32_t grid_rank() {
-    return internal::multi_grid::grid_rank();
-  }
-  __CG_QUALIFIER__ uint32_t thread_rank() const {
-    return internal::multi_grid::thread_rank();
-  }
-  __CG_QUALIFIER__ bool is_valid() const {
-    return internal::multi_grid::is_valid();
-  }
-  __CG_QUALIFIER__ void sync() const {
-    internal::multi_grid::sync();
-  }
+  __CG_QUALIFIER__ uint32_t grid_rank() { return internal::multi_grid::grid_rank(); }
+  __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::multi_grid::thread_rank(); }
+  __CG_QUALIFIER__ bool is_valid() const { return internal::multi_grid::is_valid(); }
+  __CG_QUALIFIER__ void sync() const { internal::multi_grid::sync(); }
 };
 
 /** \brief User exposed API interface to construct multi-grid cooperative
@@ -121,8 +119,7 @@ class multi_grid_group : public thread_group {
  *           `multi_grid_group`. Instead, he should construct it through this
  *           API function
  */
-__CG_QUALIFIER__ multi_grid_group
-this_multi_grid() {
+__CG_QUALIFIER__ multi_grid_group this_multi_grid() {
   return multi_grid_group(internal::multi_grid::size());
 }
 
@@ -139,19 +136,12 @@ class grid_group : public thread_group {
 
  protected:
   // Construct grid thread group (through the API this_grid())
-  explicit __CG_QUALIFIER__ grid_group(uint32_t size)
-      : thread_group(internal::cg_grid, size) { }
+  explicit __CG_QUALIFIER__ grid_group(uint32_t size) : thread_group(internal::cg_grid, size) {}
 
  public:
-  __CG_QUALIFIER__ uint32_t thread_rank() const {
-    return internal::grid::thread_rank();
-  }
-  __CG_QUALIFIER__ bool is_valid() const {
-    return internal::grid::is_valid();
-  }
-  __CG_QUALIFIER__ void sync() const {
-    internal::grid::sync();
-  }
+  __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::grid::thread_rank(); }
+  __CG_QUALIFIER__ bool is_valid() const { return internal::grid::is_valid(); }
+  __CG_QUALIFIER__ void sync() const { internal::grid::sync(); }
 };
 
 /** \brief User exposed API interface to construct grid cooperative group type
@@ -161,60 +151,112 @@ class grid_group : public thread_group {
  *           `multi_grid_group`. Instead, he should construct it through this
  *           API function
  */
-__CG_QUALIFIER__ grid_group
-this_grid() {
-  return grid_group(internal::grid::size());
-}
+__CG_QUALIFIER__ grid_group this_grid() { return grid_group(internal::grid::size()); }
 
-/** \brief The workgroup (thread-block in CUDA terminology) cooperative group
- *         type
+/** \brief   The workgroup (thread-block in CUDA terminology) cooperative group
+ *           type
  *
  *  \details Represents an intra-workgroup cooperative group type where the
- *           participating threads within the group are exctly the same threads
+ *           participating threads within the group are exactly the same threads
  *           which are participated in the currently executing `workgroup`
  */
 class thread_block : public thread_group {
-  // Only these friend functions are allowed to construct an object of this
+  // Only these friend functions are allowed to construct an object of thi
   // class and access its resources
   friend __CG_QUALIFIER__ thread_block this_thread_block();
+  friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent,
+                                                       unsigned int tile_size);
+  friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_block& parent,
+                                                       unsigned int tile_size);
 
  protected:
   // Construct a workgroup thread group (through the API this_thread_block())
   explicit __CG_QUALIFIER__ thread_block(uint32_t size)
-      : thread_group(internal::cg_workgroup, size) { }
+      : thread_group(internal::cg_workgroup, size) {}
+
+  __CG_QUALIFIER__ thread_group new_tiled_group(unsigned int tile_size) const {
+    const bool pow2 = ((tile_size & (tile_size - 1)) == 0);
+    // Invalid tile size, assert
+    if (!tile_size || (tile_size > WAVEFRONT_SIZE) || !pow2) {
+      assert(false && "invalid tile size");
+    }
+
+    thread_group tiledGroup = thread_group(internal::cg_tiled_group, tile_size);
+    tiledGroup.tiled_info.size = tile_size;
+    tiledGroup.tiled_info.is_tiled = true;
+    return tiledGroup;
+  }
 
  public:
   // 3-dimensional block index within the grid
-  __CG_QUALIFIER__ dim3 group_index() {
-    return internal::workgroup::group_index();
-  }
+  __CG_QUALIFIER__ dim3 group_index() { return internal::workgroup::group_index(); }
   // 3-dimensional thread index within the block
-  __CG_QUALIFIER__ dim3 thread_index() {
-    return internal::workgroup::thread_index();
-  }
-  __CG_QUALIFIER__ uint32_t thread_rank() const {
-    return internal::workgroup::thread_rank();
-  }
-  __CG_QUALIFIER__ bool is_valid() const {
-    return internal::workgroup::is_valid();
-  }
-  __CG_QUALIFIER__ void sync() const {
-    internal::workgroup::sync();
-  }
+  __CG_QUALIFIER__ dim3 thread_index() { return internal::workgroup::thread_index(); }
+  __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::workgroup::thread_rank(); }
+  __CG_QUALIFIER__ bool is_valid() const { return internal::workgroup::is_valid(); }
+  __CG_QUALIFIER__ void sync() const { internal::workgroup::sync(); }
 };
 
-/** \brief User exposed API interface to construct workgroup cooperative
- *         group type object - `thread_block`
+/** \brief   User exposed API interface to construct workgroup cooperative
+ *           group type object - `thread_block`.
  *
  *  \details User is not allowed to directly construct an object of type
  *           `thread_block`. Instead, he should construct it through this API
- *           function
+ *           function.
  */
-__CG_QUALIFIER__ thread_block
-this_thread_block() {
+__CG_QUALIFIER__ thread_block this_thread_block() {
   return thread_block(internal::workgroup::size());
 }
 
+/** \brief   The tiled_group cooperative group type
+ *
+ *  \details Represents one tiled thread group in a wavefront.
+ *           This group type also supports sub-wave level intrinsics.
+ */
+
+class tiled_group : public thread_group {
+ private:
+  friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent,
+                                                       unsigned int tile_size);
+  friend __CG_QUALIFIER__ tiled_group tiled_partition(const tiled_group& parent,
+                                                      unsigned int tile_size);
+
+  __CG_QUALIFIER__ tiled_group new_tiled_group(unsigned int tile_size) const {
+    const bool pow2 = ((tile_size & (tile_size - 1)) == 0);
+
+    if (!tile_size || (tile_size > WAVEFRONT_SIZE) || !pow2) {
+      assert(false && "invalid tile size");
+    }
+
+    if (size() <= tile_size) {
+      return (*this);
+    }
+
+    tiled_group tiledGroup = tiled_group(tile_size);
+    tiledGroup.tiled_info.is_tiled = true;
+    return tiledGroup;
+  }
+
+ protected:
+  explicit __CG_QUALIFIER__ tiled_group(unsigned int tileSize)
+      : thread_group(internal::cg_tiled_group, tileSize) {
+    tiled_info.size = tileSize;
+    tiled_info.is_tiled = true;
+  }
+
+ public:
+  __CG_QUALIFIER__ unsigned int size() const { return (tiled_info.size); }
+
+  __CG_QUALIFIER__ unsigned int thread_rank() const {
+    return (internal::workgroup::thread_rank() & (tiled_info.size - 1));
+  }
+
+  __CG_QUALIFIER__ void sync() const {
+    // enforce memory ordering for memory instructions.
+    __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent");
+  }
+};
+
 /**
  *  Implemenation of all publicly exposed base class APIs
  */
@@ -229,6 +271,9 @@ __CG_QUALIFIER__ uint32_t thread_group::thread_rank() const {
     case internal::cg_workgroup: {
       return (static_cast<const thread_block*>(this)->thread_rank());
     }
+    case internal::cg_tiled_group: {
+      return (static_cast<const tiled_group*>(this)->thread_rank());
+    }
     default: {
       assert(false && "invalid cooperative group type");
       return -1;
@@ -247,6 +292,9 @@ __CG_QUALIFIER__ bool thread_group::is_valid() const {
     case internal::cg_workgroup: {
       return (static_cast<const thread_block*>(this)->is_valid());
     }
+    case internal::cg_tiled_group: {
+      return (static_cast<const tiled_group*>(this)->is_valid());
+    }
     default: {
       assert(false && "invalid cooperative group type");
       return false;
@@ -268,6 +316,10 @@ __CG_QUALIFIER__ void thread_group::sync() const {
       static_cast<const thread_block*>(this)->sync();
       break;
     }
+    case internal::cg_tiled_group: {
+      static_cast<const tiled_group*>(this)->sync();
+      break;
+    }
     default: {
       assert(false && "invalid cooperative group type");
     }
@@ -278,27 +330,181 @@ __CG_QUALIFIER__ void thread_group::sync() const {
  *  Implemenation of publicly exposed `wrapper` APIs on top of basic cooperative
  *  group type APIs
  */
-template <class CGTy>
-__CG_QUALIFIER__ uint32_t group_size(CGTy const &g) {
-  return g.size();
-}
+template <class CGTy> __CG_QUALIFIER__ uint32_t group_size(CGTy const& g) { return g.size(); }
 
-template <class CGTy>
-__CG_QUALIFIER__ uint32_t thread_rank(CGTy const &g) {
+template <class CGTy> __CG_QUALIFIER__ uint32_t thread_rank(CGTy const& g) {
   return g.thread_rank();
 }
 
-template <class CGTy>
-__CG_QUALIFIER__ bool is_valid(CGTy const &g) {
-  return g.is_valid();
+template <class CGTy> __CG_QUALIFIER__ bool is_valid(CGTy const& g) { return g.is_valid(); }
+
+template <class CGTy> __CG_QUALIFIER__ void sync(CGTy const& g) { g.sync(); }
+
+template <unsigned int tileSize> class tile_base {
+ protected:
+  _CG_STATIC_CONST_DECL_ unsigned int numThreads = tileSize;
+
+ public:
+  // Rank of the thread within this tile
+  _CG_STATIC_CONST_DECL_ unsigned int thread_rank() {
+    return (internal::workgroup::thread_rank() & (numThreads - 1));
+  }
+
+  // Number of threads within this tile
+  __CG_STATIC_QUALIFIER__ unsigned int size() { return numThreads; }
+};
+
+template <unsigned int size> class thread_block_tile_base : public tile_base<size> {
+  static_assert(is_valid_tile_size<size>::value,
+                "Tile size is either not a power of 2 or greater than the wavefront size");
+  using tile_base<size>::numThreads;
+
+ public:
+  __CG_STATIC_QUALIFIER__ void sync() {
+    // enforce ordering for memory instructions
+    __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent");
+  }
+
+  template <class T> __CG_QUALIFIER__ T shfl(T var, int srcRank) const {
+    static_assert(is_valid_type<T>::value, "Neither an integer or float type.");
+    return (__shfl(var, srcRank, numThreads));
+  }
+
+  template <class T> __CG_QUALIFIER__ T shfl_down(T var, unsigned int lane_delta) const {
+    static_assert(is_valid_type<T>::value, "Neither an integer or float type.");
+    return (__shfl_down(var, lane_delta, numThreads));
+  }
+
+  template <class T> __CG_QUALIFIER__ T shfl_up(T var, unsigned int lane_delta) const {
+    static_assert(is_valid_type<T>::value, "Neither an integer or float type.");
+    return (__shfl_up(var, lane_delta, numThreads));
+  }
+
+  template <class T> __CG_QUALIFIER__ T shfl_xor(T var, unsigned int laneMask) const {
+    static_assert(is_valid_type<T>::value, "Neither an integer or float type.");
+    return (__shfl_xor(var, laneMask, numThreads));
+  }
+};
+
+/** \brief   Group type - thread_block_tile
+ *
+ *  \details  Represents one tile of thread group.
+ */
+
+template <unsigned int tileSize, class ParentCGTy = void>
+class thread_block_tile_type : public thread_block_tile_base<tileSize>, public tiled_group {
+  _CG_STATIC_CONST_DECL_ unsigned int numThreads = tileSize;
+
+  friend class thread_block_tile_type<tileSize, ParentCGTy>;
+
+  typedef thread_block_tile_base<numThreads> tbtBase;
+
+ protected:
+  __CG_QUALIFIER__ thread_block_tile_type() : tiled_group(numThreads) {
+    tiled_info.size = numThreads;
+    tiled_info.is_tiled = true;
+  }
+
+ public:
+  using tbtBase::size;
+  using tbtBase::sync;
+  using tbtBase::thread_rank;
+};
+
+
+/** \brief   User exposed API to partition groups.
+ *
+ *  \details A collective operation that partitions the parent group into a one-dimensional,
+ *           row-major, tiling of subgroups.
+ */
+
+__CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent, unsigned int tile_size) {
+  if (parent.cg_type() == internal::cg_tiled_group) {
+    const tiled_group* cg = static_cast<const tiled_group*>(&parent);
+    return cg->new_tiled_group(tile_size);
+  } else {
+    const thread_block* tb = static_cast<const thread_block*>(&parent);
+    return tb->new_tiled_group(tile_size);
+  }
 }
 
-template <class CGTy>
-__CG_QUALIFIER__ void sync(CGTy const &g) {
-  g.sync();
+// Thread block type overload
+__CG_QUALIFIER__ thread_group tiled_partition(const thread_block& parent, unsigned int tile_size) {
+  return (parent.new_tiled_group(tile_size));
 }
 
-} // namespace cooperative_groups
+// Coalesced group type overload
+__CG_QUALIFIER__ tiled_group tiled_partition(const tiled_group& parent, unsigned int tile_size) {
+  return (parent.new_tiled_group(tile_size));
+}
 
-#endif // __cplusplus
-#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H
+template <unsigned int size, class ParentCGTy> class thread_block_tile;
+
+namespace impl {
+template <unsigned int size, class ParentCGTy> class thread_block_tile_internal;
+
+template <unsigned int size, class ParentCGTy>
+class thread_block_tile_internal : public thread_block_tile_type<size, ParentCGTy> {
+ protected:
+  template <unsigned int tbtSize, class tbtParentT>
+  __CG_QUALIFIER__ thread_block_tile_internal(
+      const thread_block_tile_internal<tbtSize, tbtParentT>& g)
+      : thread_block_tile_type<size, ParentCGTy>() {}
+
+  __CG_QUALIFIER__ thread_block_tile_internal(const thread_block& g)
+      : thread_block_tile_type<size, ParentCGTy>() {}
+};
+}  // namespace impl
+
+template <unsigned int size, class ParentCGTy>
+class thread_block_tile : public impl::thread_block_tile_internal<size, ParentCGTy> {
+ protected:
+  __CG_QUALIFIER__ thread_block_tile(const ParentCGTy& g)
+      : impl::thread_block_tile_internal<size, ParentCGTy>(g) {}
+
+ public:
+  __CG_QUALIFIER__ operator thread_block_tile<size, void>() const {
+    return thread_block_tile<size, void>(*this);
+  }
+};
+
+
+template <unsigned int size>
+class thread_block_tile<size, void> : public impl::thread_block_tile_internal<size, void> {
+  template <unsigned int, class ParentCGTy> friend class thread_block_tile;
+
+ protected:
+ public:
+  template <class ParentCGTy>
+  __CG_QUALIFIER__ thread_block_tile(const thread_block_tile<size, ParentCGTy>& g)
+      : impl::thread_block_tile_internal<size, void>(g) {}
+};
+
+template <unsigned int size, class ParentCGTy = void> class thread_block_tile;
+
+namespace impl {
+template <unsigned int size, class ParentCGTy = void> struct tiled_partition_internal;
+
+template <unsigned int size>
+struct tiled_partition_internal<size, thread_block> : public thread_block_tile<size, thread_block> {
+  __CG_QUALIFIER__ tiled_partition_internal(const thread_block& g)
+      : thread_block_tile<size, thread_block>(g) {}
+};
+
+}  // namespace impl
+
+/** \brief   User exposed API to partition groups.
+ *
+ *  \details  This constructs a templated class derieved from thread_group.
+ *            The template defines tile size of the new thread group at compile time.
+ */
+template <unsigned int size, class ParentCGTy>
+__CG_QUALIFIER__ thread_block_tile<size, ParentCGTy> tiled_partition(const ParentCGTy& g) {
+  static_assert(is_valid_tile_size<size>::value,
+                "Tiled partition with size > wavefront size. Currently not supported ");
+  return impl::tiled_partition_internal<size, ParentCGTy>(g);
+}
+}  // namespace cooperative_groups
+
+#endif  // __cplusplus
+#endif  // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H

projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h

@@ -47,12 +47,34 @@ THE SOFTWARE.
 #define __CG_STATIC_QUALIFIER__ __device__ static __forceinline__
 #endif
 
+#if !defined(_CG_STATIC_CONST_DECL_)
+#define _CG_STATIC_CONST_DECL_ static constexpr
+#endif
+
 #if !defined(WAVEFRONT_SIZE)
+#if __gfx1010__ || __gfx1011__ || __gfx1012__ || __gfx1030__ || __gfx1031__
+#define WAVEFRONT_SIZE 32
+#else
 #define WAVEFRONT_SIZE 64
 #endif
 
 namespace cooperative_groups {
 
+/* Global scope */
+template <unsigned int size>
+using is_power_of_2 = std::integral_constant<bool, (size & (size - 1)) == 0>;
+
+template <unsigned int size>
+using is_valid_wavefront = std::integral_constant<bool, (size <= WAVEFRONT_SIZE)>;
+
+template <unsigned int size>
+using is_valid_tile_size =
+    std::integral_constant<bool, is_power_of_2<size>::value && is_valid_wavefront<size>::value>;
+
+template <typename T>
+using is_valid_type =
+    std::integral_constant<bool, std::is_integral<T>::value || std::is_floating_point<T>::value>;
+
 namespace internal {
 
 /** \brief Enums representing different cooperative group types
@@ -61,7 +83,8 @@ typedef enum {
   cg_invalid,
   cg_multi_grid,
   cg_grid,
-  cg_workgroup
+  cg_workgroup,
+  cg_tiled_group
 } group_type;
 
 /**
@@ -69,31 +92,19 @@ typedef enum {
  */
 namespace multi_grid {
 
-__CG_STATIC_QUALIFIER__ uint32_t num_grids() {
-  return (uint32_t)__ockl_multi_grid_num_grids();
-}
+__CG_STATIC_QUALIFIER__ uint32_t num_grids() { return (uint32_t)__ockl_multi_grid_num_grids(); }
 
-__CG_STATIC_QUALIFIER__ uint32_t grid_rank() {
-  return (uint32_t)__ockl_multi_grid_grid_rank();
-}
+__CG_STATIC_QUALIFIER__ uint32_t grid_rank() { return (uint32_t)__ockl_multi_grid_grid_rank(); }
 
-__CG_STATIC_QUALIFIER__ uint32_t size() {
-  return (uint32_t)__ockl_multi_grid_size();
-}
+__CG_STATIC_QUALIFIER__ uint32_t size() { return (uint32_t)__ockl_multi_grid_size(); }
 
-__CG_STATIC_QUALIFIER__ uint32_t thread_rank() {
-  return (uint32_t)__ockl_multi_grid_thread_rank();
-}
+__CG_STATIC_QUALIFIER__ uint32_t thread_rank() { return (uint32_t)__ockl_multi_grid_thread_rank(); }
 
-__CG_STATIC_QUALIFIER__ bool is_valid() {
-  return (bool)__ockl_multi_grid_is_valid();
-}
+__CG_STATIC_QUALIFIER__ bool is_valid() { return (bool)__ockl_multi_grid_is_valid(); }
 
-__CG_STATIC_QUALIFIER__ void sync() {
-  __ockl_multi_grid_sync();
-}
+__CG_STATIC_QUALIFIER__ void sync() { __ockl_multi_grid_sync(); }
 
-} // namespace multi_grid
+}  // namespace multi_grid
 
 /**
  *  Functionalities related to grid cooperative group type
@@ -101,41 +112,32 @@ __CG_STATIC_QUALIFIER__ void sync() {
 namespace grid {
 
 __CG_STATIC_QUALIFIER__ uint32_t size() {
-  return (uint32_t)((hipBlockDim_z * hipGridDim_z) *
-                    (hipBlockDim_y * hipGridDim_y) *
+  return (uint32_t)((hipBlockDim_z * hipGridDim_z) * (hipBlockDim_y * hipGridDim_y) *
                     (hipBlockDim_x * hipGridDim_x));
 }
 
 __CG_STATIC_QUALIFIER__ uint32_t thread_rank() {
   // Compute global id of the workgroup to which the current thread belongs to
-  uint32_t blkIdx =
-           (uint32_t)((hipBlockIdx_z * hipGridDim_y * hipGridDim_x) +
-                      (hipBlockIdx_y * hipGridDim_x) +
-                      (hipBlockIdx_x));
+  uint32_t blkIdx = (uint32_t)((hipBlockIdx_z * hipGridDim_y * hipGridDim_x) +
+                               (hipBlockIdx_y * hipGridDim_x) + (hipBlockIdx_x));
 
   // Compute total number of threads being passed to reach current workgroup
   // within grid
   uint32_t num_threads_till_current_workgroup =
-           (uint32_t)(blkIdx * (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z));
+      (uint32_t)(blkIdx * (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z));
 
   // Compute thread local rank within current workgroup
-  uint32_t local_thread_rank =
-           (uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) +
-                      (hipThreadIdx_y * hipBlockDim_x) +
-                      (hipThreadIdx_x));
+  uint32_t local_thread_rank = (uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) +
+                                          (hipThreadIdx_y * hipBlockDim_x) + (hipThreadIdx_x));
 
   return (num_threads_till_current_workgroup + local_thread_rank);
 }
 
-__CG_STATIC_QUALIFIER__ bool is_valid() {
-  return (bool)__ockl_grid_is_valid();
-}
+__CG_STATIC_QUALIFIER__ bool is_valid() { return (bool)__ockl_grid_is_valid(); }
 
-__CG_STATIC_QUALIFIER__ void sync() {
-  __ockl_grid_sync();
-}
+__CG_STATIC_QUALIFIER__ void sync() { __ockl_grid_sync(); }
 
-} // namespace grid
+}  // namespace grid
 
 /**
  *  Functionalities related to `workgroup` (thread_block in CUDA terminology)
@@ -144,39 +146,35 @@ __CG_STATIC_QUALIFIER__ void sync() {
 namespace workgroup {
 
 __CG_STATIC_QUALIFIER__ dim3 group_index() {
-  return (dim3((uint32_t)hipBlockIdx_x, (uint32_t)hipBlockIdx_y,
-               (uint32_t)hipBlockIdx_z));
+  return (dim3((uint32_t)hipBlockIdx_x, (uint32_t)hipBlockIdx_y, (uint32_t)hipBlockIdx_z));
 }
 
 __CG_STATIC_QUALIFIER__ dim3 thread_index() {
-  return (dim3((uint32_t)hipThreadIdx_x, (uint32_t)hipThreadIdx_y,
-               (uint32_t)hipThreadIdx_z));
+  return (dim3((uint32_t)hipThreadIdx_x, (uint32_t)hipThreadIdx_y, (uint32_t)hipThreadIdx_z));
 }
 
 __CG_STATIC_QUALIFIER__ uint32_t size() {
-  return((uint32_t)(hipBlockDim_x * hipBlockDim_y * hipBlockDim_z));
+  return ((uint32_t)(hipBlockDim_x * hipBlockDim_y * hipBlockDim_z));
 }
 
 __CG_STATIC_QUALIFIER__ uint32_t thread_rank() {
- return ((uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) +
-                    (hipThreadIdx_y * hipBlockDim_x) +
-                    (hipThreadIdx_x)));
+  return ((uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) +
+                     (hipThreadIdx_y * hipBlockDim_x) + (hipThreadIdx_x)));
 }
 
 __CG_STATIC_QUALIFIER__ bool is_valid() {
-   //TODO(mahesha) any functionality need to be added here? I believe not
+  // TODO(mahesha) any functionality need to be added here? I believe not
   return true;
 }
 
-__CG_STATIC_QUALIFIER__ void sync() {
-  __syncthreads();
-}
+__CG_STATIC_QUALIFIER__ void sync() { __syncthreads(); }
 
-} // namespace workgroup
+}  // namespace workgroup
 
-} // namespace internal
+}  // namespace internal
 
-} // namespace cooperative_groups
+}  // namespace cooperative_groups
 
-#endif // __cplusplus
-#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H
+#endif  // __cplusplus
+#endif  // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_HELPER_H
+#endif

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_tiled_partition.cpp

@@ -0,0 +1,400 @@
+/*
+Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+// Test Description:
+/* This test implements sum reduction kernel, first with each threads own rank
+   as input and comparing the sum with expected sum output derieved from n(n-1)/2
+   formula. The second part, partitions this parent group into child subgroups
+   a.k.a tiles using using tiled_partition() collective operation. This can be called
+   with a static tile size, passed in templated non-type variable-tiled_partition<tileSz>,
+   or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these
+   cases.
+   This test tests functionality of cg group partitioning, (static and dynamic) and its respective
+   API's size(), thread_rank(), and sync().
+*/
+
+#include "test_common.h"
+#include <hip/hip_cooperative_groups.h>
+#include <stdio.h>
+#include <vector>
+
+using namespace cooperative_groups;
+
+#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
+
+/* Parallel reduce kernel.
+ *
+ * Step complexity: O(log n)
+ * Work complexity: O(n)
+ *
+ * Note: This kernel works only with power of 2 input arrays.
+ */
+__device__ int reduction_kernel(thread_group g, int* x, int val) {
+  int lane = g.thread_rank();
+  int sz = g.size();
+
+  for (int i = g.size() / 2; i > 0; i /= 2) {
+    // use lds to store the temporary result
+    x[lane] = val;
+    // Ensure all the stores are completed.
+    g.sync();
+
+    if (lane < i) {
+      val += x[lane + i];
+    }
+    // It must work on one tiled thread group at a time,
+    // and it must make sure all memory operations are
+    // completed before moving to the next stride.
+    // sync() here just does that.
+    g.sync();
+  }
+
+  // Choose the 0'th indexed thread that holds the reduction value to return
+  if (g.thread_rank() == 0) {
+    return val;
+  }
+  // Rest of the threads return no useful values
+  else {
+    return -1;
+  }
+}
+
+template <unsigned int tileSz>
+__global__ void kernel_cg_group_partition_static(int* result, bool isGlobalMem, int* globalMem) {
+  thread_block threadBlockCGTy = this_thread_block();
+  int threadBlockGroupSize = threadBlockCGTy.size();
+
+  int* workspace = NULL;
+
+  if (isGlobalMem) {
+    workspace = globalMem;
+  } else {
+    // Declare a shared memory
+    extern __shared__ int sharedMem[];
+    workspace = sharedMem;
+  }
+
+  int input, outputSum, expectedOutput;
+
+  // we pass its own thread rank as inputs
+  input = threadBlockCGTy.thread_rank();
+
+  expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2;
+
+  outputSum = reduction_kernel(threadBlockCGTy, workspace, input);
+
+  // Choose a leader thread to print the results
+  if (threadBlockCGTy.thread_rank() == 0) {
+    printf(" Sum of all ranks 0..%d in threadBlockCooperativeGroup is %d (expected %d)\n\n",
+           (int)threadBlockCGTy.size() - 1, outputSum, expectedOutput);
+    printf(" Creating %d groups, of tile size %d threads:\n\n",
+           (int)threadBlockCGTy.size() / tileSz, tileSz);
+  }
+
+  threadBlockCGTy.sync();
+
+  thread_block_tile<tileSz> tiledPartition = tiled_partition<tileSz>(threadBlockCGTy);
+
+  // This offset allows each group to have its own unique area in the workspace array
+  int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank();
+
+  outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input);
+
+  if (tiledPartition.thread_rank() == 0) {
+    printf(
+        "   Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread "
+        "rank: %d\n",
+        tiledPartition.size() - 1, outputSum, input);
+    result[input / (tileSz)] = outputSum;
+  }
+  return;
+}
+
+
+__global__ void kernel_cg_group_partition_dynamic(unsigned int tileSz, int* result,
+                                                  bool isGlobalMem, int* globalMem) {
+  thread_block threadBlockCGTy = this_thread_block();
+  int threadBlockGroupSize = threadBlockCGTy.size();
+
+  int* workspace = NULL;
+
+  if (isGlobalMem) {
+    workspace = globalMem;
+  } else {
+    // Declare a shared memory
+    extern __shared__ int sharedMem[];
+    workspace = sharedMem;
+  }
+
+  int input, outputSum, expectedOutput;
+
+  // input to reduction, for each thread, is its' rank in the group
+  input = threadBlockCGTy.thread_rank();
+
+  expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2;
+
+  outputSum = reduction_kernel(threadBlockCGTy, workspace, input);
+
+  if (threadBlockCGTy.thread_rank() == 0) {
+    printf(" Sum of all ranks 0..%d in threadBlockCooperativeGroup is %d\n\n",
+           (int)threadBlockCGTy.size() - 1, outputSum);
+    printf(" Creating %d groups, of tile size %d threads:\n\n",
+           (int)threadBlockCGTy.size() / tileSz, tileSz);
+  }
+
+  threadBlockCGTy.sync();
+
+  thread_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz);
+
+  // This offset allows each group to have its own unique area in the workspace array
+  int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank();
+
+  outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input);
+
+  if (tiledPartition.thread_rank() == 0) {
+    printf(
+        "   Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread "
+        "rank: %d\n",
+        tiledPartition.size() - 1, outputSum, input);
+
+    result[input / (tileSz)] = outputSum;
+  }
+  return;
+}
+
+// Search if the sum exists in the expected results array
+void verifyResults(int* hPtr, int* dPtr, int size) {
+  int i = 0, j = 0;
+  for (i = 0; i < size; i++) {
+    for (j = 0; j < size; j++) {
+      if (hPtr[i] == dPtr[j]) {
+        break;
+      }
+    }
+    if (j == size) {
+      failed(" Result verification failed!");
+    }
+  }
+}
+
+
+template <unsigned int tileSz> static void test_group_partition(bool useGlobalMem) {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = 64;
+
+  int numTiles = (blockSize * threadsPerBlock) / tileSz;
+
+  // Build an array of expected reduction sum output on the host
+  // based on the sum of their respective thread ranks for verification.
+  // eg: parent group has 64threads.
+  // child thread ranks: 0-15, 16-31, 32-47, 48-63
+  // expected sum:       120,   376,  632,  888
+  int* expectedSum = new int[numTiles];
+  int temp = 0, sum = 0;
+
+  for (int i = 1; i <= numTiles; i++) {
+    sum = temp;
+    temp = (((tileSz * i) - 1) * (tileSz * i)) / 2;
+    expectedSum[i-1] = temp - sum;
+  }
+
+  int* dResult = NULL;
+  hipMalloc((void**)&dResult, numTiles * sizeof(int));
+
+  int* globalMem = NULL;
+  if (useGlobalMem) {
+    hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int));
+  }
+
+  int* hResult = NULL;
+  hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
+  memset(hResult, 0, numTiles * sizeof(int));
+
+  if (useGlobalMem) {
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_cg_group_partition_static<tileSz>, blockSize, threadsPerBlock, 0, 0,
+                       dResult, useGlobalMem, globalMem);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  } else {
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_cg_group_partition_static<tileSz>, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, dResult, useGlobalMem, globalMem);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  }
+
+  hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost);
+
+  verifyResults(expectedSum, hResult, numTiles);
+
+  // Free all allocated memory on host and device
+  hipFree(dResult);
+  hipFree(hResult);
+  if (useGlobalMem) {
+    hipFree(globalMem);
+  }
+  delete[] expectedSum;
+
+  printf("\n...PASSED.\n\n");
+}
+
+static void test_group_partition(unsigned int tileSz, bool useGlobalMem) {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = 64;
+
+  int numTiles = (blockSize * threadsPerBlock) / tileSz;
+  // Build an array of expected reduction sum output on the host
+  // based on the sum of their respective thread ranks to use for verification
+  int* expectedSum = new int[numTiles];
+  int temp = 0, sum = 0;
+  for (int i = 1; i <= numTiles; i++) {
+    sum = temp;
+    temp = (((tileSz * i) - 1) * (tileSz * i)) / 2;
+    expectedSum[i-1] = temp - sum;
+  }
+
+  int* dResult = NULL;
+  hipMalloc(&dResult, sizeof(int) * numTiles);
+
+  int* globalMem = NULL;
+  if (useGlobalMem) {
+    hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int));
+  }
+
+  int* hResult = NULL;
+  hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
+  memset(hResult, 0, numTiles * sizeof(int));
+
+  // Launch Kernel
+  if (useGlobalMem) {
+    hipLaunchKernelGGL(kernel_cg_group_partition_dynamic, blockSize, threadsPerBlock, 0, 0, tileSz,
+                       dResult, useGlobalMem, globalMem);
+
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  } else {
+    hipLaunchKernelGGL(kernel_cg_group_partition_dynamic, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, tileSz, dResult, useGlobalMem, globalMem);
+
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  }
+
+  hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost);
+
+  verifyResults(expectedSum, hResult, numTiles);
+
+  // Free all allocated memory on host and device
+  hipFree(dResult);
+  hipFree(hResult);
+  if (useGlobalMem) {
+    hipFree(globalMem);
+  }
+  delete[] expectedSum;
+
+  printf("\n...PASSED.\n\n");
+}
+
+int main() {
+  // Use default device for validating the test
+  int deviceId;
+  ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
+  hipDeviceProp_t deviceProperties;
+  ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
+  int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
+
+  if (!deviceProperties.cooperativeLaunch) {
+    std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
+    if (hip_skip_tests_enabled()) {
+      return hip_skip_retcode();
+    } else {
+      passed();
+    }
+  }
+
+  bool useGlobalMem = true;
+  std::cout << "Testing static tiled_partition for different tile sizes" << std::endl;
+  std::cout << "\nUsing global memory for computation\n";
+  /* Test static tile_partition */
+  std::cout << "TEST 1:" << '\n' << std::endl;
+  test_group_partition<2>(useGlobalMem);
+  std::cout << "TEST 2:" << '\n' << std::endl;
+  test_group_partition<4>(useGlobalMem);
+  std::cout << "TEST 3:" << '\n' << std::endl;
+  test_group_partition<8>(useGlobalMem);
+  std::cout << "TEST 4:" << '\n' << std::endl;
+  test_group_partition<16>(useGlobalMem);
+  std::cout << "TEST 5:" << '\n' << std::endl;
+  test_group_partition<32>(useGlobalMem);
+
+  useGlobalMem = false;
+  std::cout << "Testing static tiled_partition for different tile sizes" << std::endl;
+  std::cout << "\nUsing shared memory for computation\n";
+  /* Test static tile_partition */
+  std::cout << "TEST 1:" << '\n' << std::endl;
+  test_group_partition<2>(useGlobalMem);
+  std::cout << "TEST 2:" << '\n' << std::endl;
+  test_group_partition<4>(useGlobalMem);
+  std::cout << "TEST 3:" << '\n' << std::endl;
+  test_group_partition<8>(useGlobalMem);
+  std::cout << "TEST 4:" << '\n' << std::endl;
+  test_group_partition<16>(useGlobalMem);
+  std::cout << "TEST 5:" << '\n' << std::endl;
+  test_group_partition<32>(useGlobalMem);
+
+
+  std::cout << "Now testing dynamic tiled_partition for different tile sizes" << '\n' << std::endl;
+
+  /* Test dynamic group partition*/
+  useGlobalMem = true;
+  int testNo = 1;
+  std::vector<unsigned int> tileSizes = {2, 4, 8, 16, 32};
+  std::cout << "\nUsing global memory for computation\n";
+  for (auto i : tileSizes) {
+    std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
+    test_group_partition(i, useGlobalMem);
+    testNo++;
+  }
+
+  useGlobalMem = false;
+  testNo = 1;
+  std::cout << "\nUsing shared memory for computation\n";
+  for (auto i : tileSizes) {
+    std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
+    test_group_partition(i, useGlobalMem);
+    testNo++;
+  }
+
+  passed();
+  return 0;
+}

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/thread_block_tile_shfl_ops.cpp

@@ -0,0 +1,231 @@
+/*
+Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+// Test Description:
+/* This test implements sum reduction kernel, first with each threads own rank
+   as input and comparing the sum with expected sum output derieved from n(n-1)/2
+   formula.
+   This sample tests functionality of intrinsics provided by thread_block_tile type,
+   shfl_down and shfl_xor.
+*/
+
+#include "test_common.h"
+#include <hip/hip_cooperative_groups.h>
+#include <stdio.h>
+#include <vector>
+
+using namespace cooperative_groups;
+
+#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
+
+template <unsigned int tileSz>
+__device__ int reduction_kernel_shfl_down(thread_block_tile<tileSz> const& g, volatile int val) {
+  int sz = g.size();
+
+  for (int i = sz / 2; i > 0; i >>= 1) {
+    val += g.shfl_down(val, i);
+  }
+
+  // Choose the 0'th indexed thread that holds the reduction value to return
+  if (g.thread_rank() == 0) {
+    return val;
+  }
+  // Rest of the threads return no useful values
+  else {
+    return -1;
+  }
+}
+
+template <unsigned int tileSz>
+__device__ int reduction_kernel_shfl_xor(thread_block_tile<tileSz> const& g, int val) {
+  int sz = g.size();
+
+  for (int i = sz / 2; i > 0; i >>= 1) {
+    val += g.shfl_xor(val, i);
+  }
+
+  // Choose the 0'th indexed thread that holds the reduction value to return
+  if (g.thread_rank() == 0) {
+    return val;
+  }
+  // Rest of the threads return no useful values
+  else {
+    return -1;
+  }
+}
+
+template <unsigned int tileSz>
+__global__ void kernel_cg_group_partition_static(int* result, bool runShflDown) {
+  thread_block threadBlockCGTy = this_thread_block();
+  int threadBlockGroupSize = threadBlockCGTy.size();
+  int input, outputSum, expectedSum;
+
+  // Choose a leader thread to print the results
+  if (threadBlockCGTy.thread_rank() == 0) {
+    printf(" Creating %d groups, of tile size %d threads:\n\n",
+           (int)threadBlockCGTy.size() / tileSz, tileSz);
+  }
+
+  threadBlockCGTy.sync();
+
+  thread_block_tile<tileSz> tiledPartition = tiled_partition<tileSz>(threadBlockCGTy);
+  int threadRank = tiledPartition.thread_rank();
+
+  input = tiledPartition.thread_rank();
+
+  // (n-1)(n)/2
+  expectedSum = ((tileSz - 1) * tileSz / 2);
+
+  if (runShflDown) {
+    outputSum = reduction_kernel_shfl_down(tiledPartition, input);
+
+    if (tiledPartition.thread_rank() == 0) {
+      printf(
+          "   Sum of all ranks 0..%d in this tiledPartition group using shfl_down is %d (expected "
+          "%d)\n",
+          tiledPartition.size() - 1, outputSum, expectedSum);
+      result[threadBlockCGTy.thread_rank() / (tileSz)] = outputSum;
+    }
+  } else {
+    outputSum = reduction_kernel_shfl_xor(tiledPartition, input);
+
+    if (tiledPartition.thread_rank() == 0) {
+      printf(
+          "   Sum of all ranks 0..%d in this tiledPartition group using shfl_xor is %d (expected "
+          "%d)\n",
+          tiledPartition.size() - 1, outputSum, expectedSum);
+      result[threadBlockCGTy.thread_rank() / (tileSz)] = outputSum;
+    }
+  }
+
+  return;
+}
+
+void verifyResults(int* ptr, int expectedResult, int numTiles) {
+  for (int i = 0; i < numTiles; i++) {
+    if (ptr[i] != expectedResult) {
+      failed(" Results do not match! ");
+    }
+  }
+}
+
+template <unsigned int tileSz> static void test_group_partition(bool runShflDown) {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = 64;
+
+  int numTiles = (blockSize * threadsPerBlock) / tileSz;
+  int expectedSum = ((tileSz - 1) * tileSz / 2);
+  int* expectedResult = new int[numTiles];
+
+  for (int i = 0; i < numTiles; i++) {
+    expectedResult[i] = expectedSum;
+  }
+
+  int* dResult = NULL;
+  int* hResult = NULL;
+
+  hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
+  memset(hResult, 0, numTiles * sizeof(int));
+
+  hipMalloc(&dResult, numTiles * sizeof(int));
+
+  if (runShflDown) {
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_cg_group_partition_static<tileSz>, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, dResult, runShflDown);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  } else {
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_cg_group_partition_static<tileSz>, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, dResult, runShflDown);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  }
+
+  hipMemcpy(hResult, dResult, sizeof(int) * numTiles, hipMemcpyDeviceToHost);
+
+  verifyResults(hResult, expectedSum, numTiles);
+
+  // Free all allocated memory on host and device
+  hipFree(dResult);
+  hipFree(hResult);
+  delete[] expectedResult;
+
+  printf("\n...PASSED.\n\n");
+}
+
+
+int main() {
+  // Use default device for validating the test
+  int deviceId;
+  ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
+  hipDeviceProp_t deviceProperties;
+  ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
+  int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
+
+  if (!deviceProperties.cooperativeLaunch) {
+    std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
+    if (hip_skip_tests_enabled()) {
+      return hip_skip_retcode();
+    } else {
+      passed();
+    }
+    return 0;
+  }
+
+  bool runShflDown = true;
+  std::cout << "Testing static tiled_partition for different tile sizes using shfl_down"
+            << std::endl;
+  /* Test static tile_partition */
+  std::cout << "TEST 1:" << '\n' << std::endl;
+  test_group_partition<2>(runShflDown);
+  std::cout << "TEST 2:" << '\n' << std::endl;
+  test_group_partition<4>(runShflDown);
+  std::cout << "TEST 3:" << '\n' << std::endl;
+  test_group_partition<8>(runShflDown);
+  std::cout << "TEST 4:" << '\n' << std::endl;
+  test_group_partition<16>(runShflDown);
+  std::cout << "TEST 5:" << '\n' << std::endl;
+  test_group_partition<32>(runShflDown);
+
+  runShflDown = false;
+  std::cout << "Testing static tiled_partition for different tile sizes using shfl_xor"
+            << std::endl;
+  /* Test static tile_partition */
+  std::cout << "TEST 1:" << '\n' << std::endl;
+  test_group_partition<2>(runShflDown);
+  std::cout << "TEST 2:" << '\n' << std::endl;
+  test_group_partition<4>(runShflDown);
+  std::cout << "TEST 3:" << '\n' << std::endl;
+  test_group_partition<8>(runShflDown);
+  std::cout << "TEST 4:" << '\n' << std::endl;
+  test_group_partition<16>(runShflDown);
+  std::cout << "TEST 5:" << '\n' << std::endl;
+  test_group_partition<32>(runShflDown);
+
+  passed();
+}

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/thread_block_tiled_shfl_up.cpp

@@ -0,0 +1,179 @@
+/*
+Copyright (c) 2020 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+// Test Description:
+/* This test implements prefix sum(scan) kernel, first with each threads own rank
+   as input and comparing the sum with expected serial summation output on CPU.
+
+   This sample tests functionality of intrinsics provided by thread_block_tile type,
+   shfl_up.
+*/
+#include "test_common.h"
+#include <hip/hip_cooperative_groups.h>
+#include <stdio.h>
+#include <vector>
+
+using namespace cooperative_groups;
+
+#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
+
+template <unsigned int tileSz>
+__device__ int prefix_sum_kernel(thread_block_tile<tileSz> const& g, volatile int val) {
+  int sz = g.size();
+#pragma unroll
+  for (int i = 1; i < sz; i <<= 1) {
+    int temp = g.shfl_up(val, i);
+
+    if (g.thread_rank() >= i) {
+      val += temp;
+    }
+  }
+  return val;
+}
+
+template <unsigned int tileSz> __global__ void kernel_cg_group_partition_static(int* dPtr) {
+  thread_block threadBlockCGTy = this_thread_block();
+  int threadBlockGroupSize = threadBlockCGTy.size();
+
+  int input, outputSum;
+
+  // we pass its own thread rank as inputs
+  input = threadBlockCGTy.thread_rank();
+
+  // Choose a leader thread to print the results
+  if (threadBlockCGTy.thread_rank() == 0) {
+    printf(" Creating %d groups, of tile size %d threads:\n\n",
+           (int)threadBlockCGTy.size() / tileSz, tileSz);
+  }
+
+  threadBlockCGTy.sync();
+
+  thread_block_tile<tileSz> tiledPartition = tiled_partition<tileSz>(threadBlockCGTy);
+
+  input = tiledPartition.thread_rank();
+
+  outputSum = prefix_sum_kernel(tiledPartition, input);
+
+  // Update the result array with the corresponsing prefix sum
+  dPtr[threadBlockCGTy.thread_rank()] = outputSum;
+  return;
+}
+
+void serialScan(int* ptr, int size) {
+  // Fill up the array
+  for (int i = 0; i < size; i++) {
+    ptr[i] = i;
+  }
+
+  int acc = 0;
+  for (int i = 0; i < size; i++) {
+    acc = acc + ptr[i];
+    ptr[i] = acc;
+  }
+}
+
+void printResults(int* ptr, int size) {
+  for (int i = 0; i < size; i++) {
+    std::cout << ptr[i] << " ";
+  }
+  std::cout << '\n';
+}
+
+void verifyResults(int* cpu, int* gpu, int size) {
+  for (unsigned int i = 0; i < size / sizeof(int); i++) {
+    if (cpu[i] != gpu[i]) {
+      failed(" Prefix sum results do not match.");
+    }
+  }
+}
+
+template <unsigned int tileSz> static void test_group_partition() {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = 64;
+
+  int* hPtr = NULL;
+  int* dPtr = NULL;
+  int* cpuPrefixSum = NULL;
+
+  int arrSize = blockSize * threadsPerBlock * sizeof(int);
+
+  hipHostMalloc(&hPtr, arrSize);
+  hipMalloc(&dPtr, arrSize);
+
+  // Launch Kernel
+  hipLaunchKernelGGL(kernel_cg_group_partition_static<tileSz>, blockSize, threadsPerBlock,
+                     threadsPerBlock * sizeof(int), 0, dPtr);
+  hipMemcpy(hPtr, dPtr, arrSize, hipMemcpyDeviceToHost);
+  err = hipDeviceSynchronize();
+  if (err != hipSuccess) {
+    fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+  }
+
+  cpuPrefixSum = new int[tileSz];
+  serialScan(cpuPrefixSum, tileSz);
+  std::cout << "\nPrefix sum results on CPU\n";
+  printResults(cpuPrefixSum, tileSz);
+
+  std::cout << "\nPrefix sum results on GPU\n";
+  printResults(hPtr, tileSz);
+  std::cout << "\n";
+  verifyResults(hPtr, cpuPrefixSum, tileSz);
+  std::cout << "Results verified!\n";
+
+  delete[] cpuPrefixSum;
+  hipFree(hPtr);
+  hipFree(dPtr);
+}
+
+int main() {
+  // Use default device for validating the test
+  int deviceId;
+  ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess);
+  hipDeviceProp_t deviceProperties;
+  ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess);
+  int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock;
+
+  if (!deviceProperties.cooperativeLaunch) {
+    std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
+    if (hip_skip_tests_enabled()) {
+      return hip_skip_retcode();
+    } else {
+      passed();
+    }
+    return 0;
+  }
+  std::cout << "Testing static tiled_partition for different tile sizes" << std::endl;
+  /* Test static tile_partition */
+  std::cout << "TEST 1:" << '\n' << std::endl;
+  test_group_partition<2>();
+  std::cout << "TEST 2:" << '\n' << std::endl;
+  test_group_partition<4>();
+  std::cout << "TEST 3:" << '\n' << std::endl;
+  test_group_partition<8>();
+  std::cout << "TEST 4:" << '\n' << std::endl;
+  test_group_partition<16>();
+  std::cout << "TEST 5:" << '\n' << std::endl;
+  test_group_partition<32>();
+  passed();
+}
+
+/* Kogge-Stone algorithm */