From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from mails.dpdk.org (mails.dpdk.org [217.70.189.124]) by inbox.dpdk.org (Postfix) with ESMTP id C20CDA0C53; Wed, 3 Nov 2021 12:06:02 +0100 (CET) Received: from [217.70.189.124] (localhost [127.0.0.1]) by mails.dpdk.org (Postfix) with ESMTP id 32F6D41203; Wed, 3 Nov 2021 12:05:09 +0100 (CET) Received: from NAM04-MW2-obe.outbound.protection.outlook.com (mail-mw2nam08on2072.outbound.protection.outlook.com [40.107.101.72]) by mails.dpdk.org (Postfix) with ESMTP id 703FE411FE for ; Wed, 3 Nov 2021 12:05:02 +0100 (CET) ARC-Seal: i=1; a=rsa-sha256; s=arcselector9901; d=microsoft.com; cv=none; b=JY/s8kdgY7akIEHfHGu667/G0z5usf+pLHtKVFpm9ZGHgMDU/OTy0pK6o0P87pxcaL63ujQQfZGbLOdCLFNDe+SpvmeN21G59PKA6y3FVTP/uMxG58waXxVwQO734Idvl3wxectYltZZUszKB3vrXjf4NJgxtiLkG6qSSC4CEr7MPa1yWGw7pVuI5gEAPhWctiPE+9u7E8LSJ9xRv0JRKww+X9BScs5yrMs6b/07byPIgvm9iLIv4fiGECJ9pvrcL8QyE4o0RyfZWXfSiKR/LfQRH53p495z6OGb4y2HioAODajimIwnK9a6qhYVr8gRGQmv4AJMcyP9TOGx3VIIHg== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=microsoft.com; s=arcselector9901; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-AntiSpam-MessageData-ChunkCount:X-MS-Exchange-AntiSpam-MessageData-0:X-MS-Exchange-AntiSpam-MessageData-1; bh=OxtiHhow+oMMccuwlq9p6l5sIJT9EIjR/wpVQcG82Kw=; b=jGpjLdoiAhvRDq0onCgVOT/fPAGBb4Npe2VpvlmmrQXgfolAsBTa/gWESreaN/+51HP83B9QJ3Wbu3aU4AB16CQCvp3hPTXalZz8ByczCF7racXcTl4l2u7uW2aC70T0mDmhpNZl/4OIHtwqjLL74Q+CCeKIho0opz3KRb8K1pOxYEyurrfQBauo8XcUIgGYlHfaxbnMV2ZhF0xA1EJplycj3Iph1xU1En1e5lXdY/MA4KlCd4P0j5I1nLsPEytCl9xmA5T1Z+zYHMHKx4587e0h/VDb4uFmpap1LwDZWhKVeDoI0P4cHLNA8lgRVv/ZgPNbCHSKFquUjxtvPknX4g== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass (sender ip is 216.228.112.34) smtp.rcpttodomain=dpdk.org smtp.mailfrom=nvidia.com; dmarc=pass (p=quarantine sp=quarantine pct=100) action=none header.from=nvidia.com; dkim=none (message not signed); arc=none DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=Nvidia.com; s=selector2; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-SenderADCheck; bh=OxtiHhow+oMMccuwlq9p6l5sIJT9EIjR/wpVQcG82Kw=; b=kFLcdkyPpysBxvj0ST+uSMJ8sNbAtcj27MoGsp3z6ier7kEtWLrruWvpxskk7Yc3iP2HGZv0MBz0+QNgRT7JBoudVLPbh6MfCkaSEiikbRItUZmUuBBIxzAHBFaAiAcMLwu+SwmCcpZypG/ITEPPrXMcacGFtbwh8MY5ie7vTKxzAVG/JsLIsLzRhzG/wfOk16FeopYXfh1KWVPs0lx5RZeRjlCB1aOWrKmwW8JTTPJOc0s10p7HtDw4djdKYvzRy0b1Z6h1HPdSV8H8tMP1kjcqVi8OtqSyDLgg9pU2WMI9qJY0/Sva506va4nFu4IvusJP2Uh6cemy6c51OooSpA== Received: from DM6PR03CA0087.namprd03.prod.outlook.com (2603:10b6:5:333::20) by CH2PR12MB4294.namprd12.prod.outlook.com (2603:10b6:610:a9::11) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.4649.15; Wed, 3 Nov 2021 11:05:00 +0000 Received: from DM6NAM11FT064.eop-nam11.prod.protection.outlook.com (2603:10b6:5:333:cafe::7b) by DM6PR03CA0087.outlook.office365.com (2603:10b6:5:333::20) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.4649.15 via Frontend Transport; Wed, 3 Nov 2021 11:05:00 +0000 X-MS-Exchange-Authentication-Results: spf=pass (sender IP is 216.228.112.34) smtp.mailfrom=nvidia.com; dpdk.org; dkim=none (message not signed) header.d=none;dpdk.org; dmarc=pass action=none header.from=nvidia.com; Received-SPF: Pass (protection.outlook.com: domain of nvidia.com designates 216.228.112.34 as permitted sender) receiver=protection.outlook.com; client-ip=216.228.112.34; helo=mail.nvidia.com; Received: from mail.nvidia.com (216.228.112.34) by DM6NAM11FT064.mail.protection.outlook.com (10.13.172.234) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_CBC_SHA384) id 15.20.4669.10 via Frontend Transport; Wed, 3 Nov 2021 11:05:00 +0000 Received: from nvidia.com (172.20.187.5) by HQMAIL107.nvidia.com (172.20.187.13) with Microsoft SMTP Server (TLS) id 15.0.1497.18; Wed, 3 Nov 2021 11:04:56 +0000 From: To: CC: Elena Agostini Date: Wed, 3 Nov 2021 19:15:54 +0000 Message-ID: <20211103191554.16449-10-eagostini@nvidia.com> X-Mailer: git-send-email 2.17.1 In-Reply-To: <20211103191554.16449-1-eagostini@nvidia.com> References: <20210602203531.2288645-1-thomas@monjalon.net> <20211103191554.16449-1-eagostini@nvidia.com> MIME-Version: 1.0 Content-Type: text/plain X-Originating-IP: [172.20.187.5] X-ClientProxiedBy: HQMAIL111.nvidia.com (172.20.187.18) To HQMAIL107.nvidia.com (172.20.187.13) X-EOPAttributedMessage: 0 X-MS-PublicTrafficType: Email X-MS-Office365-Filtering-Correlation-Id: c6941581-9763-4710-46f1-08d99eb9c776 X-MS-TrafficTypeDiagnostic: CH2PR12MB4294: X-Microsoft-Antispam-PRVS: X-MS-Oob-TLC-OOBClassifiers: OLM:6430; X-MS-Exchange-SenderADCheck: 1 X-MS-Exchange-AntiSpam-Relay: 0 X-Microsoft-Antispam: BCL:0; X-Microsoft-Antispam-Message-Info: 0QI3brHn0JG53dTIoy9ZHdcB9+OwjAsSJEDcoq8zRRaCmOOJ5zJP77TprDIFHe84AE2gk80+RdzmaUt+3b8R4RNhh+s+QeS2U30QrRacbqXfTJc5qDAvHpwMnvViRioRX/9yKu0lho/aORWjsDoV0hItGRV1nfMJdrEGW0dqG//y/Ke55uTSVvMpIq4Kp+smJrQTj+yloZZBqB+JTzLeAcTeHRVIJOK2N/V00H+tZXILgGEgdPjwXzKCiXF20DOlEotsP+eLqebmpRU9M8SfCdi+26jpFle0iPDzWwhtsSYWGuIEJ0FxyvB/HoqncJxm56DCf87NLAmXT4eZzOp1oesqy1RQB0eOdDqTv4UNXuLkTheHJ1KuMS9oq9huTgVkVVMSQFuqZ1u4N32ZuwNlS6hsrf2oRCUnXieLVbbsRCLRNDkC5AzjTzZkc0ulV+RTa2+3XfIEPddIl0wEI9xCaOeMKAqCuS7/6hNhO6v4whY8hgiAZ9FfMQT7VMhlY+o/pDwL6hCqgbNkTscorAu+3u88H51B20B726gLWulnyel5gmVAxIrGjQHzlIFK/Z2h7DCqJKMbfSFp6VlnYz3t1ujkuIDGuflewXTMB9MoFxIGmzjArl2zpdIDBvp7CoeU6dJNf0nwjUVzF5A1Q1chPNoR7po5+QRM0rTd5btVRV4VtScQefPtLZnOk5decnWWPr2J53tieTjO6vZKQ6yXPQ== X-Forefront-Antispam-Report: CIP:216.228.112.34; CTRY:US; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:mail.nvidia.com; PTR:schybrid03.nvidia.com; CAT:NONE; SFS:(4636009)(46966006)(36840700001)(2876002)(1076003)(356005)(55016002)(70206006)(86362001)(6916009)(316002)(8936002)(508600001)(70586007)(16526019)(26005)(82310400003)(336012)(186003)(36906005)(2616005)(6286002)(426003)(36756003)(36860700001)(7696005)(2906002)(5660300002)(83380400001)(7636003)(4326008)(107886003)(47076005)(8676002); DIR:OUT; SFP:1101; X-OriginatorOrg: Nvidia.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 03 Nov 2021 11:05:00.0312 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: c6941581-9763-4710-46f1-08d99eb9c776 X-MS-Exchange-CrossTenant-Id: 43083d15-7273-40c1-b7db-39efd9ccc17a X-MS-Exchange-CrossTenant-OriginalAttributedTenantConnectingIp: TenantId=43083d15-7273-40c1-b7db-39efd9ccc17a; Ip=[216.228.112.34]; Helo=[mail.nvidia.com] X-MS-Exchange-CrossTenant-AuthSource: DM6NAM11FT064.eop-nam11.prod.protection.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: CH2PR12MB4294 Subject: [dpdk-dev] [PATCH v4 9/9] doc: add CUDA example in GPU guide X-BeenThere: dev@dpdk.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: DPDK patches and discussions List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: dev-bounces@dpdk.org Sender: "dev" From: Elena Agostini Signed-off-by: Elena Agostini --- doc/guides/prog_guide/gpudev.rst | 122 +++++++++++++++++++++++++++++++ 1 file changed, 122 insertions(+) diff --git a/doc/guides/prog_guide/gpudev.rst b/doc/guides/prog_guide/gpudev.rst index cbaec5a1e4..1baf0c6772 100644 --- a/doc/guides/prog_guide/gpudev.rst +++ b/doc/guides/prog_guide/gpudev.rst @@ -102,3 +102,125 @@ the list of mbuf payload addresses where received packet have been stored. The ``rte_gpu_comm_*()`` functions are responsible to create a list of packets that can be populated with receive mbuf payload addresses and communicated to the task running on the GPU. + + +CUDA Example +------------ + +In the example below, there is a pseudo-code to give an example +about how to use functions in this library in case of a CUDA application. + +.. code-block:: c + + ////////////////////////////////////////////////////////////////////////// + ///// gpudev library + CUDA functions + ////////////////////////////////////////////////////////////////////////// + #define GPU_PAGE_SHIFT 16 + #define GPU_PAGE_SIZE (1UL << GPU_PAGE_SHIFT) + + int main() { + struct rte_gpu_flag quit_flag; + struct rte_gpu_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; + int16_t port_id = 0; + + /** Initialize CUDA objects (cstream, context, etc..). */ + /** Use gpudev 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_gpu_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[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_gpu_comm_create_flag(dev_id, &quit_flag, RTE_GPU_COMM_FLAG_CPU); + rte_gpu_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_gpu_comm_create_list(dev_id, num_entries); + + /** A very simple CUDA kernel with just 1 CUDA block and RTE_GPU_COMM_LIST_PKTS_MAX CUDA threads. */ + cuda_kernel_packet_processing<<<1, RTE_GPU_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_gpu_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_gpu_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_gpu_comm_cleanup_list(comm_list[0])); + while(rte_gpu_comm_cleanup_list(comm_list[1])); + + /** CPU notifies the CUDA kernel that it has to terminate */ + rte_gpu_comm_set_flag(&quit_flag, 1); + + /** gpudev objects cleanup/destruction */ + /** CUDA cleanup */ + + rte_gpu_free(dev_id, ext_mem.buf_len); + + /** DPDK cleanup */ + + return 0; + } + + ////////////////////////////////////////////////////////////////////////// + ///// CUDA kernel + ////////////////////////////////////////////////////////////////////////// + + void cuda_kernel(uint32_t * quit_flag_ptr, struct rte_gpu_comm_list *comm_list, int comm_list_entries) { + int comm_list_index = 0; + struct rte_gpu_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_GPU_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.17.1