 
     // Phase 2
 
-    blockSize = 64;
-    gridSize = (n + blockSize - 1) / blockSize;
-
-    bias_batch_sum4_kernel_second<Mean><<<gridSize, blockSize>>>(b, n, w, h, tmp_first, y);
+    blockSize = 128;
+    gridSize = (8 * n + blockSize - 1) / blockSize;
+
+    bias_batch_sum4_kernel_second<8, Mean><<<gridSize, blockSize>>>(b, n, w, h, 8 * n, tmp_first, y);
 
     cuda_check(cudaFree(tmp_zero));
     cuda_check(cudaFree(tmp_first));

    }
}

template <bool Mean, typename T>
__global__ void bias_batch_sum4_kernel_second(size_t B, size_t N, size_t W, size_t H, const T* x, T* y) {
    auto n  = threadIdx.x + blockIdx.x * blockDim.x;

    if (n < N) {
        T sum = 0;

        for (size_t b = 0; b < B; ++b) {
            sum += x[b * N + n];
        }

        if (Mean) {
            y[n] = sum / (B * W * H);
        } else {
            y[n] = sum;
        }
    }
}


    cudaMemset(tmp_zero, 0, b * n * w * sizeof(T));
    cudaMemset(tmp_first, 0, b * n * sizeof(T));
    cudaMemset(y, 0, n * sizeof(T));

    // Phase 0 (Bottleneck)

