From patchwork Thu Sep 18 18:01:02 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 390877 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org 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 DDAC31400DD for ; Fri, 19 Sep 2014 04:01:23 +1000 (EST) 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:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=TEYM6vVfIdTkY+AY yF974T+sUCV9ZvyZ3J3H/JhA0HUY07KCDu5URNeMUDFrInh9MEMMHGVumIEBJv1N xAxKhPgGafmwNh7UvzRKY4EmBkkBxfHAF4HkDVWX3EGBkDsxP0Y60f6SUTleLN3C p9+2yX063o08Gmdp0l3Te3Cjfew= 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:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=eKexMrEaFsaZLNIysKPKgT Q6dkA=; b=BkkqaqFQ+IVzfy1ql4oJfYz6EHr1rFzVLCCebN3qMb6UXuhdZA1qiq bK9+gS0i83jbiWgE+2ooU+h8mNVSiJ4SRIfHJFse/LmTgG7J+GBCrgVLeraxuBXo abErrQzBzO6SONe1HcxQuW0QqDHGgR7y3ezhi2mYKSu2npSStF7zg= Received: (qmail 12078 invoked by alias); 18 Sep 2014 18:01:15 -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 12054 invoked by uid 89); 18 Sep 2014 18:01:14 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.2 required=5.0 tests=AWL, BAYES_00, KAM_STOCKGEN autolearn=no version=3.3.2 X-Spam-User: qpsmtpd, 2 recipients X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 18 Sep 2014 18:01:11 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1XUg0w-0002w4-JQ from Thomas_Schwinge@mentor.com ; Thu, 18 Sep 2014 11:01:07 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.3.181.6; Thu, 18 Sep 2014 19:01:04 +0100 From: Thomas Schwinge To: Jakub Jelinek , Tobias Burnus , Cesar Philippidis CC: , Subject: [gomp4] OpenACC acc_on_device (was: various OpenACC/PTX built-ins and a reduction tweak) In-Reply-To: <20140917084954.GB17454@tucnak.redhat.com> References: <5418D6B6.40801@codesourcery.com> <20140917084411.GA12930@physik.fu-berlin.de> <20140917084954.GB17454@tucnak.redhat.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.3.1 (x86_64-pc-linux-gnu) Date: Thu, 18 Sep 2014 20:01:02 +0200 Message-ID: <87k350n6zl.fsf@schwinge.name> MIME-Version: 1.0 Hi! Here is my OpenACC acc_on_device patch, in a more complete form, with test cases and all that. Thanks, Cesar, for getting the ball rolling! On Wed, 17 Sep 2014 10:49:54 +0200, Jakub Jelinek wrote: > On Wed, Sep 17, 2014 at 10:44:12AM +0200, Tobias Burnus wrote: > > Cesar Philippidis wrote: > > > The patch introduces the following OpenACC/PTX-specific built-ins: > > ... > > > > It is not completely clear how they are supposed to get used. Should the > > user call them directly in some cases? Or are they only used internally? > > > > acc_on_device sounds like a function which would be in C/C++ made available > > to the user via #define acc_on_device __builtin_acc_on_device. > > And not just providing acc_on_device prototype in some header? Yes, just a prototype. And next to DEF_GOACC_BUILTIN (configured the same as DEF_GOMP_BUILTIN), I add a new DEF_GOACC_BUILTIN_COMPILER that is configured to always provide the __builtin_[...] variant, but the un-prefixed [...] only if -fopenacc is in effect. Does that look alright? > Without > looking at the OpenACC standard, it sounds like this function could be > similar to omp_is_initial_device, so can and should be handled supposedly > similarly. I think we've been talking about this at the Cauldron, where you agreed that omp_is_initial_device should also be implemented as a builtin. (Or am I confusing things?) > > However, the rest looks as if it should rather be an internal function > > instead of a builtin. Or should the user really ever call the builtin > > directly? > > GOMP_* functions are builtins and not internal functions too, all those > functions are library functions, while the user typically doesn't call them > directly, they still are implemented in the library. Internal functions are > used for something that doesn't have a library implementation and is not > something user can call directly. > > Regarding Fortran: Builtins aren't directly available to the user. You have to > > wrap them into an intrinsic to make them available. If they have to be made > > available via a module (e.g. via "module acc) - you have to create a virtual > > module, which provides the intrinsic. If you don't want to convert the whole > > module, you could create an auxiliar module (e.g. acc_internal_) which provides > > only those bits - and then include it ("use,intrinsic :: ...") it in the > > main module - written in normal Fortran. This I have not yet addressed -- please see the TODO comments in the gcc/fortran/ files as well as Fortran test cases. > For the user callable fortran functions, for OpenMP libgomp just provides > *_ entrypoints to * functions. Perhaps acc_on_device_ could be provided > too. This is what I had done already. Does that patch look good? (With the Fortran things still to be addressed.) (And, obviously this is not yet based on the Tobias/Jim Fortran module/header rewrite.) commit 8efbd08ed058d7ed3c43e10fbff0eac35b4defc9 Author: Thomas Schwinge Date: Fri Jul 4 11:45:05 2014 +0000 OpenACC acc_on_device. gcc/ * builtins.def (DEF_GOACC_BUILTIN_COMPILER): New macro. * oacc-builtins.def (BUILT_IN_GOACC_UPDATE): New builtin. * builtins.c (expand_builtin_acc_on_device): New function. (expand_builtin): Use it to handle BUILT_IN_ACC_ON_DEVICE. (is_inexpensive_builtin): Handle BUILT_IN_ACC_ON_DEVICE. gcc/fortran/ * f95-lang.c (DEF_GOACC_BUILTIN_COMPILER): New macro. * types.def (BT_FN_INT_INT): New type. gcc/testsuite/ * c-c++-common/goacc/acc_on_device-1.c: New file. * c-c++-common/goacc/acc_on_device-2.c: Likewise. * c-c++-common/goacc/acc_on_device-2-off.c: Likewise. * gfortran.dg/goacc/acc_on_device-1.f95: Likewise. * gfortran.dg/goacc/acc_on_device-2.f95: Likewise. * gfortran.dg/goacc/acc_on_device-2-off.f95: Likewise. libgomp/ * libgomp.map (OACC_2.0): Add acc_on_device, acc_on_device_. * fortran.c: Include "openacc.h". (acc_on_device_): New function. * oacc-parallel.c: Include "openacc.h". (acc_on_device): New function. * openacc.f90 (acc_device_kind, acc_device_none) (acc_device_default, acc_device_host, acc_device_not_host): New parameters. (acc_on_device): New function declaration. * openacc_lib.h (acc_device_kind, acc_device_none) (acc_device_default, acc_device_host, acc_device_not_host): New parameters. (acc_on_device): New function declaration. * openacc.h (acc_device_t): New enum. (acc_on_device): New function declaration. * testsuite/libgomp.oacc-c/acc_on_device-1.c: New file. * testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise. * testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise. --- gcc/ChangeLog.gomp | 8 ++++ gcc/builtins.c | 50 ++++++++++++++++++++ gcc/builtins.def | 8 +++- gcc/fortran/ChangeLog.gomp | 5 ++ gcc/fortran/f95-lang.c | 5 ++ gcc/fortran/types.def | 1 + gcc/oacc-builtins.def | 2 + gcc/testsuite/ChangeLog.gomp | 9 ++++ gcc/testsuite/c-c++-common/goacc/acc_on_device-1.c | 20 ++++++++ .../c-c++-common/goacc/acc_on_device-2-off.c | 17 +++++++ gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c | 17 +++++++ .../gfortran.dg/goacc/acc_on_device-1.f95 | 22 +++++++++ .../gfortran.dg/goacc/acc_on_device-2-off.f95 | 39 ++++++++++++++++ .../gfortran.dg/goacc/acc_on_device-2.f95 | 40 ++++++++++++++++ libgomp/ChangeLog.gomp | 22 +++++++++ libgomp/fortran.c | 8 ++++ libgomp/libgomp.map | 3 ++ libgomp/oacc-parallel.c | 10 ++++ libgomp/openacc.f90 | 17 ++++++- libgomp/openacc.h | 13 +++++- libgomp/openacc_lib.h | 16 ++++++- libgomp/testsuite/libgomp.oacc-c/acc_on_device-1.c | 54 ++++++++++++++++++++++ .../libgomp.oacc-fortran/acc_on_device-1-1.f90 | 39 ++++++++++++++++ .../libgomp.oacc-fortran/acc_on_device-1-2.f | 39 ++++++++++++++++ .../libgomp.oacc-fortran/acc_on_device-1-3.f | 39 ++++++++++++++++ 25 files changed, 498 insertions(+), 5 deletions(-) diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index ef9a81d..0c25a27 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-09-18 Thomas Schwinge + + * builtins.def (DEF_GOACC_BUILTIN_COMPILER): New macro. + * oacc-builtins.def (BUILT_IN_GOACC_UPDATE): New builtin. + * builtins.c (expand_builtin_acc_on_device): New function. + (expand_builtin): Use it to handle BUILT_IN_ACC_ON_DEVICE. + (is_inexpensive_builtin): Handle BUILT_IN_ACC_ON_DEVICE. + 2014-09-08 Thomas Schwinge * configure.ac (offload_targets): Remove. diff --git gcc/builtins.c gcc/builtins.c index 975f696..5b2ebcc 100644 --- gcc/builtins.c +++ gcc/builtins.c @@ -5747,6 +5747,49 @@ expand_stack_save (void) return ret; } + +/* Expand OpenACC acc_on_device. + + This has to happen late (that is, not in early folding; expand_builtin_*, + rather than fold_builtin_*), as we have to act differently for host and + acceleration device (ACCEL_COMPILER conditional). */ + +static rtx +expand_builtin_acc_on_device (tree exp, rtx target ATTRIBUTE_UNUSED) +{ + if (!validate_arglist (exp, INTEGER_TYPE, VOID_TYPE)) + return NULL_RTX; + + tree arg, v1, v2, ret; + location_t loc; + + arg = CALL_EXPR_ARG (exp, 0); + arg = builtin_save_expr (arg); + loc = EXPR_LOCATION (exp); + + /* Build: (arg == v1 || arg == v2) ? 1 : 0. */ + +#ifdef ACCEL_COMPILER + v1 = build_int_cst (TREE_TYPE (arg), /* TODO: acc_device_not_host */ 3); + v2 = build_int_cst (TREE_TYPE (arg), ACCEL_COMPILER_acc_device); +#else + v1 = build_int_cst (TREE_TYPE (arg), /* TODO: acc_device_none */ 0); + v2 = build_int_cst (TREE_TYPE (arg), /* TODO: acc_device_host */ 2); +#endif + + v1 = fold_build2_loc (loc, EQ_EXPR, integer_type_node, arg, v1); + v2 = fold_build2_loc (loc, EQ_EXPR, integer_type_node, arg, v2); + + /* Can't use TRUTH_ORIF_EXPR, as that is not supported by + expand_expr_real*. */ + ret = fold_build3_loc (loc, COND_EXPR, integer_type_node, v1, v1, v2); + ret = fold_build3_loc (loc, COND_EXPR, integer_type_node, + ret, integer_one_node, integer_zero_node); + + return expand_normal (ret); +} + + /* Expand an expression EXP that calls a built-in function, with result going to TARGET if that's convenient (and in mode MODE if that's convenient). @@ -6816,6 +6859,12 @@ expand_builtin (tree exp, rtx target, rtx subtarget, enum machine_mode mode, expand_builtin_cilk_pop_frame (exp); return const0_rtx; + case BUILT_IN_ACC_ON_DEVICE: + target = expand_builtin_acc_on_device (exp, target); + if (target) + return target; + break; + default: /* just do library call, if unknown builtin */ break; } @@ -12748,6 +12797,7 @@ is_inexpensive_builtin (tree decl) case BUILT_IN_LABS: case BUILT_IN_LLABS: case BUILT_IN_PREFETCH: + case BUILT_IN_ACC_ON_DEVICE: return true; default: diff --git gcc/builtins.def gcc/builtins.def index 2ef896e..b9b8e74 100644 --- gcc/builtins.def +++ gcc/builtins.def @@ -146,12 +146,16 @@ along with GCC; see the file COPYING3. If not see DEF_BUILTIN (ENUM, NAME, BUILT_IN_NORMAL, BT_LAST, BT_LAST, false, false, \ false, ATTR_LAST, false, false) -/* Builtin used by the implementation of GNU OpenACC. None of these are - actually implemented in the compiler; they're all in libgomp. */ +/* Builtin used by the implementation of GNU OpenACC. Few of these are + actually implemented in the compiler; most are in libgomp. */ #undef DEF_GOACC_BUILTIN #define DEF_GOACC_BUILTIN(ENUM, NAME, TYPE, ATTRS) \ DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE, \ false, true, true, ATTRS, false, flag_openacc) +#undef DEF_GOACC_BUILTIN_COMPILER +#define DEF_GOACC_BUILTIN_COMPILER(ENUM, NAME, TYPE, ATTRS) \ + DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE, \ + flag_openacc, true, true, ATTRS, false, true) /* Builtin used by the implementation of GNU OpenMP. None of these are actually implemented in the compiler; they're all in libgomp. */ diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp index cc7c888..df86db7 100644 --- gcc/fortran/ChangeLog.gomp +++ gcc/fortran/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-09-18 Thomas Schwinge + + * f95-lang.c (DEF_GOACC_BUILTIN_COMPILER): New macro. + * types.def (BT_FN_INT_INT): New type. + 2014-09-08 Cesar Philippidis * gfortran.h (enum OMP_LIST_FIRST, OMP_LIST_LAST): New diff --git gcc/fortran/f95-lang.c gcc/fortran/f95-lang.c index e7c64b7..1b017b1 100644 --- gcc/fortran/f95-lang.c +++ gcc/fortran/f95-lang.c @@ -1093,7 +1093,12 @@ gfc_init_builtin_functions (void) #define DEF_GOACC_BUILTIN(code, name, type, attr) \ gfc_define_builtin ("__builtin_" name, builtin_types[type], \ code, name, attr); +#undef DEF_GOACC_BUILTIN_COMPILER + /* TODO: this is not doing the right thing. */ +#define DEF_GOACC_BUILTIN_COMPILER(code, name, type, attr) \ + gfc_define_builtin (name, builtin_types[type], code, name, attr); #include "../oacc-builtins.def" +#undef DEF_GOACC_BUILTIN_COMPILER #undef DEF_GOACC_BUILTIN } diff --git gcc/fortran/types.def gcc/fortran/types.def index 59ac4c3..6c2fdc0 100644 --- gcc/fortran/types.def +++ gcc/fortran/types.def @@ -82,6 +82,7 @@ DEF_FUNCTION_TYPE_0 (BT_FN_VOID, BT_VOID) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTR, BT_VOID, BT_PTR) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR) +DEF_FUNCTION_TYPE_1 (BT_FN_INT_INT, BT_INT, BT_INT) DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT) DEF_FUNCTION_TYPE_1 (BT_FN_PTR_PTR, BT_PTR, BT_PTR) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_INT, BT_VOID, BT_INT) diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def index dfb688c..e4bc756 100644 --- gcc/oacc-builtins.def +++ gcc/oacc-builtins.def @@ -39,3 +39,5 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel", ATTR_NOTHROW_LIST) DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update", BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) +DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device", + BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST) diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 4427521..e210c6b 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,12 @@ +2014-09-18 Thomas Schwinge + + * c-c++-common/goacc/acc_on_device-1.c: New file. + * c-c++-common/goacc/acc_on_device-2.c: Likewise. + * c-c++-common/goacc/acc_on_device-2-off.c: Likewise. + * gfortran.dg/goacc/acc_on_device-1.f95: Likewise. + * gfortran.dg/goacc/acc_on_device-2.f95: Likewise. + * gfortran.dg/goacc/acc_on_device-2-off.f95: Likewise. + 2014-09-08 Cesar Philippidis * gfortran.dg/goacc/private-1.f95: New test. diff --git gcc/testsuite/c-c++-common/goacc/acc_on_device-1.c gcc/testsuite/c-c++-common/goacc/acc_on_device-1.c new file mode 100644 index 0000000..d0e137b --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/acc_on_device-1.c @@ -0,0 +1,20 @@ +/* Have to enable optimizations, as otherwise builtins won't be expanded. */ +/* { dg-additional-options "-O -fdump-rtl-expand" } */ + +int +f (void) +{ + int r = 0; + + r |= acc_on_device (); + r |= acc_on_device (1, 2); + r |= acc_on_device (3.14); + r |= acc_on_device ("hello"); + + return r; +} + +/* Unsuitable to be handled as a builtin, so we're expecting four calls. + { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 4 "expand" } } */ + +/* { dg-final { cleanup-rtl-dump "expand" } } */ diff --git gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c new file mode 100644 index 0000000..ddc43ab --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c @@ -0,0 +1,17 @@ +/* Have to enable optimizations, as otherwise builtins won't be expanded. */ +/* { dg-additional-options "-O -fdump-rtl-expand -fno-openacc" } */ + +typedef enum acc_device_t { acc_device_X = 123 } acc_device_t; +extern int acc_on_device (acc_device_t); + +int +f (void) +{ + const int dev = acc_device_X; + return acc_on_device (dev); +} + +/* Without -fopenacc, we're expecting one call. + { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 1 "expand" } } */ + +/* { dg-final { cleanup-rtl-dump "expand" } } */ diff --git gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c new file mode 100644 index 0000000..65b4ae6 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c @@ -0,0 +1,17 @@ +/* Have to enable optimizations, as otherwise builtins won't be expanded. */ +/* { dg-additional-options "-O -fdump-rtl-expand" } */ + +typedef enum acc_device_t { acc_device_X = 123 } acc_device_t; +extern int acc_on_device (acc_device_t); + +int +f (void) +{ + const int dev = acc_device_X; + return acc_on_device (dev); +} + +/* With -fopenacc, we're expecting the builtin to be expanded, so no calls. + { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 0 "expand" } } */ + +/* { dg-final { cleanup-rtl-dump "expand" } } */ diff --git gcc/testsuite/gfortran.dg/goacc/acc_on_device-1.f95 gcc/testsuite/gfortran.dg/goacc/acc_on_device-1.f95 new file mode 100644 index 0000000..9dfde26 --- /dev/null +++ gcc/testsuite/gfortran.dg/goacc/acc_on_device-1.f95 @@ -0,0 +1,22 @@ +! Have to enable optimizations, as otherwise builtins won't be expanded. +! { dg-additional-options "-O -fdump-rtl-expand" } + +logical function f () + implicit none + + external acc_on_device + logical (4) acc_on_device + + f = .false. + f = f .or. acc_on_device () + f = f .or. acc_on_device (1, 2) + f = f .or. acc_on_device (3.14) + f = f .or. acc_on_device ("hello") + + return +end function f + +! Unsuitable to be handled as a builtin, so we're expecting four calls. +! { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 4 "expand" } } + +! { dg-final { cleanup-rtl-dump "expand" } } diff --git gcc/testsuite/gfortran.dg/goacc/acc_on_device-2-off.f95 gcc/testsuite/gfortran.dg/goacc/acc_on_device-2-off.f95 new file mode 100644 index 0000000..cf28264 --- /dev/null +++ gcc/testsuite/gfortran.dg/goacc/acc_on_device-2-off.f95 @@ -0,0 +1,39 @@ +! Have to enable optimizations, as otherwise builtins won't be expanded. +! { dg-additional-options "-O -fdump-rtl-expand -fno-openacc" } + +module openacc_kinds + implicit none + + integer, parameter :: acc_device_kind = 4 + +end module openacc_kinds + +module openacc + use openacc_kinds + implicit none + + integer (acc_device_kind), parameter :: acc_device_host = 2 + + interface + function acc_on_device (dev) + use openacc_kinds + logical (4) :: acc_on_device + integer (acc_device_kind), intent (in) :: dev + end function acc_on_device + end interface +end module openacc + +logical (4) function f () + use openacc + implicit none + + integer (4), parameter :: dev = 2 + + f = acc_on_device (dev) + return +end function f + +! Without -fopenacc, we're expecting one call. +! { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 1 "expand" } } + +! { dg-final { cleanup-rtl-dump "expand" } } diff --git gcc/testsuite/gfortran.dg/goacc/acc_on_device-2.f95 gcc/testsuite/gfortran.dg/goacc/acc_on_device-2.f95 new file mode 100644 index 0000000..7730a60 --- /dev/null +++ gcc/testsuite/gfortran.dg/goacc/acc_on_device-2.f95 @@ -0,0 +1,40 @@ +! Have to enable optimizations, as otherwise builtins won't be expanded. +! { dg-additional-options "-O -fdump-rtl-expand" } + +module openacc_kinds + implicit none + + integer, parameter :: acc_device_kind = 4 + +end module openacc_kinds + +module openacc + use openacc_kinds + implicit none + + integer (acc_device_kind), parameter :: acc_device_host = 2 + + interface + function acc_on_device (dev) + use openacc_kinds + logical (4) :: acc_on_device + integer (acc_device_kind), intent (in) :: dev + end function acc_on_device + end interface +end module openacc + +logical (4) function f () + use openacc + implicit none + + integer (4), parameter :: dev = 2 + + f = acc_on_device (dev) + return +end function f + +! With -fopenacc, we're expecting the builtin to be expanded, so no calls. +! TODO: not working. +! { dg-final { scan-rtl-dump-times "\\\(call \[^\\n\]*\\\"acc_on_device" 0 "expand" { xfail *-*-* } } } + +! { dg-final { cleanup-rtl-dump "expand" } } diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 5b2a39d..8d774ee 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,3 +1,25 @@ +2014-09-18 Thomas Schwinge + + * libgomp.map (OACC_2.0): Add acc_on_device, acc_on_device_. + * fortran.c: Include "openacc.h". + (acc_on_device_): New function. + * oacc-parallel.c: Include "openacc.h". + (acc_on_device): New function. + * openacc.f90 (acc_device_kind, acc_device_none) + (acc_device_default, acc_device_host, acc_device_not_host): New + parameters. + (acc_on_device): New function declaration. + * openacc_lib.h (acc_device_kind, acc_device_none) + (acc_device_default, acc_device_host, acc_device_not_host): New + parameters. + (acc_on_device): New function declaration. + * openacc.h (acc_device_t): New enum. + (acc_on_device): New function declaration. + * testsuite/libgomp.oacc-c/acc_on_device-1.c: New file. + * testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise. + * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise. + * testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise. + 2014-07-09 Thomas Schwinge Jakub Jelinek diff --git libgomp/fortran.c libgomp/fortran.c index 1f30c51..28c83cc 100644 --- libgomp/fortran.c +++ libgomp/fortran.c @@ -26,6 +26,7 @@ #include "libgomp.h" #include "libgomp_f.h" +#include "openacc.h" #include #include @@ -73,6 +74,7 @@ ialias_redirect (omp_get_num_devices) ialias_redirect (omp_get_num_teams) ialias_redirect (omp_get_team_num) ialias_redirect (omp_is_initial_device) +ialias_redirect (acc_on_device) #endif #ifndef LIBGOMP_GNU_SYMBOL_VERSIONING @@ -492,3 +494,9 @@ omp_is_initial_device_ (void) { return omp_is_initial_device (); } + +int32_t +acc_on_device_ (const int32_t *dev) +{ + return acc_on_device (*dev); +} diff --git libgomp/libgomp.map libgomp/libgomp.map index c575be3..69a4d83 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -234,6 +234,9 @@ GOMP_4.0.1 { } GOMP_4.0; OACC_2.0 { + global: + acc_on_device; + acc_on_device_; }; GOACC_2.0 { diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c index 79b6254..02fbb12 100644 --- libgomp/oacc-parallel.c +++ libgomp/oacc-parallel.c @@ -27,6 +27,7 @@ #include "libgomp.h" #include "libgomp_g.h" +#include "openacc.h" void GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target, @@ -128,3 +129,12 @@ GOACC_update (int device, const void *openmp_target, size_t mapnum, } GOMP_target_update (device, openmp_target, mapnum, hostaddrs, sizes, kinds_); } + +/* TODO: Move elsewhere. */ +int +acc_on_device (acc_device_t dev) +{ + /* Just rely on the compiler builtin. */ + return __builtin_acc_on_device (dev); +} +ialias (acc_on_device) diff --git libgomp/openacc.f90 libgomp/openacc.f90 index b2a79f6..70b58d6 100644 --- libgomp/openacc.f90 +++ libgomp/openacc.f90 @@ -1,6 +1,6 @@ ! OpenACC Runtime Library Definitions. -! Copyright (C) 2013 Free Software Foundation, Inc. +! Copyright (C) 2013-2014 Free Software Foundation, Inc. ! Contributed by Thomas Schwinge . @@ -28,6 +28,8 @@ module openacc_kinds implicit none + integer, parameter :: acc_device_kind = 4 + end module openacc_kinds module openacc @@ -36,4 +38,17 @@ module openacc integer, parameter :: openacc_version = 201306 + integer (acc_device_kind), parameter :: acc_device_none = 0 + integer (acc_device_kind), parameter :: acc_device_default = 1 + integer (acc_device_kind), parameter :: acc_device_host = 2 + integer (acc_device_kind), parameter :: acc_device_not_host = 3 + + interface + function acc_on_device (dev) + use openacc_kinds + logical (4) :: acc_on_device + integer (acc_device_kind), intent (in) :: dev + end function acc_on_device + end interface + end module openacc diff --git libgomp/openacc.h libgomp/openacc.h index a6f7ec94..cde7429 100644 --- libgomp/openacc.h +++ libgomp/openacc.h @@ -37,7 +37,18 @@ extern "C" { #else # define __GOACC_NOTHROW __attribute__ ((__nothrow__)) #endif - + +typedef enum acc_device_t + { + acc_device_none = 0, + acc_device_default, /* This has to be a distinct value, as no + return value can match it. */ + acc_device_host = 2, + acc_device_not_host = 3 + } acc_device_t; + +int acc_on_device (acc_device_t __dev) __GOACC_NOTHROW; + #ifdef __cplusplus } #endif diff --git libgomp/openacc_lib.h libgomp/openacc_lib.h index d19c95c..be49100 100644 --- libgomp/openacc_lib.h +++ libgomp/openacc_lib.h @@ -1,6 +1,6 @@ ! OpenACC Runtime Library Definitions. -*- mode: fortran -*- -! Copyright (C) 2013 Free Software Foundation, Inc. +! Copyright (C) 2013-2014 Free Software Foundation, Inc. ! Contributed by Thomas Schwinge . @@ -27,3 +27,17 @@ integer openacc_version parameter (openacc_version = 201306) + + integer acc_device_kind + parameter (acc_device_kind = 4) + integer (acc_device_kind) acc_device_none + parameter (acc_device_none = 0) + integer (acc_device_kind) acc_device_default + parameter (acc_device_default = 1) + integer (acc_device_kind) acc_device_host + parameter (acc_device_host = 2) + integer (acc_device_kind) acc_device_not_host + parameter (acc_device_not_host = 3) + + external acc_on_device + logical (4) acc_on_device diff --git libgomp/testsuite/libgomp.oacc-c/acc_on_device-1.c libgomp/testsuite/libgomp.oacc-c/acc_on_device-1.c new file mode 100644 index 0000000..f216587 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/acc_on_device-1.c @@ -0,0 +1,54 @@ +/* Disable the acc_on_device builtin; we want to test the libgomp library + function. */ +/* TODO: Remove -DACC_DEVICE_TYPE_host once that is set by the test harness. */ +/* { dg-additional-options "-fno-builtin-acc_on_device -DACC_DEVICE_TYPE_host" } */ + +#include +#include + +int +main (int argc, char *argv[]) +{ + /* Host. */ + + { + if (!acc_on_device (acc_device_none)) + abort (); + if (!acc_on_device (acc_device_host)) + abort (); + if (acc_on_device (acc_device_not_host)) + abort (); + } + + + /* Host via offloading fallback mode. */ + +#pragma acc parallel if(0) + { + if (!acc_on_device (acc_device_none)) + abort (); + if (!acc_on_device (acc_device_host)) + abort (); + if (acc_on_device (acc_device_not_host)) + abort (); + } + + +#if !ACC_DEVICE_TYPE_host + + /* Offloaded. */ + +#pragma acc parallel + { + if (acc_on_device (acc_device_none)) + abort (); + if (acc_on_device (acc_device_host)) + abort (); + if (!acc_on_device (acc_device_not_host)) + abort (); + } + +#endif + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 new file mode 100644 index 0000000..c4597a6 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 @@ -0,0 +1,39 @@ +! TODO: Remove -DACC_DEVICE_TYPE_host once that is set by the test harness. +! { dg-additional-options "-cpp -DACC_DEVICE_TYPE_host" } +! TODO: Have to disable the acc_on_device builtin for we want to test the +! libgomp library function? The command line option +! '-fno-builtin-acc_on_device' is valid for C/C++/ObjC/ObjC++ but not for +! Fortran. + +use openacc +implicit none + +! Host. + +if (.not. acc_on_device (acc_device_none)) call abort +if (.not. acc_on_device (acc_device_host)) call abort +if (acc_on_device (acc_device_not_host)) call abort + + +! Host via offloading fallback mode. + +!$acc parallel if(.false.) +if (.not. acc_on_device (acc_device_none)) call abort +if (.not. acc_on_device (acc_device_host)) call abort +if (acc_on_device (acc_device_not_host)) call abort +!$acc end parallel + + +#if !ACC_DEVICE_TYPE_host + +! Offloaded. + +!$acc parallel +if (acc_on_device (acc_device_none)) call abort +if (acc_on_device (acc_device_host)) call abort +if (.not. acc_on_device (acc_device_not_host)) call abort +!$acc end parallel + +#endif + +end diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f new file mode 100644 index 0000000..3787e1e --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f @@ -0,0 +1,39 @@ +! TODO: Remove -DACC_DEVICE_TYPE_host once that is set by the test harness. +! { dg-additional-options "-cpp -DACC_DEVICE_TYPE_host" } +! TODO: Have to disable the acc_on_device builtin for we want to test +! the libgomp library function? The command line option +! '-fno-builtin-acc_on_device' is valid for C/C++/ObjC/ObjC++ but not +! for Fortran. + + USE OPENACC + IMPLICIT NONE + +!Host. + + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT + IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT + + +!Host via offloading fallback mode. + +!$ACC PARALLEL IF(.FALSE.) + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT + IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT +!$ACC END PARALLEL + + +#if !ACC_DEVICE_TYPE_host + +! Offloaded. + +!$ACC PARALLEL + IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT + IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT +!$ACC END PARALLEL + +#endif + + END diff --git libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f new file mode 100644 index 0000000..1ee5926 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f @@ -0,0 +1,39 @@ +! TODO: Remove -DACC_DEVICE_TYPE_host once that is set by the test harness. +! { dg-additional-options "-cpp -DACC_DEVICE_TYPE_host" } +! TODO: Have to disable the acc_on_device builtin for we want to test +! the libgomp library function? The command line option +! '-fno-builtin-acc_on_device' is valid for C/C++/ObjC/ObjC++ but not +! for Fortran. + + IMPLICIT NONE + INCLUDE "openacc_lib.h" + +!Host. + + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT + IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT + + +!Host via offloading fallback mode. + +!$ACC PARALLEL IF(.FALSE.) + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT + IF (ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT +!$ACC END PARALLEL + + +#if !ACC_DEVICE_TYPE_host + +! Offloaded. + +!$ACC PARALLEL + IF (ACC_ON_DEVICE (ACC_DEVICE_NONE)) CALL ABORT + IF (ACC_ON_DEVICE (ACC_DEVICE_HOST)) CALL ABORT + IF (.NOT. ACC_ON_DEVICE (ACC_DEVICE_NOT_HOST)) CALL ABORT +!$ACC END PARALLEL + +#endif + + END