MapExpr supports concat op, WIP MapExpr2Ir

[jiahy0825]

Corner case:
前端的 fusion pass 会将其拆分为两个 Group

        x = builder.relu(x)
        y = builder.elementwise_add(x, y)
        out = builder.concat([x, y], axis=1)

__global__
void __launch_bounds__(1024) fn_fill_constant_0_max_1_3_kernel(const float* __restrict__ x, float* __restrict__ var_1)
{
  if (((int)threadIdx.x < 1024)) {
    var_1[(int)threadIdx.x] = max(x[(int)threadIdx.x], 0.00000000f);
  };
}

__global__
void __launch_bounds__(1024) fn_elementwise_add_2_concat_3_4_kernel(const float* __restrict__ var_1, const float* __restrict__ y, float* __restrict__ var_3)
{
  if (((int)blockIdx.x < 2)) {
    if (((int)threadIdx.x < 1024)) {
      var_3[((1024 * (int)blockIdx.x) + (int)threadIdx.x)] = ((((((1024 * (int)blockIdx.x) + (int)threadIdx.x) & 31) < 16)) ? var_1[((16 * ((int)threadIdx.x / 32)) + (((int)threadIdx.x & 31) + (512 * (int)blockIdx.x)))] : (var_1[(-16 + ((16 * ((int)threadIdx.x / 32)) + (((int)threadIdx.x & 31) + (512 * (int)blockIdx.x))))] + y[(-16 + ((16 * ((int)threadIdx.x / 32)) + (((int)threadIdx.x & 31) + (512 * (int)blockIdx.x))))]));
    };
  };
}