From patchwork Mon Jan 29 17:48:47 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Kwok Cheung Yeung X-Patchwork-Id: 1892442 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; dkim=pass (2048-bit key; unprotected) header.d=baylibre-com.20230601.gappssmtp.com header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256 header.s=20230601 header.b=uvEJjnF6; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4TNwm501Ntz23fN for ; Tue, 30 Jan 2024 04:49:24 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 115F93858422 for ; Mon, 29 Jan 2024 17:49:22 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x32a.google.com (mail-wm1-x32a.google.com [IPv6:2a00:1450:4864:20::32a]) by sourceware.org (Postfix) with ESMTPS id 9423B3858C60 for ; Mon, 29 Jan 2024 17:48:52 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 9423B3858C60 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 9423B3858C60 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::32a ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1706550536; cv=none; b=hSHwlOW3oZVnz09JYh2b2apBCR3JC2wEUvLKi9n+gwsmRPpkEUUkiozMSsjQoZUgdmn6hdk+xdZLN7Ai9GjUzKZX3Iy8xwOrHBxThF2yseJBdAJyO/7kG85398OpNqoDn7hxMnVxRiuq5Om2LWFOeuPOHdJjyOjyaAhj8A2GDa8= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1706550536; c=relaxed/simple; bh=X9kw/s9HW8xtoMK3Xk3sVg9KULwFmtkG46mdBlho6xI=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=kNe27rWZ5GBpcjlK+v2wMeYj9hOEHDC5/ATy56XLAdTzK1rtQNwdCKcWvHQMJalIokGtQHIJk4zkNVL1J8caom4AHzMbyEBc+lkKfMJhjXXImczorJKd8KcJB6eED4vCnavOL25HCAiD4xjO+n08yBhpU0JpnQUVfmirkJM0eYE= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x32a.google.com with SMTP id 5b1f17b1804b1-40eb033c1b0so36899815e9.2 for ; Mon, 29 Jan 2024 09:48:52 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1706550531; x=1707155331; darn=gcc.gnu.org; h=in-reply-to:from:cc:content-language:subject:references:to :user-agent:mime-version:date:message-id:from:to:cc:subject:date :message-id:reply-to; bh=QE0BSiWOcPvwEaRLScIYTKGEijnY38ZhTD+5937APJ4=; b=uvEJjnF6YCJETZSytuxDZrwAxDoW7VpClfB2rqD7yAwmtyYvdQEIaUjjVghaI6ks/X uGH1YsDM20KMTOPz2TkeKO6N2KyYj5MjuYZAPv9AotF7E0XDPvMftyogU7QQWYmmA9t7 tbDJ2f8KXKN/AORjZitrFH2asLPLuFBlGvKFiLwQFTnz7ipqVHN3sg+9WEr/hkaQohQX JLh1fKAiPJxqVcyHk8ABPeASuL9yDsBqfLADQTnKwBfkhPIcfG5GrcRoz9OhR6dOwYEM PaNFSXDC9yxQSGG/FtOANIu/mT9ivu3gz9ZO+8iw8IcIn8TVdcRyuzPqwdFeh9SVooVs /D3A== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1706550531; x=1707155331; h=in-reply-to:from:cc:content-language:subject:references:to :user-agent:mime-version:date:message-id:x-gm-message-state:from:to :cc:subject:date:message-id:reply-to; bh=QE0BSiWOcPvwEaRLScIYTKGEijnY38ZhTD+5937APJ4=; b=I9LarXfG8K1lrvx18VCLLGjtC8mCHihPAfFeXL0U/scbIkNcd+4SIxMs8+iVW1Cx8t pX65Xyy4Zj7wUa7fzvy7cXmuJEdnRrFxfA5yga8nCpklFMAzqg6z+4jmU/Ku4RT6tJuH nXfzosqXLY2D6Gvppt1zK3xWA+AzIYUPxpuSyTDlE7L9hZUoyhlgjiX9XDBDtDYJHGIp Lf25YmjVu6Ge/7PRSSfcb1Di7daahT5WDfSxOzT5a+wUsZ7VeXJm0wtFvWZkHrrqiiaM 20Dj/qA5vGVBOrXWVpQeX4wEvNLnJAEYXYiBTlFE3CScbw+KexS0pr7nBBOZWOKv7CKn kdvQ== X-Gm-Message-State: AOJu0Yzlgq4MpESjOCTYmC0n/z72l5dzS6fMLhJdbMpFgOJsa3n8ULcO EhUrc0xVv3iOQgKQ/k0MPcaemS2Hbld3aPoQgoceeMNOZ34sVsdNJmgo/hrj+TY= X-Google-Smtp-Source: AGHT+IE2v98KqmtuNTW71uAsMLguiT1QOvAEjSjdDgCLiY7mZeTG23Zy1Eu/amm0hveAQeP0tbeYgg== X-Received: by 2002:a05:600c:3c91:b0:40e:e834:3d86 with SMTP id bg17-20020a05600c3c9100b0040ee8343d86mr5112853wmb.37.1706550530926; Mon, 29 Jan 2024 09:48:50 -0800 (PST) Received: from ?IPV6:2a00:23c6:88e4:c502:ade8:41f3:51b6:ce93? ([2a00:23c6:88e4:c502:ade8:41f3:51b6:ce93]) by smtp.gmail.com with ESMTPSA id t18-20020a05600c451200b0040e880ac6ecsm14693528wmo.35.2024.01.29.09.48.50 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Mon, 29 Jan 2024 09:48:50 -0800 (PST) Message-ID: <679889de-bf47-4a01-887e-db96f7fad427@baylibre.com> Date: Mon, 29 Jan 2024 17:48:47 +0000 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird To: rep.dot.nop@gmail.com, gcc-patches References: Subject: [PATCH v2] openmp: Change to using a hashtab to lookup offload target addresses for indirect function calls Content-Language: en-GB Cc: Tobias Burnus , Jakub Jelinek From: Kwok Cheung Yeung In-Reply-To: X-Spam-Status: No, score=-13.5 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GB_TO_NAME_FREEMAIL, GIT_PATCH_0, KAM_SHORT, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org > Can you please akso update the comments to talk about hashtab instead of splay? > Hello This version has the comments updated and removes a stray 'volatile' in the #ifdefed out code. Thanks Kwok From 5737298f4f5e5471667b05e207b22c9c91b94ca0 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Mon, 29 Jan 2024 17:40:04 +0000 Subject: [PATCH 1/2] openmp: Change to using a hashtab to lookup offload target addresses for indirect function calls A splay-tree was previously used to lookup equivalent target addresses for a given host address on offload targets. However, as splay-trees can modify their structure on lookup, they are not suitable for concurrent access from separate teams/threads without some form of locking. This patch changes the lookup data structure to a hashtab instead, which does not have these issues. The call to build_indirect_map to initialize the data structure is now called from just the first thread of the first team to avoid redundant calls to this function. 2024-01-29 Kwok Cheung Yeung libgomp/ * config/accel/target-indirect.c: Include string.h and hashtab.h. Remove include of splay-tree.h. Update comments. (splay_tree_prefix, splay_tree_c): Delete. (struct indirect_map_t): New. (hash_entry_type, htab_alloc, htab_free, htab_hash, htab_eq): New. (GOMP_INDIRECT_ADD_MAP): Remove volatile qualifier. (USE_SPLAY_TREE_LOOKUP): Rename to... (USE_HASHTAB_LOOKUP): ..this. (indirect_map, indirect_array): Delete. (indirect_htab): New. (build_indirect_map): Remove locking. Build indirect map using hashtab. (GOMP_target_map_indirect_ptr): Use indirect_htab to lookup target address. (GOMP_target_map_indirect_ptr): Remove volatile qualifier. * config/gcn/team.c (gomp_gcn_enter_kernel): Call build_indirect_map from first thread of first team only. * config/nvptx/team.c (gomp_nvptx_main): Likewise. * testsuite/libgomp.c-c++-common/declare-target-indirect-2.c (main): Add missing break statements. --- libgomp/config/accel/target-indirect.c | 83 ++++++++++--------- libgomp/config/gcn/team.c | 7 +- libgomp/config/nvptx/team.c | 9 +- .../declare-target-indirect-2.c | 14 ++-- 4 files changed, 63 insertions(+), 50 deletions(-) diff --git a/libgomp/config/accel/target-indirect.c b/libgomp/config/accel/target-indirect.c index c60fd547cb6..cfef1ddbc49 100644 --- a/libgomp/config/accel/target-indirect.c +++ b/libgomp/config/accel/target-indirect.c @@ -25,60 +25,73 @@ . */ #include +#include #include "libgomp.h" -#define splay_tree_prefix indirect -#define splay_tree_c -#include "splay-tree.h" +struct indirect_map_t +{ + void *host_addr; + void *target_addr; +}; + +typedef struct indirect_map_t *hash_entry_type; + +static inline void * htab_alloc (size_t size) { return gomp_malloc (size); } +static inline void htab_free (void *ptr) { free (ptr); } + +#include "hashtab.h" + +static inline hashval_t +htab_hash (hash_entry_type element) +{ + return hash_pointer (element->host_addr); +} -volatile void **GOMP_INDIRECT_ADDR_MAP = NULL; +static inline bool +htab_eq (hash_entry_type x, hash_entry_type y) +{ + return x->host_addr == y->host_addr; +} -/* Use a splay tree to lookup the target address instead of using a - linear search. */ -#define USE_SPLAY_TREE_LOOKUP +void **GOMP_INDIRECT_ADDR_MAP = NULL; -#ifdef USE_SPLAY_TREE_LOOKUP +/* Use a hashtab to lookup the target address instead of using a linear + search. */ +#define USE_HASHTAB_LOOKUP -static struct indirect_splay_tree_s indirect_map; -static indirect_splay_tree_node indirect_array = NULL; +#ifdef USE_HASHTAB_LOOKUP -/* Build the splay tree used for host->target address lookups. */ +static htab_t indirect_htab = NULL; + +/* Build the hashtab used for host->target address lookups. */ void build_indirect_map (void) { size_t num_ind_funcs = 0; - volatile void **map_entry; - static int lock = 0; /* == gomp_mutex_t lock; gomp_mutex_init (&lock); */ + void **map_entry; if (!GOMP_INDIRECT_ADDR_MAP) return; - gomp_mutex_lock (&lock); - - if (!indirect_array) + if (!indirect_htab) { /* Count the number of entries in the NULL-terminated address map. */ for (map_entry = GOMP_INDIRECT_ADDR_MAP; *map_entry; map_entry += 2, num_ind_funcs++); - /* Build splay tree for address lookup. */ - indirect_array = gomp_malloc (num_ind_funcs * sizeof (*indirect_array)); - indirect_splay_tree_node array = indirect_array; + /* Build hashtab for address lookup. */ + indirect_htab = htab_create (num_ind_funcs); map_entry = GOMP_INDIRECT_ADDR_MAP; - for (int i = 0; i < num_ind_funcs; i++, array++) + for (int i = 0; i < num_ind_funcs; i++, map_entry += 2) { - indirect_splay_tree_key k = &array->key; - k->host_addr = (uint64_t) *map_entry++; - k->target_addr = (uint64_t) *map_entry++; - array->left = NULL; - array->right = NULL; - indirect_splay_tree_insert (&indirect_map, array); + struct indirect_map_t element = { *map_entry, NULL }; + hash_entry_type *slot = htab_find_slot (&indirect_htab, &element, + INSERT); + *slot = (hash_entry_type) map_entry; } } - - gomp_mutex_unlock (&lock); } void * @@ -88,15 +101,11 @@ GOMP_target_map_indirect_ptr (void *ptr) if (!ptr) return ptr; - assert (indirect_array); - - struct indirect_splay_tree_key_s k; - indirect_splay_tree_key node = NULL; - - k.host_addr = (uint64_t) ptr; - node = indirect_splay_tree_lookup (&indirect_map, &k); + assert (indirect_htab); - return node ? (void *) node->target_addr : ptr; + struct indirect_map_t element = { ptr, NULL }; + hash_entry_type entry = htab_find (indirect_htab, &element); + return entry ? entry->target_addr : ptr; } #else @@ -115,7 +124,7 @@ GOMP_target_map_indirect_ptr (void *ptr) assert (GOMP_INDIRECT_ADDR_MAP); - for (volatile void **map_entry = GOMP_INDIRECT_ADDR_MAP; *map_entry; + for (void **map_entry = GOMP_INDIRECT_ADDR_MAP; *map_entry; map_entry += 2) if (*map_entry == ptr) return (void *) *(map_entry + 1); diff --git a/libgomp/config/gcn/team.c b/libgomp/config/gcn/team.c index 61e9c616b67..bd3df448b52 100644 --- a/libgomp/config/gcn/team.c +++ b/libgomp/config/gcn/team.c @@ -52,14 +52,15 @@ gomp_gcn_enter_kernel (void) { int threadid = __builtin_gcn_dim_pos (1); - /* Initialize indirect function support. */ - build_indirect_map (); - if (threadid == 0) { int numthreads = __builtin_gcn_dim_size (1); int teamid = __builtin_gcn_dim_pos(0); + /* Initialize indirect function support. */ + if (teamid == 0) + build_indirect_map (); + /* Set up the global state. Every team will do this, but that should be harmless. */ gomp_global_icv.nthreads_var = 16; diff --git a/libgomp/config/nvptx/team.c b/libgomp/config/nvptx/team.c index 0cf5dad39ca..d5361917a24 100644 --- a/libgomp/config/nvptx/team.c +++ b/libgomp/config/nvptx/team.c @@ -60,9 +60,6 @@ gomp_nvptx_main (void (*fn) (void *), void *fn_data) asm ("mov.u32 %0, %%tid.y;" : "=r" (tid)); asm ("mov.u32 %0, %%ntid.y;" : "=r" (ntids)); - /* Initialize indirect function support. */ - build_indirect_map (); - if (tid == 0) { gomp_global_icv.nthreads_var = ntids; @@ -74,6 +71,12 @@ gomp_nvptx_main (void (*fn) (void *), void *fn_data) nvptx_thrs = alloca (ntids * sizeof (*nvptx_thrs)); memset (nvptx_thrs, 0, ntids * sizeof (*nvptx_thrs)); + /* Initialize indirect function support. */ + unsigned int block_id; + asm ("mov.u32 %0, %%ctaid.x;" : "=r" (block_id)); + if (block_id == 0) + build_indirect_map (); + /* Find the low-latency heap details .... */ uint32_t *shared_pool; uint32_t shared_pool_size = 0; diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c b/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c index 9fe190efce8..545f1a9fcbf 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c @@ -17,17 +17,17 @@ int main (void) { switch (i % 3) { - case 0: fn_ptr[i] = &foo; - case 1: fn_ptr[i] = &bar; - case 2: fn_ptr[i] = &baz; + case 0: fn_ptr[i] = &foo; break; + case 1: fn_ptr[i] = &bar; break; + case 2: fn_ptr[i] = &baz; break; } expected += (*fn_ptr[i]) (); } -#pragma omp target teams distribute parallel for reduction(+: x) \ - map (to: fn_ptr) map (tofrom: x) - for (int i = 0; i < N; i++) - x += (*fn_ptr[i]) (); + #pragma omp target teams distribute parallel for \ + reduction (+: x) map (to: fn_ptr) map (tofrom: x) + for (int i = 0; i < N; i++) + x += (*fn_ptr[i]) (); return x - expected; }