                                   dim3* gb_block_dim, dim3* gb_grid_dim,
                                   int* row_repeat, int* npart) {
   *npart = 16;
-  *part_grad_block_dim = dim3(32, 4);
+  *part_grad_block_dim = dim3(32, 16);
   *part_grad_grid_dim = dim3((nchannel + 32 - 1) / 32, *npart);
   *row_repeat = 4;
   *gb_block_dim = dim3(32, *npart);

          DType local_mean = mean_data[r];
          DType local_std = std_data[r];
          int read_idx = r * nchannel + c;
          int write_idx = r_offset * row_stride + threadIdx.x;
          buf_gamma_grad[write_idx] += (in_data[read_idx] - local_mean) / local_std
                                                                        * out_grad[read_idx];
          buf_beta_grad[write_idx] += out_grad[read_idx];
        }
      }
    }
  }
  __syncthreads();
  DType local_gamma_grad = 0;
  DType local_beta_grad = 0;
  for(int i = 0; i < row_repeat; ++i) {
    int idx = (i * blockDim.y + threadIdx.y) * row_stride + threadIdx.x;
    local_gamma_grad += buf_gamma_grad[idx];
    local_beta_grad += buf_beta_grad[idx];
  }
  buf_gamma_grad[threadIdx.y * row_stride + threadIdx.x] = local_gamma_grad;
  buf_beta_grad[threadIdx.y * row_stride + threadIdx.x] = local_beta_grad;
  __syncthreads();
  for(int offset = blockDim.y/2;  offset > 0;  offset >>= 1) {
    if(threadIdx.y < offset) {
      int idx1 = threadIdx.y * row_stride + threadIdx.x;
      int idx2 = (threadIdx.y + offset) * row_stride + threadIdx.x;
      buf_gamma_grad[idx1] += buf_gamma_grad[idx2];
      buf_beta_grad[idx1] += buf_beta_grad[idx2];
    }

        ctx.requested[0].get_space_typed<gpu, 1, DType>(Shape1(2 * npart * nchannel), s);
      DType* part_gamma_grad_ptr = workspace.dptr_;
      DType* part_beta_grad_ptr = workspace.dptr_ + npart * nchannel;
      const int nshared_K1 = 2 * part_grad_block_dim.x * part_grad_block_dim.y * sizeof(DType);
      const int nshared_K2 = 2 * gb_block_dim.x * gb_block_dim.y * sizeof(DType);
      DType* gamma_grad_ptr = (gamma_grad_req != kNullOp) ? gamma_grad.dptr<DType>() : nullptr;
      DType* beta_grad_ptr = (beta_grad_req != kNullOp) ? beta_grad.dptr<DType>() : nullptr;
