[dpdk-dev] [RFC PATCH v2 0/7] heterogeneous computing library

[Thomas Monjalon <thomas@monjalon.net>]

From: Elena Agostini <eagostini@nvidia.com>

In heterogeneous computing system, processing is not only in the CPU.
Some tasks can be delegated to devices working in parallel.

The goal of this new library is to enhance the collaboration between
DPDK, that's primarily a CPU framework, and other type of devices like GPUs.

When mixing network activity with task processing on a non-CPU device,
there may be the need to put in communication the CPU with the device
in order to manage the memory, synchronize operations, exchange info, etc..

This library provides a number of new features:
- Interoperability with device specific library with generic handlers
- Possibility to allocate and free memory on the device
- Possibility to allocate and free memory on the CPU but visible from the device
- Communication functions to enhance the dialog between the CPU and the device

The infrastructure is prepared to welcome drivers in drivers/hc/
as the upcoming NVIDIA one, implementing the hcdev API.

Some parts are not complete:
  - locks
  - memory allocation table
  - memory freeing
  - guide documentation
  - integration in devtools/check-doc-vs-code.sh
  - unit tests
  - integration in testpmd to enable Rx/Tx to/from GPU memory.

Below is a pseudo-code to give an example about how to use functions
in this library in case of a CUDA application.


Elena Agostini (4):
  hcdev: introduce heterogeneous computing device library
  hcdev: add memory API
  hcdev: add communication flag
  hcdev: add communication list

Thomas Monjalon (3):
  hcdev: add event notification
  hcdev: add child device representing a device context
  hcdev: support multi-process

 .gitignore                             |   1 +
 MAINTAINERS                            |   6 +
 doc/api/doxy-api-index.md              |   1 +
 doc/api/doxy-api.conf.in               |   1 +
 doc/guides/conf.py                     |   8 +
 doc/guides/hcdevs/features/default.ini |  13 +
 doc/guides/hcdevs/index.rst            |  11 +
 doc/guides/hcdevs/overview.rst         |  11 +
 doc/guides/index.rst                   |   1 +
 doc/guides/prog_guide/hcdev.rst        |   5 +
 doc/guides/prog_guide/index.rst        |   1 +
 doc/guides/rel_notes/release_21_08.rst |   5 +
 drivers/hc/meson.build                 |   4 +
 drivers/meson.build                    |   1 +
 lib/hcdev/hcdev.c                      | 789 +++++++++++++++++++++++++
 lib/hcdev/hcdev_driver.h               |  96 +++
 lib/hcdev/meson.build                  |  12 +
 lib/hcdev/rte_hcdev.h                  | 592 +++++++++++++++++++
 lib/hcdev/version.map                  |  35 ++
 lib/meson.build                        |   1 +
 20 files changed, 1594 insertions(+)
 create mode 100644 doc/guides/hcdevs/features/default.ini
 create mode 100644 doc/guides/hcdevs/index.rst
 create mode 100644 doc/guides/hcdevs/overview.rst
 create mode 100644 doc/guides/prog_guide/hcdev.rst
 create mode 100644 drivers/hc/meson.build
 create mode 100644 lib/hcdev/hcdev.c
 create mode 100644 lib/hcdev/hcdev_driver.h
 create mode 100644 lib/hcdev/meson.build
 create mode 100644 lib/hcdev/rte_hcdev.h
 create mode 100644 lib/hcdev/version.map



////////////////////////////////////////////////////////////////////////
///// HCDEV library + CUDA functions
////////////////////////////////////////////////////////////////////////
#define GPU_PAGE_SHIFT 16
#define GPU_PAGE_SIZE (1UL << GPU_PAGE_SHIFT)

int main() {
    struct rte_hcdev_flag quit_flag;
    struct rte_hcdev_comm_list *comm_list;
    int nb_rx = 0;
    int comm_list_entry = 0;
    struct rte_mbuf * rx_mbufs[max_rx_mbufs];
    cudaStream_t cstream;
    struct rte_mempool *mpool_payload, *mpool_header;
    struct rte_pktmbuf_extmem ext_mem;
    int16_t dev_id;

    /* Initialize CUDA objects (cstream, context, etc..). */
    /* Use hcdev library to register a new CUDA context if any */
    /* Let's assume the application wants to use the default context of the GPU device 0 */
    dev_id = 0;

    /* Create an external memory mempool using memory allocated on the GPU. */
    ext_mem.elt_size = mbufs_headroom_size;
                ext_mem.buf_len = RTE_ALIGN_CEIL(mbufs_num * ext_mem.elt_size, GPU_PAGE_SIZE);
    ext_mem.buf_iova = RTE_BAD_IOVA;
    ext_mem.buf_ptr = rte_hcdev_malloc(dev_id, ext_mem.buf_len, 0);
    rte_extmem_register(ext_mem.buf_ptr, ext_mem.buf_len, NULL, ext_mem.buf_iova, GPU_PAGE_SIZE);
    rte_dev_dma_map(rte_eth_devices[l2fwd_port_id].device, ext_mem.buf_ptr, ext_mem.buf_iova, ext_mem.buf_len);
    mpool_payload = rte_pktmbuf_pool_create_extbuf("gpu_mempool", mbufs_num,
                                                    0, 0, ext_mem.elt_size,
                                                    rte_socket_id(), &ext_mem, 1);

    /*
     * Create CPU - device communication flag. With this flag, the CPU can tell to the CUDA kernel
     * to exit from the main loop.
     */
    rte_hcdev_comm_create_flag(dev_id, &quit_flag, RTE_HCDEV_COMM_FLAG_CPU);
    rte_hcdev_comm_set_flag(&quit_flag, 0);

    /*
     * Create CPU - device communication list. Each entry of this list will be populated by the CPU
     * with a new set of received mbufs that the CUDA kernel has to process.
     */
    comm_list = rte_hcdev_comm_create_list(dev_id, num_entries);

    /* A very simple CUDA kernel with just 1 CUDA block and RTE_HCDEV_COMM_LIST_PKTS_MAX CUDA threads. */
    cuda_kernel_packet_processing<<<1, RTE_HCDEV_COMM_LIST_PKTS_MAX, 0, cstream>>>(quit_flag->ptr, comm_list, num_entries, ...);

    /*
     * For simplicity, the CPU here receives only 2 bursts of mbufs.
     * In a real application, network activity and device processing should overlap.
     */
    nb_rx = rte_eth_rx_burst(port_id, queue_id, &(rx_mbufs[0]), max_rx_mbufs);
    rte_hcdev_comm_populate_list_pkts(comm_list[0], rx_mbufs, nb_rx);
    nb_rx = rte_eth_rx_burst(port_id, queue_id, &(rx_mbufs[0]), max_rx_mbufs);
    rte_hcdev_comm_populate_list_pkts(comm_list[1], rx_mbufs, nb_rx);

    /*
     * CPU waits for the completion of the packets' processing on the CUDA kernel
     * and then it does a cleanup of the received mbufs.
     */
    while (rte_hcdev_comm_cleanup_list(comm_list[0]));
    while (rte_hcdev_comm_cleanup_list(comm_list[1]));

    /* CPU notifies the CUDA kernel that it has to terminate */
    rte_hcdev_comm_set_flag(&quit_flag, 1);

    /* hcdev objects cleanup/destruction */
    /* CUDA cleanup */
    /* DPDK cleanup */

    return 0;
}

////////////////////////////////////////////////////////////////////////
///// CUDA kernel
////////////////////////////////////////////////////////////////////////

void cuda_kernel(uint32_t * quit_flag_ptr, struct rte_hcdev_comm_list *comm_list, int comm_list_entries) {
    int comm_list_index = 0;
    struct rte_hcdev_comm_pkt *pkt_list = NULL;

    /* Do some pre-processing operations. */

    /* GPU kernel keeps checking this flag to know if it has to quit or wait for more packets. */
    while (*quit_flag_ptr == 0)
    {
        if (comm_list[comm_list_index]->status != RTE_HCDEV_COMM_LIST_READY)
            continue;

        if (threadIdx.x < comm_list[comm_list_index]->num_pkts)
        {
            /* Each CUDA thread processes a different packet. */
            packet_processing(comm_list[comm_list_index]->addr, comm_list[comm_list_index]->size, ..);
        }
        __threadfence();
        __syncthreads();

        /* Wait for new packets on the next communication list entry. */
        comm_list_index = (comm_list_index+1) % comm_list_entries;
    }

    /* Do some post-processing operations. */
}


-- 
2.31.1