     for(int l = tid; l < nchannel; l += nthread) {
       DType ele_out_grad = out_grad[bid * nchannel + l];
       DType ele_x = in_data[bid * nchannel + l];
-      if(!no_gamma) {
-        DType ele_gamma = gamma[l];
+      if(data_addto) {
         data_grad[bid * nchannel + l] +=
-          ele_out_grad * ele_gamma * invstd_eps - sum_val0 - (ele_x - mean) * invstd_eps * sum_val1;
+          ele_out_grad * gamma[l] * invstd_eps - sum_val0
+                                               - (ele_x - mean) * invstd_eps * sum_val1;
       } else {
-        data_grad[bid * nchannel + l] +=
-          ele_out_grad * invstd_eps - sum_val0 - (ele_x - mean) * invstd_eps * sum_val1;
+        data_grad[bid * nchannel + l] =
+          ele_out_grad * gamma[l] * invstd_eps - sum_val0
+                                               - (ele_x - mean) * invstd_eps * sum_val1;
       }
     }
   }

      for(int i = 0; i < LOAD_UNROLL; ++i) {
        DType ele_og = out_grad[bid * nchannel + l + i];
        DType ele_x = in_data[bid * nchannel + l + i];
        if(!no_gamma) {
          DType ele_gamma = gamma[l + i];
          sum_val0 += ele_og * ele_gamma * invstd_eps;
          sum_val1 += ele_og * ele_gamma * (ele_x - mean) * invstd_eps * invstd_eps;
        } else {
          sum_val0 += ele_og * invstd_eps;
          sum_val1 += ele_og * (ele_x - mean) * invstd_eps * invstd_eps;
        }
      }
    }
    for(; l < nchannel; ++l) {
      DType ele_og = out_grad[bid * nchannel + l];
      DType ele_x = in_data[bid * nchannel + l];
      if(!no_gamma) {
        DType ele_gamma = gamma[l];
        sum_val0 += ele_og * ele_gamma * invstd_eps;
        sum_val1 += ele_og * ele_gamma * (ele_x - mean) * invstd_eps * invstd_eps;
      } else {
        sum_val0 += ele_og * invstd_eps;
        sum_val1 += ele_og * (ele_x - mean) * invstd_eps * invstd_eps;
      }
    }
    // Intra-warp reduction (all-reduce)
    for(int mask = blockDim.x / 2; mask > 0; mask /= 2) {

  CHECK_EQ(mean_data.CheckContiguous(), true);
  CHECK_EQ(std_data.CheckContiguous(), true);

  // Lauch the kernel. The dynamic shared memory size is
  // sizeof(DType) * blockDim.y + sizeof(DType) * blockDim.y / 2
  int nbatch = data_shape[0];
  int nchannel = data_shape[1];
  float eps = param.eps;
  int ngrid_x = (nbatch > kMaxGridDim) ? (nbatch + kBaseGridNum - 1) / kBaseGridNum : nbatch;
  int ngrid_y = (nbatch > kMaxGridDim) ? kBaseGridNum : 1;
  int nthread_y = 0;
  const dim3 dimGrid(ngrid_x, ngrid_y);
  if(nchannel <= 32) {
    nthread_y = 1;
  } else {
    nthread_y = 2;
  }
  cudaStream_t stream = Stream<gpu>::GetStream(ctx.get_stream<gpu>());
  const dim3 dimBlock(32, nthread_y);
  MSHADOW_REAL_TYPE_SWITCH(in_data.type_flag_, DType, {
    int nshared = nthread_y > 1 ? nthread_y * sizeof(DType) + (nthread_y / 2) * sizeof(DType) : 0;
    CheckLaunchParam(dimGrid, dimBlock);
