         ConcatKernel<<<grid_dims, block_dims, 0, context.stream()>>>(
             inputs_data[0], inputs_data[1], in_col, out_row, out_col,
             output->data<T>());
+      } else if (in_num == 3) {
+        ConcatKernel<<<grid_dims, block_dims, 0, context.stream()>>>(
+            inputs_data[0], inputs_data[1], inputs_data[2], in_col, out_row,
+            out_col, output->data<T>());
+      } else if (in_num == 4) {
+        ConcatKernel<<<grid_dims, block_dims, 0, context.stream()>>>(
+            inputs_data[0], inputs_data[1], inputs_data[2], inputs_data[3],
+            in_col, out_row, out_col, output->data<T>());
       } else {
         ConcatKernel<<<grid_dims, block_dims, 0, context.stream()>>>(
             dev_ins_data, in_num, in_col, out_row, out_col, output->data<T>());

  SplitKernelDetail<T>(input_data, in_row, in_col, fixed_out_col, outputs_data);
}

static inline void GetBlockDims(const platform::CUDADeviceContext& context,
                                int num_rows, int num_cols, dim3* block_dims,
                                dim3* grid_dims) {


    memory::allocation::AllocationPtr tmp_dev_ins_data;
    const T** dev_ins_data = nullptr;
    if (!has_same_shape || in_num < 2 || in_num > 4) {
      tmp_dev_ins_data =
          platform::DeviceTemporaryAllocator::Instance().Get(context).Allocate(
              inputs_data.size() * sizeof(T*));
