From patchwork Fri Feb 7 08:56:38 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 1234795 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-519112-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.a=rsa-sha1 header.s=default header.b=CprBFg9S; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 48DTh30bBJz9sSD for ; Fri, 7 Feb 2020 19:57:05 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=q3Z0CFQoHq5jXLY8LbN7AXjvR/Hs4QO4csmRTqX2KZbcVEwrPH vwEUrDDSf76ScXp/OoFoVVgvCJyLcCHCk+mJB5fui2ZWq4hRCxtJqrpdqBm977tS /79SHb++H4p3fNqJpbaJv9PD0V4dAmvHJwC61UeKoL9jmai6FeiuB29PQ= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:message-id:date:mime-version:content-type; s= default; bh=ovrho14QzXR/9sllcIjvuF0c5iw=; b=CprBFg9SFKrmXMexG+Ct pVkBlKQnorxA17m7H7YYdi0kKMTbN+nSit/ORUkvfNyq54yYVD3uZtgvVageu+hi uV6kTwqkZMZYy5lQ0Zp7cSfy6PkgTb/xQj1Gi3RiwMFQNq20hq+X3ukcTf2FjbHt RmBt20EP5RT0CAH5pfQ4mXU= Received: (qmail 130394 invoked by alias); 7 Feb 2020 08:56:56 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 130385 invoked by uid 89); 7 Feb 2020 08:56:56 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-20.9 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, KAM_SHORT autolearn=ham version=3.3.1 spammy=libgompc, libgomp.c, UD:libgomp.c, standing X-HELO: esa2.mentor.iphmx.com Received: from esa2.mentor.iphmx.com (HELO esa2.mentor.iphmx.com) (68.232.141.98) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 07 Feb 2020 08:56:53 +0000 IronPort-SDR: hHmpuKXXxv0vkURrEG8WgDYtbDVmyokRRfvtH+odUwTklgqNTwYl/ptVx0esjbtPhit0GcC5yt 4tvrUqqDWVtcXI/p8yznsgWDq9bZCfwlLx5xEew0qkPVAIa/LwnhoYF32gGPkKM/K3QE46Bw3d bcIkJPQFJg4ANvcxsGSIJ7o1cKmhF1RL/v/Sr3X7JSZsQVIRrcRBcNgPZUaj/JWN4/JV8LV17B U5fdsd/i1eSIPhCqWm2G5jBAIgZARF0EHBnWND8Z00J5Pn+3CB0ZpEKO+YYxqGMveedc5nOnHR pjM= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 07 Feb 2020 00:56:51 -0800 IronPort-SDR: 9sgUO6W71FOzwq+hx/irAIyax//iCxW5CnjMdJNaNcr7DZnKQynShHjOEKjR2VGIaWpdkYiXs7 iSlI2wEGZhE0w6VrvFaSH0LG7IVsDCSm3VzVehYgi3wPwPEli2sFs5liUSjzgQjdomKoyX/MqI SkKKvAK+o7TSyqsZEdANNJ5BTKGRkBaTYZQwkuU3lEbAuWcZJuk+iXweIAYceH32EfFFMc1PAi 1RX2RIzYbmp2EswuZBF2aplj/GFS+rUOI2SPf54Cam+XuA8Z8mb3g8Af40bwRGnTsQzLdebw4f rXs= From: "Harwath, Frederik" To: Jakub Jelinek , GCC Patches CC: Tobias Burnus , "Stubbs, Andrew" , , Thomas Schwinge Subject: [PATCH] xfail and improve some failing libgomp tests X-Pep-Version: 2.0 Message-ID: <1ea4c14f-f975-5170-46ac-ada4588c4969@codesourcery.com> Date: Fri, 7 Feb 2020 09:56:38 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.4.2 MIME-Version: 1.0 X-IsSubscribed: yes Hi, the libgomp testsuite contains some test cases (all in /libgomp/testsuite/libgomp.c/) which fail with nvptx offloading because of some long standing issues: * {target-32.c, thread-limit-2.c}: no "usleep" implemented for nvptx. Cf. https://gcc.gnu.org/PR81690 * target-{33,34}.c: no "GOMP_OFFLOAD_async_run" implemented in plugin-nvptx.c. Cf. https://gcc.gnu.org/PR81688 * target-link-1.c: omp "target link" not implemented for nvptx. Cf. https://gcc.gnu.org/PR81689 All these issues have been known, at least, since 2016: https://gcc.gnu.org/ml/gcc-patches/2016-11/msg00972.html As suggested in this mail: "Short term, it should be possible to implement something like -foffload=^nvptx to skip PTX (and only PTX) offloading on those tests." Well, we can now skip/xfail tests for nvptx offloading using the effective target "offload_target_nvptx" and the present patch uses this to xfail the tests for which no short-term solution is in sight, i.e. the GOMP_OFFLOAD_async_run and the "target link" related failures. Regarding the "usleep" issue, I have decided to follow Jakub's suggestion (cf. https://gcc.gnu.org/ml/gcc-patches/2016-11/msg01026.html) to replace usleep by busy waiting. As noted by Tobias (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=81690#c4), this involves creating separate test files for the cases with and without usleep. This solution is a bit cumbersome but I think we can live with it, in particular, since the actual test case implementations do not get duplicated (they have been moved into auxiliary header files which are shared by both variants of the corresponding tests). Since the "usleep" issue also concerns amdgcn, I have introduced an effective target "offload_target_amdgcn" to add xfails for this offloading target, too. This behaves like "offload_target_nvptx" but for amdgcn. Note that the existing amdgcn effective targets cannot be used for our purpose since they are OpenACC-specific. The new thread-limit-2-nosleep.c should now pass for both nvptx and amdgcn offloading whereas thread-limit-2.c should xfail. The new target-32-nosleep.c passes with amdgcn offloading, but xfails with nvptx offloading, because it also needs the unimplemented GOMP_OFFLOAD_async_run. With the patch, the detailed test summary now looks as follows for me: nvptx offloading: // Expected execution failures due to missing usleep PASS: libgomp.c/target-32-nosleep.c (test for excess errors) XFAIL: libgomp.c/target-32-nosleep.c execution test // missing GOMP_OFFLOAD_async_run XFAIL: libgomp.c/target-32.c (test for excess errors) UNRESOLVED: libgomp.c/target-32.c compilation failed to produce executable PASS: libgomp.c/thread-limit-2-nosleep.c (test for excess errors) PASS: libgomp.c/thread-limit-2-nosleep.c execution test XFAIL: libgomp.c/thread-limit-2.c (test for excess errors) UNRESOLVED: libgomp.c/thread-limit-2.c compilation failed to produce executable // Expected execution failures due to missing GOMP_OFFLOAD_async_run PASS: libgomp.c/target-33.c (test for excess errors) XFAIL: libgomp.c/target-33.c execution test PASS: libgomp.c/target-34.c (test for excess errors) XFAIL: libgomp.c/target-34.c execution test // Expected compilation failures due to missing target link XFAIL: libgomp.c/target-link-1.c (test for excess errors) UNRESOLVED: libgomp.c/target-link-1.c compilation failed to produce executable amdgcn offloading: // Tests using usleep PASS: libgomp.c/target-32-nosleep.c (test for excess errors) PASS: libgomp.c/target-32-nosleep.c execution test XFAIL: libgomp.c/target-32.c 7 blank line(s) in output XFAIL: libgomp.c/target-32.c (test for excess errors) UNRESOLVED: libgomp.c/target-32.c compilation failed to produce executable PASS: libgomp.c/thread-limit-2-nosleep.c (test for excess errors) PASS: libgomp.c/thread-limit-2-nosleep.c execution test XFAIL: libgomp.c/thread-limit-2.c 1 blank line(s) in output XFAIL: libgomp.c/thread-limit-2.c (test for excess errors) // No failures since GOMP_OFFLOAD_async_run works on amdgcn PASS: libgomp.c/target-33.c (test for excess errors) PASS: libgomp.c/target-33.c execution test PASS: libgomp.c/target-34.c (test for excess errors) PASS: libgomp.c/target-34.c execution test // No xfail here PASS: libgomp.c/target-link-1.c (test for excess errors) FAIL: libgomp.c/target-link-1.c execution test Note that target-link-1.c execution does also fail on amdgcn. Since - in contrast to nvptx - it seems that the cause of this failure has not yet been investigated and discussed, I have not added an xfail for amdgcn to this test. All testing has been done with a x86_64-linux-gnu host and target. Ok to commit this patch? Best regards, Frederik From 6e5e2d45f02235a0f72e6130dcd8d52f88f7b126 Mon Sep 17 00:00:00 2001 From: Frederik Harwath Date: Fri, 7 Feb 2020 08:03:00 +0100 Subject: [PATCH] xfail and improve some failing libgomp tests * libgomp.c/{target-32.c,thread-limit-2.c} Regarding failures because "no usleep implemented for nvptx." (cf. https://gcc.gnu.org/PR81690): Create test copies using busy wait instead of usleep, add xfails for nvptx and amdgcn (introduce new effective target for the latter) to original tests. * libgomp.c/target-{33,34}.c Regarding "no GOMP_OFFLOAD_async_run implemented in plugin-nvptx.c." (cf. https://gcc.gnu.org/PR81688): Add xfails for nvptx. * libgomp.c/target-link-1.c: Regarding "omp target link not implemented for nvptx." (cf. https://gcc.gnu.org/PR81689): Add xfail for nvptx. libgomp/ * testsuite/lib/libgomp.exp (proc match_effective_offload_target): New proc extracted from check_effective_target_offload_target_nvptx. (proc check_effective_target_offload_target_nvptx): Change to use match_effective_offload_target. (proc check_effective_target_offload_target_amgcn): New proc. * testsuite/libgomp.c/target-32-aux.h: New file, extracted from target-32.c. * testsuite/libgomp.c/target-32-nosleep.c: New test, like target-32.c but with busy waiting instead of usleep. * testsuite/libgomp.c/target-32.c: Use target-32-aux.h. * testsuite/libgomp.c/target-33.c Add xfail for execution on offload_target_nvptx. * testsuite/libgomp.c/target-34.c: Add xfail for execution on offload_target_nvptx. * testsuite/libgomp.c/target-link-1.c: Add xfail for offload_target_nvptx. * testsuite/libgomp.c/thread-limit-2-aux.h: New file, extracted from thread-limit-2.c. * testsuite/libgomp.c/thread-limit-2-nosleep.c: New test, like thread-limit-2.c, but with busy waiting instead of usleep. * testsuite/libgomp.c/thread-limit-2.c: Use thread-limit-2-aux.h. --- libgomp/testsuite/lib/libgomp.exp | 17 +++-- .../testsuite/libgomp.c/target-32-nosleep.c | 21 +++++++ libgomp/testsuite/libgomp.c/target-32.c | 61 ++++-------------- libgomp/testsuite/libgomp.c/target-33.c | 3 + libgomp/testsuite/libgomp.c/target-34.c | 3 + libgomp/testsuite/libgomp.c/target-link-1.c | 3 + .../libgomp.c/thread-limit-2-nosleep.c | 22 +++++++ libgomp/testsuite/libgomp.c/thread-limit-2.c | 63 ++++--------------- 8 files changed, 89 insertions(+), 104 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c/target-32-nosleep.c create mode 100644 libgomp/testsuite/libgomp.c/thread-limit-2-nosleep.c diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index e7ce784314d..3f4ced6fe7a 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -338,9 +338,8 @@ proc offload_target_to_openacc_device_type { offload_target } { } } } - -# Return 1 if compiling for offload target nvptx. -proc check_effective_target_offload_target_nvptx { } { +# Return 1 if compiling for an offload target matching the given target pattern. +proc match_effective_offload_target { target } { # Consider all actual options, including the flags passed to # 'gcc-dg-runtest', or 'gfortran-dg-runtest' (see the 'libgomp.*/*.exp' # files; in particular, '-foffload', 'libgomp.oacc-*/*.exp'), which don't @@ -353,13 +352,23 @@ proc check_effective_target_offload_target_nvptx { } { set gcc_output [libgomp_target_compile "" "" "none" $options] if [regexp "(?n)^OFFLOAD_TARGET_NAMES=(.*)" $gcc_output dummy offload_targets] { verbose "compiling for offload targets: $offload_targets" - return [string match "*:nvptx*:*" ":$offload_targets:"] + return [string match "*:$target:*" ":$offload_targets:"] } verbose "not compiling for any offload targets" return 0 } +# Return 1 if compiling for offload target nvptx. +proc check_effective_target_offload_target_nvptx { } { + return [match_effective_offload_target "nvptx*"] +} + +# Return 1 if compiling for offload target amdgcn +proc check_effective_target_offload_target_amdgcn { } { + return [match_effective_offload_target "amdgcn*"] +} + # Return 1 if offload device is available. proc check_effective_target_offload_device { } { return [check_runtime_nocache offload_device_available_ { diff --git a/libgomp/testsuite/libgomp.c/target-32-nosleep.c b/libgomp/testsuite/libgomp.c/target-32-nosleep.c new file mode 100644 index 00000000000..7534a2fdf44 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-32-nosleep.c @@ -0,0 +1,21 @@ +/* This is a variation of test-32.c for offloading targets which do not support + usleep. */ +/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } } + Cf. https://gcc.gnu.org/PR81688. */ + +#pragma omp declare target + +/* This function serves as a replacement for usleep in + this test case. It does not even attempt to be functionally + equivalent - we just want some sort of delay. */ + +void delay (int d) +{ + int i; + int N = d * 2000; + for (i = 0; i < N; i++) + asm volatile ("" : : : "memory"); +} +#pragma omp end declare target + +#include "target-32-aux.h" diff --git a/libgomp/testsuite/libgomp.c/target-32.c b/libgomp/testsuite/libgomp.c/target-32.c index 233877b702b..3366d06779d 100644 --- a/libgomp/testsuite/libgomp.c/target-32.c +++ b/libgomp/testsuite/libgomp.c/target-32.c @@ -1,54 +1,17 @@ -#include -#include - -int main () -{ - int a = 0, b = 0, c = 0, d[7]; - - #pragma omp parallel - #pragma omp single - { - #pragma omp task depend(out: d[0]) - a = 2; - - #pragma omp target enter data nowait map(to: a,b,c) depend(in: d[0]) depend(out: d[1]) - - #pragma omp target nowait map(alloc: a) depend(in: d[1]) depend(out: d[2]) - a++; +/* { dg-xfail-if "usleep not implemented" { offload_target_amdgcn || offload_target_nvptx } } + Cf. https://gcc.gnu.org/PR81690. */ +/* { dg-excess-errors "usleep not implemented" { xfail { offload_target_amdgcn || offload_target_nvptx } } } */ - #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[3]) - { - usleep (1000); - #pragma omp atomic update - b |= 4; - } - - #pragma omp target nowait map(alloc: b) depend(in: d[2]) depend(out: d[4]) - { - usleep (5000); - #pragma omp atomic update - b |= 1; - } - - #pragma omp target nowait map(alloc: c) depend(in: d[3], d[4]) depend(out: d[5]) - { - usleep (5000); - #pragma omp atomic update - c |= 8; - } +#include - #pragma omp target nowait map(alloc: c) depend(in: d[3], d[4]) depend(out: d[6]) - { - usleep (1000); - #pragma omp atomic update - c |= 2; - } +/* Use usleep for delays in the test case. + See also target-32-nosleep.c. */ - #pragma omp target exit data map(always,from: a,b,c) depend(in: d[5], d[6]) - } +void delay (int microseconds) +{ + usleep (microseconds); +} - if (a != 3 || b != 5 || c != 10) - abort (); +/* Include the actual test case definition. */ - return 0; -} +#include "target-32-aux.h" diff --git a/libgomp/testsuite/libgomp.c/target-33.c b/libgomp/testsuite/libgomp.c/target-33.c index 1bed4b6bc67..15d2d7e38ab 100644 --- a/libgomp/testsuite/libgomp.c/target-33.c +++ b/libgomp/testsuite/libgomp.c/target-33.c @@ -1,3 +1,6 @@ +/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } } + Cf. https://gcc.gnu.org/PR81688. */ + extern void abort (void); int diff --git a/libgomp/testsuite/libgomp.c/target-34.c b/libgomp/testsuite/libgomp.c/target-34.c index 66d9f54202b..5a3596424d8 100644 --- a/libgomp/testsuite/libgomp.c/target-34.c +++ b/libgomp/testsuite/libgomp.c/target-34.c @@ -1,3 +1,6 @@ +/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } } + Cf. https://gcc.gnu.org/PR81688. */ + extern void abort (void); int diff --git a/libgomp/testsuite/libgomp.c/target-link-1.c b/libgomp/testsuite/libgomp.c/target-link-1.c index 681677cc2aa..99ce33bc9b4 100644 --- a/libgomp/testsuite/libgomp.c/target-link-1.c +++ b/libgomp/testsuite/libgomp.c/target-link-1.c @@ -1,3 +1,6 @@ +/* { dg-xfail-if "#pragma omp target link not implemented" { offload_target_nvptx } } + Cf. https://gcc.gnu.org/PR81689. */ + struct S { int s, t; }; int a = 1, b = 1; diff --git a/libgomp/testsuite/libgomp.c/thread-limit-2-nosleep.c b/libgomp/testsuite/libgomp.c/thread-limit-2-nosleep.c new file mode 100644 index 00000000000..606db53d701 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/thread-limit-2-nosleep.c @@ -0,0 +1,22 @@ +/* This is a variation of thread-limit-2.c for offloading targets which do not support + usleep. */ + +#pragma omp declare target + +/* This function serves as a replacement for usleep in + this test case. It does not even attempt to be functionally + equivalent - we just want some sort of delay. */ + +void delay (int d) +{ + int i; + int N = d * 2000; + for (i = 0; i < N; i++) + asm volatile ("" : : : "memory"); +} +#pragma omp end declare target + + +/* Include the actual test case definition. */ + +#include "thread-limit-2-aux.h" diff --git a/libgomp/testsuite/libgomp.c/thread-limit-2.c b/libgomp/testsuite/libgomp.c/thread-limit-2.c index 1a97fb62985..c64781dbbb4 100644 --- a/libgomp/testsuite/libgomp.c/thread-limit-2.c +++ b/libgomp/testsuite/libgomp.c/thread-limit-2.c @@ -1,58 +1,19 @@ /* { dg-do run } */ /* { dg-set-target-env-var OMP_THREAD_LIMIT "9" } */ +/* { dg-xfail-if "usleep not implemented" { offload_target_amdgcn || offload_target_nvptx } } + Cf. https://gcc.gnu.org/PR81690. */ +/* { dg-excess-errors "usleep not implemented" { xfail { offload_target_amdgcn || offload_target_nvptx } } } */ -#include #include -#include -int -main () +/* Use usleep for delays in the test case. + See also thread-limit-2-nosleep.c. */ + +void delay (int microseconds) { - if (omp_get_thread_limit () != 9) - return 0; - omp_set_dynamic (0); - #pragma omp parallel num_threads (8) - if (omp_get_num_threads () != 8) - abort (); - #pragma omp parallel num_threads (16) - if (omp_get_num_threads () > 9) - abort (); - #pragma omp target if (0) - #pragma omp teams thread_limit (6) - { - if (omp_get_thread_limit () > 6) - abort (); - if (omp_get_thread_limit () == 6) - { - omp_set_dynamic (0); - omp_set_nested (1); - #pragma omp parallel num_threads (3) - if (omp_get_num_threads () != 3) - abort (); - #pragma omp parallel num_threads (3) - if (omp_get_num_threads () != 3) - abort (); - #pragma omp parallel num_threads (8) - if (omp_get_num_threads () > 6) - abort (); - #pragma omp parallel num_threads (6) - if (omp_get_num_threads () != 6) - abort (); - int cnt = 0; - #pragma omp parallel num_threads (5) - #pragma omp parallel num_threads (5) - #pragma omp parallel num_threads (2) - { - int v; - #pragma omp atomic capture - v = ++cnt; - if (v > 6) - abort (); - usleep (10000); - #pragma omp atomic - --cnt; - } - } - } - return 0; + usleep (microseconds); } + +/* Include the actual test case definition. */ + +#include "thread-limit-2-aux.h" -- 2.17.1