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 4DBD1A0C4D; Mon, 8 Nov 2021 11:48:15 +0100 (CET) Received: from [217.70.189.124] (localhost [127.0.0.1]) by mails.dpdk.org (Postfix) with ESMTP id 8FEE241160; Mon, 8 Nov 2021 11:47:16 +0100 (CET) Received: from NAM12-DM6-obe.outbound.protection.outlook.com (mail-dm6nam12on2082.outbound.protection.outlook.com [40.107.243.82]) by mails.dpdk.org (Postfix) with ESMTP id 2206D41135 for ; Mon, 8 Nov 2021 11:47:12 +0100 (CET) ARC-Seal: i=1; a=rsa-sha256; s=arcselector9901; d=microsoft.com; cv=none; b=YcLK8JYYE23Sf384Hufl3pkOdhu/q7UI5dTd0vzpteNV8IfYScUzfOcDhjbSgK8AZEHZ8ubXjOVDs71neViEh5p6I0uAv3ZrJGWeyAxKbWJCWr2+T36NZ9Z88TSueUDoyR0creJZZHnuuAdGM6sMy0pE4VTnhe/aHEdp7/OYlEnCDZYemRrfycWrnxgesgfMd2sfGr/+vc2Lqxn6QwOmxL/kD9X0BZwrBUntAxvP7lOjkmCHFQSszK9pXUBj1VaOKYwVj6zNuA9DpuJ4n4AAleQ0uXUqd0ZOwIV0otAmGxPJwcqT54+DQ+TdMR+wCMcXTQclqg8ymJ8z8VHAJqapug== 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=ab0b/L2cjd7jwKCx2i+8/ZLSTTodqJ3S23PKykTRaxlPULHn2L/0VUM5xjfz4qm/uAacJManN1+edSbYkbFok+XJEoxxeAh9UO/aMJpnSFsv5Wdg72vmrWcf/6Rsm62lcZyWPJsZfFhgwQOpI6aO9sSYTLGXJrkPAvbG5taQ9DaXUojtCDLl49fqW60yUZHSAiXR5WMxZazffEnZF/yJOBAe8I6Y90M7Xk0WDHbn9xXScA3yk6OngiS0u8+8ouZwufQXyRhEnRklaot5EWOaMzvlktgBZGWI4JY4k51e3Sq4JAyaG759Ga3z9s91d562S3YThLiviqxjSc0kd6WH+g== 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=mNQ2s0gO3/+0HlGrzjc1TY2N9fVZPOscc/+TkqbfzOdnLvMvdkU3LDLAvU+3VQ4JoAu6/eBxZonKw4FoKB32Iv40kLr2aqhokHzSCczM6R1j+O3SlevIiOtUp7Msn6OYAMC88bB4SSkhdbkvB4o+YFaBu/h1h95qOLLeIYqiqYoqPYOGnOVmEASS/DVdaoEH+15tFZ7Imup37Vnn9ARLg52rEwMiPP1bq2UDFx1CvIFHswY1Dw8mRQLN73P0N7tUFgMAQathP3Z0UZHYOhT2YgBWQQdCMjG7yu0u/WCsAlHwOp1NsSWn4apk6Z28cxjkqLAZV1VBO1e/HZ4/KhcGFQ== Received: from MW2PR16CA0052.namprd16.prod.outlook.com (2603:10b6:907:1::29) by MN2PR12MB4286.namprd12.prod.outlook.com (2603:10b6:208:199::22) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.4669.11; Mon, 8 Nov 2021 10:47:11 +0000 Received: from CO1NAM11FT064.eop-nam11.prod.protection.outlook.com (2603:10b6:907:1:cafe::1f) by MW2PR16CA0052.outlook.office365.com (2603:10b6:907:1::29) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.4669.10 via Frontend Transport; Mon, 8 Nov 2021 10:47:10 +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 CO1NAM11FT064.mail.protection.outlook.com (10.13.175.77) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_CBC_SHA384) id 15.20.4669.10 via Frontend Transport; Mon, 8 Nov 2021 10:47:10 +0000 Received: from nvidia.com (172.20.187.6) by HQMAIL107.nvidia.com (172.20.187.13) with Microsoft SMTP Server (TLS) id 15.0.1497.18; Mon, 8 Nov 2021 10:47:02 +0000 From: To: CC: Elena Agostini Date: Mon, 8 Nov 2021 18:58:05 +0000 Message-ID: <20211108185805.3887-10-eagostini@nvidia.com> X-Mailer: git-send-email 2.17.1 In-Reply-To: <20211108185805.3887-1-eagostini@nvidia.com> References: <20210602203531.2288645-1-thomas@monjalon.net> <20211108185805.3887-1-eagostini@nvidia.com> MIME-Version: 1.0 Content-Type: text/plain X-Originating-IP: [172.20.187.6] X-ClientProxiedBy: HQMAIL105.nvidia.com (172.20.187.12) To HQMAIL107.nvidia.com (172.20.187.13) X-EOPAttributedMessage: 0 X-MS-PublicTrafficType: Email X-MS-Office365-Filtering-Correlation-Id: e9d997b5-059e-4539-3f2f-08d9a2a51de7 X-MS-TrafficTypeDiagnostic: MN2PR12MB4286: 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: x5lmRAxraCZ+CvfCzB+rfg3862p9FIyBYqB/2I9c097MhlHqMxr38elTNATXFwrCxlfKHFXWGyUDe7/P+NkmqU3dIwSaDzp1Lc7ntPj0OrDEoJFaop2febJ+z38+qST2THS0077k1MqgrqX8cG2QjHecbfr2SwzDwWNYuf6mIYD46cXKiksF8zOiYsU+G3njTq27VnFhhX2JKh3M0fbFtfXBTNAALJ8jKrGOHx4XXZ4revOdIM3CWVUMEHEmGUxXE6lmfCnOj3rNlBcUfeIJXfq5ck5rZRftsybxCiUP2/lS16rFyiTUmhnNbt77UU8wS0d5/xeSS3RWedTOwAEjeyPUNuvXZ7l+fy3GeJ/FoFKn/JBgXnmzjPExnD7MA5sqWKwWX2rdH1zTOkahZko7M34zdFSJEFwAPa7uPfNlwCPKfsTKFFRvNlzwuY2FW4Dqw2N2ul2UNdH6pg0ZF8t7kGNMHfXcseQbbkEEYG2rFHurm8wVN+Tkw/8vDgU5Wd8XbUEJ8VR7h+dWHKLxPthlRqsZ9A51v6t21DXFfndlMMrEPjx15OXehRFoPSQXdIAMTmwGWmrFrrEQxqmV+YBtNPzCSle+u2x24tdSHB7kJfyTNJVQDBDSnbKrpcxfHsXOALYSQLp8gyaqK7ZvJ5h7kz6TsDeLFpz/2vLjnRHg+KbxwApSumAm4C0/aqgzjP4mkEFybkP8eUNOj8UaY+OE9A== 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)(336012)(55016002)(86362001)(7696005)(36860700001)(16526019)(186003)(2616005)(8936002)(6916009)(426003)(36756003)(82310400003)(26005)(316002)(8676002)(70586007)(70206006)(508600001)(83380400001)(47076005)(356005)(1076003)(7636003)(6286002)(2906002)(107886003)(4326008)(5660300002); DIR:OUT; SFP:1101; X-OriginatorOrg: Nvidia.com X-MS-Exchange-CrossTenant-OriginalArrivalTime: 08 Nov 2021 10:47:10.3554 (UTC) X-MS-Exchange-CrossTenant-Network-Message-Id: e9d997b5-059e-4539-3f2f-08d9a2a51de7 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: CO1NAM11FT064.eop-nam11.prod.protection.outlook.com X-MS-Exchange-CrossTenant-AuthAs: Anonymous X-MS-Exchange-CrossTenant-FromEntityHeader: HybridOnPrem X-MS-Exchange-Transport-CrossTenantHeadersStamped: MN2PR12MB4286 Subject: [dpdk-dev] [PATCH v5 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