fixes hip compatability with new modules

Author: Skyler Ruiter

include/backend/atomics.h

@@ -12,27 +12,52 @@
  * known to be in the same block, so scoping the atomic down from "device"
  * saves the extra memory-fence cost of a full device-wide atomic.
  *
- * ROCm's HIP declares `atomicAdd_block`/`atomicOr_block` with identical names
- * and signatures to CUDA's (`hip/amd_detail/amd_hip_atomic.h`) — unlike the
- * warp shuffle/ballot family (see warp.h), there is no warp/wavefront-width
- * dependency in a block-scope atomic, so naively this should "just hipify".
- * That said, this codebase has been burned before by CUDA intrinsics that
- * looked source-portable but silently diverged under HIP (see warp.h's
- * `__shfl_down_sync` width story), and — unlike every other intrinsic family
- * touched during the HIP port — these have **not yet been exercised or
- * verified on real AMD hardware** (no existing call site in this codebase
- * uses any block-scoped atomic). This header exists so there is exactly one
- * place to patch if that verification turns up a divergence, rather than
- * scattering raw `atomicAdd_block`/`atomicOr_block` calls through the RARE/RAZE
- * kernels. Treat the HIP branch below as "expected to work, not yet confirmed
- * on hardware" until it has been built and run on the project's MI100 target.
+ * The `_block` suffix family is CUDA-only spelling. Contrary to this header's
+ * original expectation ("ROCm declares them with identical names"), ROCm 6.4.1
+ * declares neither `atomicAdd_block` nor `atomicOr_block` anywhere under
+ * `hip/` — verified by grep against the installed toolchain, and by the
+ * `use of undeclared identifier 'atomicAdd_block'` error the first HIP build
+ * of the RARE/RAZE stages produced. This is exactly the divergence the header
+ * was created to absorb, so the patch lands here and the call sites are
+ * untouched.
+ *
+ * HIP's equivalent is the scoped-atomic builtin family
+ * (`__hip_atomic_fetch_add`/`_or` + `__HIP_MEMORY_SCOPE_WORKGROUP`), where
+ * "workgroup" is AMD's name for a CUDA block. Memory ordering is
+ * `__ATOMIC_RELAXED` to match CUDA's `atomic*_block`, which carry no ordering
+ * guarantees beyond the atomicity of the read-modify-write itself — the RARE/
+ * RAZE call sites (per-block histogram accumulation; OR-ing a straddling
+ * partial word into shared memory) separately `__syncthreads()` before reading
+ * the accumulated result, so they rely on that barrier for ordering, not on
+ * the atomic.
+ *
+ * Unlike the warp shuffle/ballot family (see warp.h), a block-scoped atomic has
+ * no wavefront-width dependency, so there is no 32-vs-64-lane subtlety here.
  */
 
 #include "backend/api.h"
 
 namespace fz {
 namespace backend {
 
+#if defined(FZGMOD_BACKEND_HIP)
+
+/** Block-scoped atomic add. `T` must be an arithmetic type the HIP scoped-atomic builtin accepts (int, unsigned, unsigned long long, float, double). */
+template <typename T>
+__device__ inline T atomicAddBlock(T* addr, T val) {
+    return __hip_atomic_fetch_add(addr, val, __ATOMIC_RELAXED,
+                                  __HIP_MEMORY_SCOPE_WORKGROUP);
+}
+
+/** Block-scoped atomic bitwise-OR. `T` must be an integral type (int, unsigned, unsigned long long). */
+template <typename T>
+__device__ inline T atomicOrBlock(T* addr, T val) {
+    return __hip_atomic_fetch_or(addr, val, __ATOMIC_RELAXED,
+                                 __HIP_MEMORY_SCOPE_WORKGROUP);
+}
+
+#else // CUDA
+
 /** Block-scoped atomic add. `T` must be one of the types CUDA's `atomicAdd_block` overloads on (int, unsigned, unsigned long long, float, double). */
 template <typename T>
 __device__ inline T atomicAddBlock(T* addr, T val) {
@@ -45,5 +70,7 @@ __device__ inline T atomicOrBlock(T* addr, T val) {
     return atomicOr_block(addr, val);
 }
 
+#endif
+
 } // namespace backend
 } // namespace fz

modules/coders/adaptive_bitpack/adaptive_bitpack_kernels.cu

@@ -49,7 +49,12 @@ __device__ __forceinline__ uint32_t warpBitTranspose32(uint32_t v, unsigned lane
             case 3:  m = 0xFF00FF00u; break;
             default: m = 0xFFFF0000u; break;
         }
-        const uint32_t t = __shfl_xor_sync(0xFFFFFFFFu, v, s);
+        // Width pinned to 32: this butterfly is a 32-lane algorithm, and that is
+        // what CUDA's implicit `width = warpSize` already meant. Under HIP the same
+        // default is 64, which would pull lanes 32-63 into each exchange and corrupt
+        // the transpose — see backend/warp.h failure mode 2. The wrapper also
+        // supplies a 64-bit-wide mask, which HIP's __shfl_xor_sync static_asserts on.
+        const uint32_t t = fz::backend::shflXor(v, static_cast<int>(s), 32);
         v = (lane & s) ? ((v & m) | ((t & m) >> s))
                        : ((v & ~m) | ((t & ~m) << s));
     }

modules/coders/clog/clog_stage.h

@@ -29,7 +29,7 @@
 
 #include "stage/stage.h"
 #include "fzm_format.h"
-#include <cuda_runtime.h>
+#include "backend/types.h"
 #include <cstdint>
 #include <cstring>
 #include <stdexcept>

modules/coders/hclog/hclog_stage.h

@@ -32,7 +32,7 @@
 
 #include "stage/stage.h"
 #include "fzm_format.h"
-#include <cuda_runtime.h>
+#include "backend/types.h"
 #include <cstdint>
 #include <cstring>
 #include <stdexcept>

modules/coders/rare/rare_stage.h

@@ -33,7 +33,7 @@
 
 #include "stage/stage.h"
 #include "fzm_format.h"
-#include <cuda_runtime.h>
+#include "backend/types.h"
 #include <cstdint>
 #include <cstring>
 #include <stdexcept>

modules/coders/raze/raze_stage.h

@@ -32,7 +32,7 @@
 
 #include "stage/stage.h"
 #include "fzm_format.h"
-#include <cuda_runtime.h>
+#include "backend/types.h"
 #include <cstdint>
 #include <cstring>
 #include <stdexcept>