         SplitKernel<<<grid_dims, block_dims, 0, context.stream()>>>(
             input.data<T>(), in_row, in_col, out0_col, outputs_data[0],
             outputs_data[1]);
+      } else if (o_num == 3) {
+        SplitKernel<<<grid_dims, block_dims, 0, context.stream()>>>(
+            input.data<T>(), in_row, in_col, out0_col, outputs_data[0],
+            outputs_data[1], outputs_data[2]);
+      } else if (o_num == 4) {
+        SplitKernel<<<grid_dims, block_dims, 0, context.stream()>>>(
+            input.data<T>(), in_row, in_col, out0_col, outputs_data[0],
+            outputs_data[1], outputs_data[2], outputs_data[3]);
       } else {
         SplitKernel<<<grid_dims, block_dims, 0, context.stream()>>>(
             input.data<T>(), in_row, in_col, out0_col, dev_out_gpu_data);

  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) {

        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>());
