fix: coalesced_threads size-1 fallback so warp-aggregated patterns work on SPIR-V

Author: pvelesko

include/hip/spirv_hip_cooperative_groups.h

@@ -434,6 +434,12 @@ class coalesced_group : public thread_group {
   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.");
 
+    // Singleton coalesced group (chipStar SPIR-V fallback when the active-
+    // lane mask is unavailable): the only valid source rank is 0 and the
+    // result is the caller's own value.
+    if (size() == 1)
+      return var;
+
     srcRank = srcRank % static_cast<int>(size());
 
     int lane = (size() == __AMDGCN_WAVEFRONT_SIZE) ? srcRank
@@ -508,7 +514,12 @@ class coalesced_group : public thread_group {
  */
 
 __CG_QUALIFIER__ coalesced_group coalesced_threads() {
-  return cooperative_groups::coalesced_group(__builtin_amdgcn_read_exec());
+  // chipStar's SPIR-V backend has no way to query the active-lane mask
+  // (the AMDGCN exec register), so we conservatively model each thread
+  // as its own coalesced group of size 1. This keeps semantics correct
+  // for common patterns such as warp-aggregated atomics (atomicAggInc)
+  // at the cost of the warp-level optimisation.
+  return cooperative_groups::coalesced_group(static_cast<lane_mask>(1));
 }
 
 /**