fix: remove redundant load (#495)

Bruce-Lee-LY · web-flow · commit 2de16b0f4afb · 2024-09-10T23:00:38.000-07:00
diff --git a/include/flashinfer/norm.cuh b/include/flashinfer/norm.cuh
@@ -178,7 +178,6 @@ __global__ void FusedAddRMSNormKernel(T* __restrict__ input, T* __restrict__ res
     weight_vec.fill(0.f);
     residual_vec.fill(0.f);
     if ((i * num_threads + thread_id) * VEC_SIZE < d) {
-      input_vec.load(input + bx * d + i * num_threads * VEC_SIZE + thread_id * VEC_SIZE);
       weight_vec.load(weight + i * num_threads * VEC_SIZE + thread_id * VEC_SIZE);
       residual_vec.load(residual + bx * d + i * num_threads * VEC_SIZE + thread_id * VEC_SIZE);
     }

Original file line number	Diff line number	Diff line change
`@@ -178,7 +178,6 @@ __global__ void FusedAddRMSNormKernel(T* __restrict__ input, T* __restrict__ res`
`178`	`178`	`weight_vec.fill(0.f);`
`179`	`179`	`residual_vec.fill(0.f);`
`180`	`180`	`if ((i * num_threads + thread_id) * VEC_SIZE < d) {`
`181`		`- input_vec.load(input + bx * d + i * num_threads * VEC_SIZE + thread_id * VEC_SIZE);`
`182`	`181`	`weight_vec.load(weight + i * num_threads * VEC_SIZE + thread_id * VEC_SIZE);`
`183`	`182`	`residual_vec.load(residual + bx * d + i * num_threads * VEC_SIZE + thread_id * VEC_SIZE);`
`184`	`183`	`}`