From patchwork Thu Feb 24 10:01:22 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1597103 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: bilbo.ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.a=rsa-sha256 header.s=default header.b=aB9Y4lUi; dkim-atps=neutral Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by bilbo.ozlabs.org (Postfix) with ESMTPS id 4K47jv6G2rz9sGH for ; Thu, 24 Feb 2022 21:02:10 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 97FCD3858020 for ; Thu, 24 Feb 2022 10:02:06 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 97FCD3858020 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1645696926; bh=qJ/VzMdApNueSyIOPZxojFIcCJ5PeOp+CTroMvA4XCc=; h=Date:Subject:To:References:In-Reply-To:List-Id:List-Unsubscribe: List-Archive:List-Post:List-Help:List-Subscribe:From:Reply-To:Cc: From; b=aB9Y4lUiRcrW581eofMs5Q/wgyVC5NiMEqQ5er+yTXYMZKRF/QgeAQ6DGqFRyGz7C BNJMrsbxYjAa13/BRfMT3Egp4YNWdW16u8eBAX9fELIu2/f6x7to83mQ/CAu6SZZvN rGLdM1yqxiFET2JRdGJKjgxi7O7oGB4g5XYXongY= X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from smtp-out1.suse.de (smtp-out1.suse.de [195.135.220.28]) by sourceware.org (Postfix) with ESMTPS id A2C563858402 for ; Thu, 24 Feb 2022 10:01:24 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org A2C563858402 Received: from imap2.suse-dmz.suse.de (imap2.suse-dmz.suse.de [192.168.254.74]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (P-521) server-digest SHA512) (No client certificate requested) by smtp-out1.suse.de (Postfix) with ESMTPS id 60FD6212BC; Thu, 24 Feb 2022 10:01:23 +0000 (UTC) Received: from imap2.suse-dmz.suse.de (imap2.suse-dmz.suse.de [192.168.254.74]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (P-521) server-digest SHA512) (No client certificate requested) by imap2.suse-dmz.suse.de (Postfix) with ESMTPS id 3D422139F3; Thu, 24 Feb 2022 10:01:23 +0000 (UTC) Received: from dovecot-director2.suse.de ([192.168.254.65]) by imap2.suse-dmz.suse.de with ESMTPSA id QSyvDXNXF2JAMAAAMHmgww (envelope-from ); Thu, 24 Feb 2022 10:01:23 +0000 Message-ID: Date: Thu, 24 Feb 2022 11:01:22 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:91.0) Gecko/20100101 Thunderbird/91.4.1 Subject: [PATCH][libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c Content-Language: en-US To: Jakub Jelinek References: <6b0f4fe5-06d4-5ede-5bd1-a53ed82f6d36@suse.de> <221e5215-3b8b-fa9f-4e7b-2faa86b8b991@codesourcery.com> <71e9c008-7ad1-c24f-76eb-6a03180525c4@suse.de> In-Reply-To: <71e9c008-7ad1-c24f-76eb-6a03180525c4@suse.de> X-Spam-Status: No, score=-12.6 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, GIT_PATCH_0, SPF_HELO_NONE, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-Patchwork-Original-From: Tom de Vries via Gcc-patches From: Tom de Vries Reply-To: Tom de Vries Cc: Tobias Burnus , gcc-patches Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" [ was: Re: [Patch] nvptx: Add -mptx=6.0 + -misa=sm_70 ] On 2/24/22 09:29, Tom de Vries wrote: > I'll try to submit a patch with one or more test-cases. Hi, These test-cases exercise the omp declare variant construct using the available nvptx isas. OK for trunk? Thanks, - Tom [libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c Add openmp test-cases that test the omp declare variant construct: ... #pragma omp declare variant (f30) match (device={isa("sm_30")}) ... using the available nvptx isas. On a Pascal board GT 1030 with sm_61, we have these unsupported: ... UNSUPPORTED: libgomp.c/declare-variant-3-sm70.c UNSUPPORTED: libgomp.c/declare-variant-3-sm75.c UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c ... and on a Turing board T400 with sm_75, we have this only this one: ... UNSUPPORTED: libgomp.c/declare-variant-3-sm80.c ... Tested on x86_64 with nvptx accelerator. libgomp/ChangeLog: 2022-02-24 Tom de Vries * testsuite/lib/libgomp.exp (check_effective_target_offload_device_nvptx_sm_xx) (check_effective_target_offload_device_nvptx_sm_30) (check_effective_target_offload_device_nvptx_sm_35) (check_effective_target_offload_device_nvptx_sm_53) (check_effective_target_offload_device_nvptx_sm_70) (check_effective_target_offload_device_nvptx_sm_75) (check_effective_target_offload_device_nvptx_sm_80): New proc. * testsuite/libgomp.c/declare-variant-3-sm30.c: New test. * testsuite/libgomp.c/declare-variant-3-sm35.c: New test. * testsuite/libgomp.c/declare-variant-3-sm53.c: New test. * testsuite/libgomp.c/declare-variant-3-sm70.c: New test. * testsuite/libgomp.c/declare-variant-3-sm75.c: New test. * testsuite/libgomp.c/declare-variant-3-sm80.c: New test. * testsuite/libgomp.c/declare-variant-3.h: New header file. --- libgomp/testsuite/lib/libgomp.exp | 46 +++++++++++++++ .../testsuite/libgomp.c/declare-variant-3-sm30.c | 5 ++ .../testsuite/libgomp.c/declare-variant-3-sm35.c | 5 ++ .../testsuite/libgomp.c/declare-variant-3-sm53.c | 5 ++ .../testsuite/libgomp.c/declare-variant-3-sm70.c | 5 ++ .../testsuite/libgomp.c/declare-variant-3-sm75.c | 5 ++ .../testsuite/libgomp.c/declare-variant-3-sm80.c | 5 ++ libgomp/testsuite/libgomp.c/declare-variant-3.h | 66 ++++++++++++++++++++++ 8 files changed, 142 insertions(+) diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index 8c5ecfff0ac..d664863b15c 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -426,6 +426,52 @@ proc check_effective_target_offload_device_nvptx { } { } ] } +# Return 1 if using nvptx offload device which supports -misa=sm_$SM. +proc check_effective_target_offload_device_nvptx_sm_xx { sm } { + if { ![check_effective_target_offload_device_nvptx] } { + return 0 + } + return [check_runtime_nocache offload_device_nvptx_sm_$sm { + int main () + { + int x = 1; + #pragma omp target map(tofrom: x) + x--; + return x; + } + } "-foffload=-misa=sm_$sm" ] +} + +# See check_effective_target_offload_device_nvptx_sm_xx. +proc check_effective_target_offload_device_nvptx_sm_30 { } { + return [check_effective_target_offload_device_nvptx_sm_xx 30] +} + +# See check_effective_target_offload_device_nvptx_sm_xx. +proc check_effective_target_offload_device_nvptx_sm_35 { } { + return [check_effective_target_offload_device_nvptx_sm_xx 35] +} + +# See check_effective_target_offload_device_nvptx_sm_xx. +proc check_effective_target_offload_device_nvptx_sm_53 { } { + return [check_effective_target_offload_device_nvptx_sm_xx 53] +} + +# See check_effective_target_offload_device_nvptx_sm_xx. +proc check_effective_target_offload_device_nvptx_sm_70 { } { + return [check_effective_target_offload_device_nvptx_sm_xx 70] +} + +# See check_effective_target_offload_device_nvptx_sm_xx. +proc check_effective_target_offload_device_nvptx_sm_75 { } { + return [check_effective_target_offload_device_nvptx_sm_xx 75] +} + +# See check_effective_target_offload_device_nvptx_sm_xx. +proc check_effective_target_offload_device_nvptx_sm_80 { } { + return [check_effective_target_offload_device_nvptx_sm_xx 80] +} + # Return 1 if at least one Nvidia GPU is accessible. proc check_effective_target_openacc_nvidia_accel_present { } { diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c new file mode 100644 index 00000000000..7c680b07a94 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm30.c @@ -0,0 +1,5 @@ +/* { dg-do run { target { offload_target_nvptx } } } */ +/* { dg-require-effective-target offload_device_nvptx_sm_30 } */ +/* { dg-additional-options "-foffload=-misa=sm_30" } */ + +#include "declare-variant-3.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c new file mode 100644 index 00000000000..b8b2a714248 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm35.c @@ -0,0 +1,5 @@ +/* { dg-do run { target { offload_target_nvptx } } } */ +/* { dg-require-effective-target offload_device_nvptx_sm_35 } */ +/* { dg-additional-options "-foffload=-misa=sm_35" } */ + +#include "declare-variant-3.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c new file mode 100644 index 00000000000..cfc7ee6f137 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm53.c @@ -0,0 +1,5 @@ +/* { dg-do run { target { offload_target_nvptx } } } */ +/* { dg-require-effective-target offload_device_nvptx_sm_53 } */ +/* { dg-additional-options "-foffload=-misa=sm_53" } */ + +#include "declare-variant-3.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c new file mode 100644 index 00000000000..4527cc1376c --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm70.c @@ -0,0 +1,5 @@ +/* { dg-do run { target { offload_target_nvptx } } } */ +/* { dg-require-effective-target offload_device_nvptx_sm_70 } */ +/* { dg-additional-options "-foffload=-misa=sm_70" } */ + +#include "declare-variant-3.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c new file mode 100644 index 00000000000..8e7da369fc5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm75.c @@ -0,0 +1,5 @@ +/* { dg-do run { target { offload_target_nvptx } } } */ +/* { dg-require-effective-target offload_device_nvptx_sm_75 } */ +/* { dg-additional-options "-foffload=-misa=sm_75" } */ + +#include "declare-variant-3.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c new file mode 100644 index 00000000000..5cf5b2867a2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-3-sm80.c @@ -0,0 +1,5 @@ +/* { dg-do run { target { offload_target_nvptx } } } */ +/* { dg-require-effective-target offload_device_nvptx_sm_80 } */ +/* { dg-additional-options "-foffload=-misa=sm_80" } */ + +#include "declare-variant-3.h" diff --git a/libgomp/testsuite/libgomp.c/declare-variant-3.h b/libgomp/testsuite/libgomp.c/declare-variant-3.h new file mode 100644 index 00000000000..772fc20a519 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/declare-variant-3.h @@ -0,0 +1,66 @@ +#pragma omp declare target +int +f30 (void) +{ + return 30; +} + +int +f35 (void) +{ + return 35; +} + +int +f53 (void) +{ + return 53; +} + +int +f70 (void) +{ + return 70; +} + +int +f75 (void) +{ + return 75; +} + +int +f80 (void) +{ + return 80; +} + +#pragma omp declare variant (f30) match (device={isa("sm_30")}) +#pragma omp declare variant (f35) match (device={isa("sm_35")}) +#pragma omp declare variant (f53) match (device={isa("sm_53")}) +#pragma omp declare variant (f70) match (device={isa("sm_70")}) +#pragma omp declare variant (f75) match (device={isa("sm_75")}) +#pragma omp declare variant (f80) match (device={isa("sm_80")}) +int +f (void) +{ + return 0; +} + +#pragma omp end declare target + +int +main (void) +{ + int v = 0; + + #pragma omp target map(from:v) + v = f (); + + if (v == 0) + __builtin_abort (); + + __builtin_printf ("Nvptx accelerator: sm_%d\n", v); + + return 0; +}