From 1529608b7f2cfba8b71853bed3da0e465f8745dd Mon Sep 17 00:00:00 2001
From: PanZezhong <panzezhong@qiyuanlab.com>
Date: Mon, 18 Aug 2025 14:23:24 +0800
Subject: [PATCH] issue/6/fix type convertion

---
 src/infiniop/ops/rms_norm/cuda/kernel.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/infiniop/ops/rms_norm/cuda/kernel.cuh b/src/infiniop/ops/rms_norm/cuda/kernel.cuh
index 53c4a5587..86bd519e4 100644
--- a/src/infiniop/ops/rms_norm/cuda/kernel.cuh
+++ b/src/infiniop/ops/rms_norm/cuda/kernel.cuh
@@ -22,7 +22,7 @@ __device__ void rmsnormBlock(
     // Thread_0 computes RMS=1/sqrt(ss/dim+epsilon) and stores in shared memory
     __shared__ Tcompute rms;
     if (threadIdx.x == 0) {
-        rms = Tdata(rsqrtf(ss / Tcompute(dim) + epsilon));
+        rms = Tcompute(rsqrtf(ss / Tcompute(dim) + epsilon));
     }
     __syncthreads();