Merge pull request #384 from InfiniTensor/issue/6/fix

PanZezhong1725 · web-flow · commit 831021b8dfe4 · 2025-08-18T14:30:15.000+08:00
issue/6/fix type convertion
diff --git 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();
 

Original file line number	Diff line number	Diff line change
`@@ -22,7 +22,7 @@ __device__ void rmsnormBlock(`
`22`	`22`	`// Thread_0 computes RMS=1/sqrt(ss/dim+epsilon) and stores in shared memory`
`23`	`23`	`__shared__ Tcompute rms;`
`24`	`24`	`if (threadIdx.x == 0) {`
`25`		`- rms = Tdata(rsqrtf(ss / Tcompute(dim) + epsilon));`
	`25`	`+ rms = Tcompute(rsqrtf(ss / Tcompute(dim) + epsilon));`
`26`	`26`	`}`
`27`	`27`	`__syncthreads();`
`28`	`28`