feat(runtime): add fp32 support to ReduceSum kernel

lib/Runtime/Kernels/hip/reduce_sum_kernel.hip

@@ -243,6 +243,50 @@ __global__ void reduce_sum_f16_kernel(
   }
 }
 
+// =============================================================================
+// FP32 reduction
+// =============================================================================
+
+// Parallel sum reduction for fp32. Some models upcast to fp32 before a sum for
+// numerical stability (e.g. a fp32 softmax-denominator / RMSNorm chain exported
+// as Cast(fp16->fp32) -> ... -> ReduceSum), so ReduceSum must accept fp32 as
+// well as fp16. Accumulator and storage are both float; structure is otherwise
+// identical to the fp16 kernel.
+__global__ void reduce_sum_f32_kernel(
+    const float* __restrict__ data,
+    float* __restrict__ output,
+    int64_t reduce_size,
+    int64_t inner,
+    int64_t num_output_elements) {
+  extern __shared__ float sdata_f32[];
+
+  int64_t out_idx = blockIdx.x;
+  if (out_idx >= num_output_elements) return;
+
+  int64_t oo = out_idx / inner;
+  int64_t ii = out_idx - oo * inner;
+  const float* slice = data + oo * reduce_size * inner + ii;
+  int tid = threadIdx.x;
+
+  float sum = 0.0f;
+  for (int64_t i = tid; i < reduce_size; i += blockDim.x) {
+    sum += slice[i * inner];
+  }
+  sdata_f32[tid] = sum;
+  __syncthreads();
+
+  for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
+    if (static_cast<unsigned int>(tid) < s) {
+      sdata_f32[tid] += sdata_f32[tid + s];
+    }
+    __syncthreads();
+  }
+
+  if (tid == 0) {
+    output[out_idx] = sdata_f32[0];
+  }
+}
+
 // =============================================================================
 // Host launcher
 // =============================================================================
@@ -362,6 +406,26 @@ extern "C" int hip_reduce_sum(
                        reduce_size, inner_size, num_output_elements);
       break;
     }
+    case HIP_DTYPE_FLOAT32: {
+      int block_size = pick_block_size();
+      size_t shared_mem = block_size * sizeof(float);
+
+      CUSTOM_KERNELS_DEBUG_LOG("[custom_kernels] hip_reduce_sum: dtype=FLOAT32, "
+              "input=%lld, output=%lld, reduce_size=%lld, "
+              "grid=%lld, block=%d, smem=%zu\n",
+              (long long)num_input_elements, (long long)num_output_elements,
+              (long long)reduce_size, (long long)num_output_elements,
+              block_size, shared_mem);
+
+      hipLaunchKernelGGL(reduce_sum_f32_kernel,
+                       dim3(static_cast<unsigned>(num_output_elements)),
+                       dim3(block_size),
+                       shared_mem, hip_stream,
+                       static_cast<const float*>(data),
+                       static_cast<float*>(output),
+                       reduce_size, inner_size, num_output_elements);
+      break;
+    }
     default:
       fprintf(stderr,
               "[custom_kernels] hip_reduce_sum: unsupported dtype=%d\n",

lib/Runtime/Kernels/include/hip_custom_kernels.h

@@ -984,10 +984,13 @@ HIP_KERNEL_API int hip_one_hot(
  *   num_output_elements - total output elements
  *   hip_dtype           - data type (hip_dtype_t value cast to int)
  *
- * Currently supported types: HIP_DTYPE_INT64, HIP_DTYPE_INT32, HIP_DTYPE_FLOAT16
+ * Currently supported types: HIP_DTYPE_INT64, HIP_DTYPE_INT32, HIP_DTYPE_FLOAT16,
+ * HIP_DTYPE_FLOAT32
  *   - INT32 accumulates in int64 internally to avoid overflow on large slices.
  *   - FLOAT16 accumulates in float internally to preserve precision; the
  *     final result is narrowed back to half.
+ *   - FLOAT32 accumulates and stores in float; required by models that upcast
+ *     to fp32 before the sum for numerical stability.
  * Returns: 0 on success, non-zero on failure
  */
 /* `inner_size` = product of input dims AFTER the reduced axis (1 for a

lib/Runtime/real/reduce_sum.cpp

@@ -24,11 +24,16 @@
 
 // Map HIPDNN_EP_DATATYPE_* -> hip_dtype_t for hip_reduce_sum.
 // The two enum systems use different orderings; only types implemented in
-// reduce_sum_kernel.hip are listed here.
+// reduce_sum_kernel.hip are listed here. fp32 is required by models that upcast
+// to fp32 before the sum for numerical stability (e.g. a fp32 softmax-
+// denominator / RMSNorm chain exported as Cast(fp16->fp32) -> ... ->
+// ReduceSum).
 static int hipdnn_to_hip_dtype(int64_t hipdnn_type) {
   switch (hipdnn_type) {
   case HIPDNN_EP_DATATYPE_HALF:
     return HIP_DTYPE_FLOAT16;
+  case HIPDNN_EP_DATATYPE_FLOAT:
+    return HIP_DTYPE_FLOAT32;
   case HIPDNN_EP_DATATYPE_INT32:
     return HIP_DTYPE_INT32;
   case HIPDNN_EP_DATATYPE_INT64:
@@ -85,7 +90,7 @@ int wrap_reduce_sum(RuntimeState *state, void *data, void *axes, void *output,
   if (hip_dtype < 0) {
     fprintf(stderr,
             "[REAL] wrap_reduce_sum: unsupported data_type=%s(%lld) "
-            "(supported: f16, i32, i64)\n",
+            "(supported: f16, f32, i32, i64)\n",
             hipdnn_ep_datatype_name(data_type), (long long)data_type);
     return -1;
   }