From 11d3c4164ec62155d739391a571a4b66bb333a0b Mon Sep 17 00:00:00 2001
From: Daniel Maddern <daniel.maddern@illuvium.io>
Date: Sat, 15 Aug 2026 02:45:16 +0700
Subject: [PATCH] Match NVFP4 midpoint rounding

---
 src/h3_blackwell_runtime/csrc/nvfp4_scale.cu | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu b/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu
index de86db8..6fb519c 100644
--- a/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu
+++ b/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu
@@ -58,19 +58,19 @@ __inline__ __device__ uint8_t encode_fp4_e2m1(float value) {
   const bool negative = signbit(value);
   float abs_value = fabsf(value);
   uint8_t code = 0;
-  if (abs_value >= 5.0f) {
+  if (abs_value > 5.0f) {
     code = 7;
-  } else if (abs_value >= 3.5f) {
+  } else if (abs_value > 3.5f) {
     code = 6;
-  } else if (abs_value >= 2.5f) {
+  } else if (abs_value > 2.5f) {
     code = 5;
-  } else if (abs_value >= 1.75f) {
+  } else if (abs_value > 1.75f) {
     code = 4;
-  } else if (abs_value >= 1.25f) {
+  } else if (abs_value > 1.25f) {
     code = 3;
-  } else if (abs_value >= 0.75f) {
+  } else if (abs_value > 0.75f) {
     code = 2;
-  } else if (abs_value >= 0.25f) {
+  } else if (abs_value > 0.25f) {
     code = 1;
   }
   return negative ? static_cast<uint8_t>(code | 0x8u) : code;