       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;
-        }
+        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;
       }
     }
     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;
-      }
+      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;
     }
     // Intra-warp reduction (all-reduce)
     for(int mask = blockDim.x / 2; mask > 0; mask /= 2) {

                            const OpContext& ctx, const std::vector<TBlob>& inputs,
                            const std::vector<OpReqType>& req,
                            const std::vector<TBlob>& outputs) {

}

template<>

                               const std::vector<OpReqType>& req,
                               const std::vector<TBlob>& outputs) {
  const LayerNormParam& param = nnvm::get<LayerNormParam>(attrs.parsed);
  if (req[0] == kNullOp) return;
  CHECK_NE(req[0], kAddTo);
  int axis = param.axis;
