 
     // Phase 0
 
-    int blockSize = 64;
+    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);

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

        y[base_n] = sum;

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

    if (base_n < B * N) {
        const T * xx = x + base_n * W;

        T sum = 0;

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

        y[base_n] = sum;
