fix: add tail-case handling for elementwise_add_f16x8_pack_kernel to … (#380)

lhycms · web-flow · commit 6d8844862db2 · 2025-08-25T11:25:40.000+08:00
* fix: add tail-case handling for elementwise_add_f16x8_pack_kernel to avoid out-of-bounds access

* Rename nx to idx

* Fix bug: nx -&gt; idx

---------

Co-authored-by: Uper &lt;41718895+Hyliu-BUAA@users.noreply.github.com&gt;
diff --git a/kernels/elementwise/elementwise.cu b/kernels/elementwise/elementwise.cu
@@ -117,6 +117,10 @@ __global__ void elementwise_add_f16x8_pack_kernel(half *a, half *b, half *c,
   // reinterpret as float4 and store 128 bits in 1 memory issue.
   if ((idx + 7) < N) {
     LDST128BITS(c[idx]) = LDST128BITS(pack_c[0]);
+  } else {
+    for (int i=0; idx+i<N; i++) {
+      c[idx+i] = __hadd(a[idx+i], b[idx+i]);
+    }
   }
 }
 

Original file line number	Diff line number	Diff line change
`@@ -117,6 +117,10 @@ __global__ void elementwise_add_f16x8_pack_kernel(half a, half b, half *c,`
`117`	`117`	`// reinterpret as float4 and store 128 bits in 1 memory issue.`
`118`	`118`	`if ((idx + 7) < N) {`
`119`	`119`	`LDST128BITS(c[idx]) = LDST128BITS(pack_c[0]);`
	`120`	`+ } else {`
	`121`	`+ for (int i=0; idx+i<N; i++) {`
	`122`	`+ c[idx+i] = __hadd(a[idx+i], b[idx+i]);`
	`123`	`+ }`
`120`	`124`	`}`
`121`	`125`	`}`
`122`	`126`