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

* fix: add tail-case handling for elementwise_add_f16x8_pack_kernel to avoid out-of-bounds access

* Rename nx to idx

* Fix bug: nx -> idx

---------

Co-authored-by: Uper <[email protected]>

Author: lhycms

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]);
+    }
   }
 }