     auto base_n  = threadIdx.x + blockIdx.x * blockDim.x;
 
     if (base_n < B * N * W) {
-        T sum = 0;
-
-        for (size_t o = 0; o < H; ++o) {
-            sum += x[base_n * H + o];
+        const T * xx = x + base_n * H;
+
+        size_t o = 0;
+
+        T sum = 0;
+
+        for (; o + 7 < H; o += 8) {
+            sum += xx[o];
+            sum += xx[o + 1];
+            sum += xx[o + 2];
+            sum += xx[o + 3];
+            sum += xx[o + 4];
+            sum += xx[o + 5];
+            sum += xx[o + 6];
+            sum += xx[o + 7];
+        }
+
+        for (; o + 3 < H; o += 4) {
+            sum += xx[o];
+            sum += xx[o + 1];
+            sum += xx[o + 2];
+            sum += xx[o + 3];
+        }
+
+        for (; o + 1 < H; o += 2) {
+            sum += xx[o];
+            sum += xx[o + 1];
+        }
+
+        for (; o < H; ++o) {
+            sum += xx[o];
         }
 
         y[base_n] = sum;

    auto base_n  = threadIdx.x + blockIdx.x * blockDim.x;

    if (base_n < B * N) {
        T sum = 0;

        for (size_t o = 0; o < W; ++o) {
            sum += x[base_n * W + o];
        }

        y[base_n] = sum;


    // Phase 0

    int blockSize = 32;
    int gridSize = (b * n * w + blockSize - 1) / blockSize;

    bias_batch_sum4_kernel_zero<<<gridSize, blockSize>>>(b, n, w, h, x, tmp_zero);
