From patchwork Fri Jun 28 10:24:48 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Andrew Stubbs X-Patchwork-Id: 1953876 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=P46pluri; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; 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 [8.43.85.97]) (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 4W9Wq43xzGz20X6 for ; Fri, 28 Jun 2024 20:28:00 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id BCA7D3884573 for ; Fri, 28 Jun 2024 10:27:58 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x233.google.com (mail-lj1-x233.google.com [IPv6:2a00:1450:4864:20::233]) by sourceware.org (Postfix) with ESMTPS id 1FB15382DB0E for ; Fri, 28 Jun 2024 10:25:03 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 1FB15382DB0E 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 1FB15382DB0E Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::233 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1719570305; cv=none; b=wr+8FRe1jH6YeLgqiqC8StkAeWWkBHVBMPeoIV9q8fCiSR/+2lEPRSv2LBbLZsTDjpNWfEq1qinfK1ivPKQmRd/ldDWbHmGu7Gp6Iaep59dfGyaaXbhpB3IS4PB8941SONRl3asRXDOKz0giiKpbGJC93+Dvh/jtZwJAPe1A9kc= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1719570305; c=relaxed/simple; bh=8ttmCGyV8LcVe2VlEfNhExLucXnQ13tq18dMY96ycMM=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=Zu5l8EiKJj9gyKCJdPXx4JHlCam3IRak33+UDXsBeMVupxjHyOBjcLodUmMXsaLxQ6AIRhbDOR41XnqNIvP13Y/JVm+ud7JoGNHyCSy6T5gATDEftx8+pU/FED9RSJ807CtSwA288ar694d/bJ6ZzT8B5It8gUOvUGZoOmyUyJw= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x233.google.com with SMTP id 38308e7fff4ca-2ebe40673d8so5067271fa.3 for ; Fri, 28 Jun 2024 03:25:03 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1719570301; x=1720175101; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to; bh=ENzaiJgOXetvpnUdxxTBCkszzOLr32xjzoeZyvQJODA=; b=P46pluril2ZZ9//cnF11nGHWC0LVNrGu+VpTa5ICzcvDHRdPf/4ysz0QxWuKHNOhPx /1001C4Sz17zPgfJmUgc9H/X8WebJMrBxa/cOkFqGNaskGwL6+8cIA+JSbzQ/doozbdD WboPOBTX76ulWuMJXAWEELJ1ZSbaLqI1ZGYhZPqKqWkLchAwrp0+vJdDQ7OutDf81jI+ 3qRicMiRE1akDkJWE1QDPqP8noy1agPD24FJZg0d+7nLyG8s1MtTbzyGA9Dp3sFVCu7G mKHWEt/jMCdA+DRTUJc74YwkkvHJcelQMt6N6E7zNMcL+lLBaXSP+ZpRLpEztc3DeH9n H71Q== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1719570301; x=1720175101; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:x-gm-message-state:from:to:cc :subject:date:message-id:reply-to; bh=ENzaiJgOXetvpnUdxxTBCkszzOLr32xjzoeZyvQJODA=; b=rm6CJ6wFYeMd50gMJx66/3LD9/cQxL+GVHHsRMt8hfn5JU46embTKeG7ojwa6+WXdz NWPsvNo8TA2Of4ZiO9zh6fS/d4SBmiWk+UfYNz7hnhRYcT+F+NvFXwuEFuyjxh5fWCWV 250lfUlZRTHxX+JOUcs0uE5tIaEN+f+cUtLee+pkvQLg1fUlHil/UZ/4jWkMrGiGFMLs 56+4EAKmKee+ctwSXL8aP7NlQAb5pTQIMWqY25Kvin7PmU0zkevPo7tV4GWAtvllWgwz ae0rWmYF368IEf1zngQg9G+6KKhD0tRmaHj3Y4K1yD8D5KGG+z3YfXJjaB4gzvO19sZB Rygg== X-Gm-Message-State: AOJu0Ywth0EabL8+fm+k0IKC5xRK5MgD/u97CP4/haTgnckpX3V0Ctk8 /GE0UdyRT9e34qclN2EebX3gRXxNl1yijCv0jyl4zsVfHgrGv7VLJJ4rblPdAzdy67QUxYBQGYm lBy0= X-Google-Smtp-Source: AGHT+IE6qDwboTsvM8DgcnaBafjzgZj85GxELJPLIjXo5q1BRcHv2IMaeH+obn87EJGOzBCZlNQ58Q== X-Received: by 2002:a2e:9054:0:b0:2ed:59af:ecb7 with SMTP id 38308e7fff4ca-2ed59afeef3mr45127601fa.15.1719570300848; Fri, 28 Jun 2024 03:25:00 -0700 (PDT) Received: from arnold.baylibre (88-127-129-70.subs.proxad.net. [88.127.129.70]) by smtp.googlemail.com with ESMTPSA id 5b1f17b1804b1-4256b061006sm28014945e9.22.2024.06.28.03.25.00 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 28 Jun 2024 03:25:00 -0700 (PDT) From: Andrew Stubbs To: gcc-patches@gcc.gnu.org Cc: tburnus@baylibre.com, jakub@redhat.com Subject: [PATCH v2 7/8] openmp, libgomp: Handle unified shared memory in omp_target_is_accessible Date: Fri, 28 Jun 2024 10:24:48 +0000 Message-ID: <20240628102449.562467-8-ams@baylibre.com> X-Mailer: git-send-email 2.41.0 In-Reply-To: <20240628102449.562467-1-ams@baylibre.com> References: <20240628102449.562467-1-ams@baylibre.com> MIME-Version: 1.0 X-Spam-Status: No, score=-10.5 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP 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 From: Marcel Vollweiler This patch handles Unified Shared Memory (USM) in the OpenMP runtime routine omp_target_is_accessible. libgomp/ChangeLog: * target.c (omp_target_is_accessible): Handle unified shared memory. * testsuite/libgomp.c-c++-common/target-is-accessible-1.c: Updated. * testsuite/libgomp.fortran/target-is-accessible-1.f90: Updated. * testsuite/libgomp.c-c++-common/target-is-accessible-2.c: New test. * testsuite/libgomp.fortran/target-is-accessible-2.f90: New test. --- libgomp/target.c | 8 +++++-- .../target-is-accessible-1.c | 22 +++++++++++++------ .../target-is-accessible-2.c | 21 ++++++++++++++++++ .../target-is-accessible-1.f90 | 20 +++++++++++------ .../target-is-accessible-2.f90 | 22 +++++++++++++++++++ 5 files changed, 77 insertions(+), 16 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c create mode 100644 libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90 diff --git a/libgomp/target.c b/libgomp/target.c index 754dea4e031..f0ee2c84197 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -5281,9 +5281,13 @@ omp_target_is_accessible (const void *ptr, size_t size, int device_num) if (devicep == NULL) return false; - /* TODO: Unified shared memory must be handled when available. */ + if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return true; - return devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM; + if (devicep->is_usm_ptr_func && devicep->is_usm_ptr_func ((void *) ptr)) + return true; + + return false; } int diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c index 2e75c6300ae..e7f9cf27a42 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c @@ -1,3 +1,5 @@ +/* { dg-do run } */ + #include int @@ -6,7 +8,8 @@ main () int d = omp_get_default_device (); int id = omp_get_initial_device (); int n = omp_get_num_devices (); - void *p; + int i = 42; + void *p = &i; if (d < 0 || d >= n) d = id; @@ -26,23 +29,28 @@ main () if (omp_target_is_accessible (p, sizeof (int), n + 1)) __builtin_abort (); - /* Currently, a host pointer is accessible if the device supports shared - memory or omp_target_is_accessible is executed on the host. This - test case must be adapted when unified shared memory is avialable. */ int a[128]; for (int d = 0; d <= omp_get_num_devices (); d++) { + /* SHARED_MEM is 1 if and only if host and device share the same memory. + OMP_TARGET_IS_ACCESSIBLE should not return 0 for shared memory. */ int shared_mem = 0; #pragma omp target map (alloc: shared_mem) device (d) shared_mem = 1; - if (omp_target_is_accessible (p, sizeof (int), d) != shared_mem) + + if (shared_mem && !omp_target_is_accessible (p, sizeof (int), d)) + __builtin_abort (); + + /* USM is disabled by default. Hence OMP_TARGET_IS_ACCESSIBLE should + return 0 if shared_mem is false. */ + if (!shared_mem && omp_target_is_accessible (p, sizeof (int), d)) __builtin_abort (); - if (omp_target_is_accessible (a, 128 * sizeof (int), d) != shared_mem) + if (shared_mem && !omp_target_is_accessible (a, 128 * sizeof (int), d)) __builtin_abort (); for (int i = 0; i < 128; i++) - if (omp_target_is_accessible (&a[i], sizeof (int), d) != shared_mem) + if (shared_mem && !omp_target_is_accessible (&a[i], sizeof (int), d)) __builtin_abort (); } diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c new file mode 100644 index 00000000000..24c77232f5d --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c @@ -0,0 +1,21 @@ +/* { dg-do run } */ +/* { dg-require-effective-target omp_usm } */ + +#include + +#pragma omp requires unified_shared_memory + +int +main () +{ + int *a = (int *) omp_alloc (sizeof (int), ompx_gnu_unified_shared_mem_alloc); + if (!a) + __builtin_abort (); + + for (int d = 0; d <= omp_get_num_devices (); d++) + if (!omp_target_is_accessible (a, sizeof (int), d)) + __builtin_abort (); + + omp_free(a, ompx_gnu_unified_shared_mem_alloc); + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 index 150df6f8a4f..0df43aae095 100644 --- a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 @@ -1,3 +1,5 @@ +! { dg-do run } + program main use omp_lib use iso_c_binding @@ -28,24 +30,28 @@ program main if (omp_target_is_accessible (p, c_sizeof (d), n + 1) /= 0) & stop 5 - ! Currently, a host pointer is accessible if the device supports shared - ! memory or omp_target_is_accessible is executed on the host. This - ! test case must be adapted when unified shared memory is avialable. do d = 0, omp_get_num_devices () + ! SHARED_MEM is 1 if and only if host and device share the same memory. + ! OMP_TARGET_IS_ACCESSIBLE should not return 0 for shared memory. shared_mem = 0; !$omp target map (alloc: shared_mem) device (d) shared_mem = 1; !$omp end target - if (omp_target_is_accessible (p, c_sizeof (d), d) /= shared_mem) & + if (shared_mem == 1 .and. omp_target_is_accessible (p, c_sizeof (d), d) == 0) & stop 6; - if (omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) /= shared_mem) & + ! USM is disabled by default. Hence OMP_TARGET_IS_ACCESSIBLE should + ! return 0 if shared_mem is false. + if (shared_mem == 0 .and. omp_target_is_accessible (p, c_sizeof (d), d) /= 0) & stop 7; + if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) == 0) & + stop 8; + do i = 1, 128 - if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /= shared_mem) & - stop 8; + if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) == 0) & + stop 9; end do end do diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90 new file mode 100644 index 00000000000..66e5a632961 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90 @@ -0,0 +1,22 @@ +! { dg-do run } +! { dg-require-effective-target omp_usm } + +program main + use omp_lib + use iso_c_binding + implicit none (external, type) + integer :: d + type(c_ptr) :: p + + !$omp requires unified_shared_memory + + p = omp_alloc (sizeof (d), ompx_gnu_unified_shared_mem_alloc) + if (.not. c_associated (p)) stop 1 + + do d = 0, omp_get_num_devices () + if (omp_target_is_accessible (p, c_sizeof (d), d) == 0) & + stop 2; + end do + + call omp_free (p, ompx_gnu_unified_shared_mem_alloc); +end program main