CUDA: fix bug in rms_norm fusion

am17an · am17an · commit 5521b2ea561e · 2025-08-29T17:48:32.000+08:00
diff --git a/ggml/src/ggml-cuda/norm.cu b/ggml/src/ggml-cuda/norm.cu
@@ -187,7 +187,7 @@ static __global__ void rms_norm_f32(const float * x, float *       dst,
             dst[col] = scale * x[col] * mul[mul_col];
         } else if constexpr (do_add) {
             const int add_col = col % add_ncols;
-            dst[col] += add[add_col];
+            dst[col] = scale * x[col] + add[add_col];
         } else {
             dst[col] = scale * x[col];
         }

Original file line number	Diff line number	Diff line change
`@@ -187,7 +187,7 @@ static __global__ void rms_norm_f32(const float * x, float * dst,`
`187`	`187`	`dst[col] = scale * x[col] * mul[mul_col];`
`188`	`188`	`} else if constexpr (do_add) {`
`189`	`189`	`const int add_col = col % add_ncols;`
`190`		`- dst[col] += add[add_col];`
	`190`	`+ dst[col] = scale * x[col] + add[add_col];`
`191`	`191`	`} else {`
`192`	`192`	`dst[col] = scale * x[col];`
`193`	`193`	`}`