__global__
void __launch_bounds__(512) fn_variance_kernel(
  const float* __restrict__ var,
  float* __restrict__ var_0,
  float* __restrict__ var_1,
  float* __restrict__ var_0_rf,
  welford_fp32* __restrict__ var_1_rf,
  int32_t* __restrict__ semaphore)
{
  float var_0_rf [ 1 ];
  welford_fp32 var_1_rf [ 1 ];
  __shared__ welford_fp32 shm32__welford_fp32_reduce [ 512 ];
  __shared__ float shm32__fp32_reduce [ 512 ];
  bool is_last_block_done [ 1 ];
  var_1_rf[0] = welford_fp32(0.00000000f, 0.00000000f, 0.00000000f);
  var_0_rf[0] = 0.00000000f;
  for (int32_t k = 0; k < 49; k += 1) {
    float var_local = var[(((k * 32 + blockIdx.y) * 16 + threadIdx.y) * 2 + blockIdx.x) * 32 + threadIdx.x];
    var_1_rf[0] = var_1_rf[0] + (welford_fp32)var_local;  // 计算variance(x)
    var_0_rf[0] = var_0_rf[0] + var_local;                // 计算sum(x)，实际上是为了计算mean(x)
  }
  var_0_rf[(blockIdx.y * 2 + blockIdx.x) * 32 + threadIdx.x] = cinn_discrete_reduce_sum_fp32(var_0_rf[0], shm32__fp32_reduce);
  var_1_rf[(blockIdx.y * 2 + blockIdx.x) * 32 + threadIdx.x] = cinn_discrete_reduce_sum_welford_fp32(var_1_rf[0], shm32__welford_fp32_reduce);
  is_last_block_done[0] = cinn_grid_reduce_update_semaphore(semaphore);
  if (is_last_block_done[0]) {
    var_0[blockIdx.x * 32 + threadIdx.x] = cinn_grid_reduce_sum_fp32(var_0_rf, 64, ((blockIdx.x * 32) + threadIdx.x));
    var_1[blockIdx.x * 32 + threadIdx.x] = (float)cinn_grid_reduce_sum_welford_fp32(var_1_rf, 64, ((blockIdx.x * 32) + threadIdx.x));
  }
}