zero out last 13 bits

Signed-off-by: cliu-us <[email protected]>

Author: chichun-charlie-liu

fms_mo/custom_ext_kernels/triton_kernels.py

@@ -448,10 +448,10 @@ def round_and_trun(x, round_bit, trun_mask):
 @triton.jit
 def fp32_clamp_to_dl16(x):
     """clamp FP32 (1-8-23) TENSOR x to DL16 (1-6-9) range."""
-    # 1. rounding, add round bit to full uint representation
+    # 1. rounding: add round bit to full uint representation, zero out last 13 bits, back to float
     x = libdevice.float_as_uint(x)
     round_bit = 1 << (23 - 9 - 1)
-    x = libdevice.uint_as_float(x + round_bit)
+    x = libdevice.uint_as_float(((x + round_bit) >> 13) << 13)
 
     # 2. clamp to min/max:
     #   max = 2^32 * 1.(1111 1111 0)_base2 => 2^32*1.(1111 1111 1) will become inf