             imsize, C, group_size, groups, p1_data, p2_data, p3_data, x_data,
             dy_data, d_x_data);
       }
-
     } else {
       if (d_scale) {
         set_zero(dev_ctx, d_scale, static_cast<T>(0));

      }
      block_size_nchw = std::max(block_size_nchw, kps::details::kWarpSize);
      dim3 blocks(block_size_nchw);
      if (imsize < vec_size) {
        if (d_scale) {
          set_zero(dev_ctx, d_scale, static_cast<T>(0));
        }
        if (d_bias) {
          set_zero(dev_ctx, d_bias, static_cast<T>(0));
        }
        ScalarGetDsDbCUDAKernel<
            T><<<x_dims[0] * C, blocks, 0, dev_ctx.stream()>>>(
            imsize, x_data, dy_data, ds_data, db_data);

            x_data, mean_data, temp_var_data, size);
      }
    } else {
      set_zero(dev_ctx, mean, static_cast<T>(0));
      set_zero(dev_ctx, &temp_var, static_cast<T>(0));
      GroupNormForwardGetMeanAndVar<T><<<grid, threads, 0, dev_ctx.stream()>>>(
          x_data, x_dims[0], C, W, imsize, groups, group_size, mean_data,
          temp_var_data);
