From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 19714 invoked by alias); 14 Oct 2014 16:12:31 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 19654 invoked by uid 89); 14 Oct 2014 16:12:30 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL,BAYES_00 autolearn=ham version=3.3.2 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; Tue, 14 Oct 2014 16:12:19 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1Xe4hr-00076a-MI from Julian_Brown@mentor.com ; Tue, 14 Oct 2014 09:12:16 -0700 Received: from octopus (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.181.6; Tue, 14 Oct 2014 17:12:13 +0100 Date: Tue, 14 Oct 2014 16:12:00 -0000 From: Julian Brown To: , Jakub Jelinek , "Thomas Schwinge" Subject: [gomp] [3/3] OpenACC 2.0 support for libgomp - documentation Message-ID: <20141014171207.2f06eb00@octopus> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="MP_/_NGz4nxS4FtWC45w2Ke2fv4" X-IsSubscribed: yes X-SW-Source: 2014-10/txt/msg01261.txt.bz2 --MP_/_NGz4nxS4FtWC45w2Ke2fv4 Content-Type: text/plain; charset="US-ASCII" Content-Transfer-Encoding: 7bit Content-Disposition: inline Content-length: 335 This is a version of the patch: https://gcc.gnu.org/ml/gcc-patches/2014-09/msg02024.html against gomp4 branch instead of mainline. OK to apply? Thanks, Julian xxxx-xx-xx Thomas Schwinge James Norris libgomp/ * libgomp.texi: Outline documentation for OpenACC. --MP_/_NGz4nxS4FtWC45w2Ke2fv4 Content-Type: text/x-patch Content-Transfer-Encoding: quoted-printable Content-Disposition: attachment; filename="0002-OpenACC-documentation.patch" Content-length: 23585 =46rom c58006a7ade2a9556bd73bac9ef45b3bbd62ca37 Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Wed, 17 Sep 2014 10:26:56 -0700 Subject: [PATCH 2/3] OpenACC documentation --- libgomp/libgomp.texi | 661 ++++++++++++++++++++++++++++++++++++++++++++++= ++-- 1 file changed, 636 insertions(+), 25 deletions(-) diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 254be57..9530a2b 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -31,10 +31,12 @@ texts being (a) (see below), and with the Back-Cover Te= xts being (b) @ifinfo @dircategory GNU Libraries @direntry -* libgomp: (libgomp). GNU OpenMP runtime library +* libgomp: (libgomp). GNU OpenACC and OpenMP runtime li= brary @end direntry =20 -This manual documents the GNU implementation of the OpenMP API for=20 +This manual documents the GNU implementation of the OpenACC API for=20 +offloading of code to accelerator devices in C/C++ and Fortran and +the GNU implementation of the OpenMP API for=20 multi-platform shared-memory parallel programming in C/C++ and Fortran. =20 Published by the Free Software Foundation @@ -48,7 +50,7 @@ Boston, MA 02110-1301 USA @setchapternewpage odd =20 @titlepage -@title The GNU OpenMP Implementation +@title The GNU OpenACC and OpenMP Implementation @page @vskip 0pt plus 1filll @comment For the @value{version-GCC} Version* @@ -69,7 +71,10 @@ Boston, MA 02110-1301, USA@* @top Introduction @cindex Introduction =20 -This manual documents the usage of libgomp, the GNU implementation of the= =20 +This manual documents the usage of libgomp, the GNU implementation of the +@uref{http://www.openacc.org/, OpenACC} Application Programming Interface = (API) +for offloading of code to accelerator devices in C/C++ and Fortran, and +the GNU implementation of the=20 @uref{http://www.openmp.org, OpenMP} Application Programming Interface (AP= I) for multi-platform shared-memory parallel programming in C/C++ and Fortran. =20 @@ -81,23 +86,619 @@ for multi-platform shared-memory parallel programming = in C/C++ and Fortran. @comment better formatting. @comment @menu -* Enabling OpenMP:: How to enable OpenMP for your applications. -* Runtime Library Routines:: The OpenMP runtime application programming= =20 - interface. -* Environment Variables:: Influencing runtime behavior with environme= nt=20 - variables. -* The libgomp ABI:: Notes on the external ABI presented by libg= omp. -* Reporting Bugs:: How to report bugs in GNU OpenMP. -* Copying:: GNU general public license says - how you can copy and share libgomp. -* GNU Free Documentation License:: - How you can copy and share this manual. -* Funding:: How to help assure continued work for free= =20 - software. -* Library Index:: Index of this documentation. +* Enabling OpenACC:: How to enable OpenACC for your + applications. +* OpenACC Runtime Library Routines:: The OpenACC runtime application + programming interface. +* OpenACC Environment Variables:: Influencing OpenACC runtime behavior = with + environment variables. +* OpenACC Library Interoperability:: OpenACC library interoperability with= the + NVIDIA CUBLAS library. +* Enabling OpenMP:: How to enable OpenMP for your + applications. +* OpenMP Runtime Library Routines: Runtime Library Routines. + The OpenMP runtime application progra= mming + interface. +* OpenMP Environment Variables: Environment Variables. + Influencing OpenMP runtime behavior w= ith + environment variables. +* The libgomp ABI:: Notes on the external libgomp ABI. +* Reporting Bugs:: How to report bugs. +* Copying:: GNU general public license says how y= ou + can copy and share libgomp. +* GNU Free Documentation License:: How you can copy and share this manua= l. +* Funding:: How to help assure continued work for= free + software. +* Library Index:: Index of this documentation. @end menu =20 =20 + +@c --------------------------------------------------------------------- +@c Enabling OpenACC +@c --------------------------------------------------------------------- + +@node Enabling OpenACC +@chapter Enabling OpenACC + +To activate the OpenACC extensions for C/C++ and Fortran, the compile-time= =20 +flag @command{-fopenacc} must be specified. This enables OpenACC, and +arranges for automatic linking of the OpenACC runtime library=20 +(@ref{Runtime Library Routines}). + +A complete description of all OpenACC directives accepted may be found in= =20 +the @uref{http://www.openacc.org/, OpenMP Application Programming +Interface} manual, version 2.0. + + +@c --------------------------------------------------------------------- +@c OpenACC Runtime Library Routines +@c --------------------------------------------------------------------- + +@node OpenACC Runtime Library Routines +@chapter OpenACC Runtime Library Routines + +The runtime routines described here are defined by section 3 of the OpenACC +specifications in version 2.0. +They have C linkage, and do not throw exceptions. +Generally, they are available only for the host, with the exception of +@code{acc_on_device}, which is available for both the host and the +acceleration device. + +@menu +* acc_get_num_devices:: Get number of devices for the given device= type +* acc_set_device_type:: +* acc_get_device_type:: +* acc_set_device_num:: +* acc_get_device_num:: +* acc_init:: +* acc_shutdown:: +* acc_on_device:: Whether executing on a particular device +* acc_malloc:: +* acc_free:: +* acc_copyin:: +* acc_present_or_copyin:: +* acc_create:: +* acc_present_or_create:: +* acc_copyout:: +* acc_delete:: +* acc_update_device:: +* acc_update_self:: +* acc_map_data:: +* acc_unmap_data:: +* acc_deviceptr:: +* acc_hostptr:: +* acc_is_present:: +* acc_memcpy_to_device:: +* acc_memcpy_from_device:: +@end menu + +API routines for target platforms. + +@menu +* acc_get_current_cuda_device:: +* acc_get_current_cuda_context:: +* acc_get_cuda_stream:: +* acc_set_cuda_stream:: +@end menu + + + +@node acc_get_num_devices +@section @code{acc_get_num_devices} -- Get number of devices for given dev= ice type +@table @asis +item @emph{Description} +This routine returns a value, between 0 and @emph{n}, indicating the +number of devices available for the given device type. It determines +the number of devices in a @emph{passive} manner. In other words, it +does not alter the state within the runtime environment aside from +possibly initializing an uninitialized device. This aspect allows +the routine to be called without concern for altering the interaction +with an attached accelerator device. + +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.1. +@end table + + + +@node acc_set_device_type +@section @code{acc_set_device_type} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.2. +@end table + + + +@node acc_get_device_type +@section @code{acc_get_device_type} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.3. +@end table + + + +@node acc_set_device_num +@section @code{acc_set_device_num} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.4. +@end table + + + +@node acc_get_device_num +@section @code{acc_get_device_num} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.5. +@end table + + + +@node acc_init +@section @code{acc_init} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.12. +@end table + + + +@node acc_shutdown +@section @code{acc_shutdown} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.13. +@end table + + + +@node acc_on_device +@section @code{acc_on_device} -- Whether executing on a particular device +@table @asis +@item @emph{Description}: +This routine tells the program whether it is executing on a particular +device. Based on the argument passed, GCC tries to evaluate this to a +constant at compile time, but library functions are also provided, for +both the host and the acceleration device. + +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.14. +@end table + + + +@node acc_malloc +@section @code{acc_malloc} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.15. +@end table + + + +@node acc_free +@section @code{acc_free} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.16. +@end table + + + +@node acc_copyin +@section @code{acc_copyin} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.17. +@end table + + + +@node acc_present_or_copyin +@section @code{acc_present_or_copyin} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.18. +@end table + + + +@node acc_create +@section @code{acc_create} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.19. +@end table + + + +@node acc_present_or_create +@section @code{acc_present_or_create} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.20. +@end table + + + +@node acc_copyout +@section @code{acc_copyout} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.21. +@end table + + + +@node acc_delete +@section @code{acc_delete} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.22. +@end table + + + +@node acc_update_device +@section @code{acc_update_device} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.23. +@end table + + + +@node acc_update_self +@section @code{acc_update_self} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.24. +@end table + + + +@node acc_map_data +@section @code{acc_map_data} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.25. +@end table + + + +@node acc_unmap_data +@section @code{acc_unmap_data} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.26. +@end table + + + +@node acc_deviceptr +@section @code{acc_deviceptr} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.27. +@end table + + + +@node acc_hostptr +@section @code{acc_hostptr} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.28. +@end table + + + +@node acc_is_present +@section @code{acc_is_present} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.29. +@end table + + + +@node acc_memcpy_to_device +@section @code{acc_memcpy_to_device} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.30. +@end table + + + +@node acc_memcpy_from_device +@section @code{acc_memcpy_from_device} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +3.2.31. +@end table + + + +@node acc_get_current_cuda_device +@section @code{acc_get_current_cuda_device} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +A.2.1.1. +@end table + + + +@node acc_get_current_cuda_context +@section @code{acc_get_current_cuda_context} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +A.2.1.2. +@end table + + + +@node acc_get_cuda_stream +@section @code{acc_get_cuda_stream} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +A.2.1.3. +@end table + + + +@node acc_set_cuda_stream +@section @code{acc_set_cuda_stream} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +A.2.1.4. +@end table + + + +@c --------------------------------------------------------------------- +@c OpenACC Environment Variables +@c --------------------------------------------------------------------- + +@node OpenACC Environment Variables +@chapter OpenACC Environment Variables + +The variables @env{ACC_DEVICE_TYPE} and @env{ACC_DEVICE_NUM} +are defined by section 4 of the OpenACC specification in version 2.0. +The variable @env{GCC_ACC_NOTIFY} is used for diagnostic purposes. + +@menu +* ACC_DEVICE_TYPE:: +* ACC_DEVICE_NUM:: +* GCC_ACC_NOTIFY:: +@end menu + + + +@node ACC_DEVICE_TYPE +@section @code{ACC_DEVICE_TYPE} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +4.1. +@end table + + + +@node ACC_DEVICE_NUM +@section @code{ACC_DEVICE_NUM} +@table @asis +@item @emph{Reference}: +@uref{http://www.openacc.org/, OpenACC specification v2.0}, section +4.2. +@end table + + + +@node GCC_ACC_NOTIFY +@section @code{GCC_ACC_NOTIFY} +@table @asis +@item @emph{Description}: +Print debug information pertaining to the accelerator. +@end table + + +@c --------------------------------------------------------------------- +@c OpenACC Library Interoperability +@c --------------------------------------------------------------------- + +@node OpenACC Library Interoperability +@chapter OpenACC Library Interoperability + +@section Introduction + +As the OpenACC library is built using the CUDA Driver API, the question has +arisen on what impact does using the OpenACC library have on a program that +uses the Runtime library, or a library based on the Runtime library, e.g., +CUBLAS@footnote{Seee section 2.26, "Interactions with the CUDA Driver API"= in +"CUDA Runtime API", Version 5.5, July 2013 and section 2.27, "VDPAU +Interoperability", in "CUDA Driver API", TRM-06703-001, Version 5.5, +July 2013, for additional information on library interoperability.}. +This chapter will describe the use cases and what changes are +required in order to use both the OpenACC library and the CUBLAS and Runti= me +libraries within a program. + +@section First invocation: NVIDIA CUBLAS library API + +In this first use case (see below), a function in the CUBLAS library is ca= lled +prior to any of the functions in the OpenACC library. More specifically, t= he +function @code{cublasCreate()}. + +When invoked, the function will initialize the library and allocate the +hardware resources on the host and the device on behalf of the caller. Once +the initialization and allocation has completed, a handle is returned to t= he +caller. The OpenACC library also requires initialization and allocation of +hardware resources. Since the CUBLAS library has already allocated the +hardware resources for the device, all that is left to do is to initialize +the OpenACC library and acquire the hardware resources on the host. + +Prior to calling the OpenACC function that will initialize the library and +allocate the host hardware resources, one needs to acquire the device numb= er +that was allocated during the call to @code{cublasCreate()}. The invoking = of the +runtime library function @code{cudaGetDevice()} will accomplish this. Once +acquired, the device number is passed along with the device type as +parameters to the OpenACC library function @code{acc_set_device_num()}. + +Once the call to @code{acc_set_device_num()} has completed, the OpenACC +library will be using the context that was created during the call to +@code{cublasCreate()}. In other words, both libraries will be sharing the +same context. + +@verbatim + /* Create the handle */ + s =3D cublasCreate(&h); + if (s !=3D CUBLAS_STATUS_SUCCESS) + { + fprintf(stderr, "cublasCreate failed %d\n", s); + exit(EXIT_FAILURE); + } + + /* Get the device number */ + e =3D cudaGetDevice(&dev); + if (e !=3D cudaSuccess) + { + fprintf(stderr, "cudaGetDevice failed %d\n", e); + exit(EXIT_FAILURE); + } + + /* Initialize OpenACC library and use device 'dev' */ + acc_set_device_num(dev, acc_device_nvidia); + +@end verbatim +@center Use Case 1=20 + +@section First invocation: OpenACC library API + +In this second use case (see below), a function in the OpenACC library is +called prior to any of the functions in the CUBLAS library. More specifici= ally, +the function acc_set_device_num(). + +In the use case presented here, the function @code{acc_set_device_num()} +is used to both initialize the OpenACC library and allocate the hardware +resources on the host and the device. In the call to the function, the +call parameters specify which device to use, i.e., 'dev', and what device +type to use, i.e., @code{acc_device_nvidia}. It should be noted that this +is but one method to initialize the OpenACC library and allocate the +appropriate hardware resources. Other methods are available through the +use of environment variables and these will be discussed in the next secti= on. + +Once the call to @code{acc_set_device_num()} has completed, other OpenACC +functions can be called as seen with multiple calls being made to +@code{acc_copyin()}. In addition, calls can be made to functions in the +CUBLAS library. In the use case a call to @code{cublasCreate()} is made +subsequent to the calls to @code{acc_copyin()}. +As seen in the previous use case, a call to @code{cublasCreate()} will +initialize the CUBLAS library and allocate the hardware resources on the +host and the device. However, since the device has already been allocated, +@code{cublasCreate()} will only initialize the CUBLAS library and allocate +the appropriate hardware resources on the host. The context that was creat= ed +as part of the OpenACC initialization will be shared with the CUBLAS libra= ry, +similarly to the first use case. + +@verbatim + dev =3D 0; + + acc_set_device_num(dev, acc_device_nvidia); + + /* Copy the first set to the device */ + d_X =3D acc_copyin(&h_X[0], N * sizeof (float)); + if (d_X =3D=3D NULL) + {=20 + fprintf(stderr, "copyin error h_X\n"); + exit(EXIT_FAILURE); + } + + /* Copy the second set to the device */ + d_Y =3D acc_copyin(&h_Y1[0], N * sizeof (float)); + if (d_Y =3D=3D NULL) + {=20 + fprintf(stderr, "copyin error h_Y1\n"); + exit(EXIT_FAILURE); + } + + /* Create the handle */ + s =3D cublasCreate(&h); + if (s !=3D CUBLAS_STATUS_SUCCESS) + { + fprintf(stderr, "cublasCreate failed %d\n", s); + exit(EXIT_FAILURE); + } + + /* Perform saxpy using CUBLAS library function */ + s =3D cublasSaxpy(h, N, &alpha, d_X, 1, d_Y, 1); + if (s !=3D CUBLAS_STATUS_SUCCESS) + { + fprintf(stderr, "cublasSaxpy failed %d\n", s); + exit(EXIT_FAILURE); + } + + /* Copy the results from the device */ + acc_memcpy_from_device(&h_Y1[0], d_Y, N * sizeof (float)); + +} +@end verbatim +@center Use Case 2 + +@section OpenACC library and environment variables + +There are two environment variables associated with the OpenACC library th= at +may be used to control the device type and device number. +Namely, @env{ACC_DEVICE_TYPE} and @env{ACC_DEVICE_NUM}. In the second +use case, the device type and device number were specified using +@code{acc_set_device_num()}. However, @env{ACC_DEVICE_TYPE} and=20 +@env{ACC_DEVICE_NUM} could have been defined and the call to +@code{acc_set_device_num()} would be not be required. At the time of the +call to @code{acc_copyin()}, these two environment variables would be +sampled and their values used. + +The use of the environment variables is only relevant when an OpenACC func= tion +is called prior to a call to @code{cudaCreate()}. If @code{cudaCreate()} +is called prior to a call to an OpenACC function, then a call to +@code{acc_set_device_num()}, must be done@footnote{More complete informati= on +about @env{ACC_DEVICE_TYPE} and @env{ACC_DEVICE_NUM} can be found in +sections 4.1 and 4.2 of the =E2=80=9CThe OpenACC +Application Programming Interface=E2=80=9D, Version 2.0, June, 2013.}. + + + @c --------------------------------------------------------------------- @c Enabling OpenMP @c --------------------------------------------------------------------- @@ -120,11 +721,11 @@ version 4.0. =20 =20 @c --------------------------------------------------------------------- -@c Runtime Library Routines +@c OpenMP Runtime Library Routines @c --------------------------------------------------------------------- =20 @node Runtime Library Routines -@chapter Runtime Library Routines +@chapter OpenMP Runtime Library Routines =20 The runtime routines described here are defined by Section 3 of the OpenMP specification in version 4.0. The routines are structured in following @@ -1281,11 +1882,11 @@ guaranteed not to change during the execution of th= e program. =20 =20 @c --------------------------------------------------------------------- -@c Environment Variables +@c OpenMP Environment Variables @c --------------------------------------------------------------------- =20 @node Environment Variables -@chapter Environment Variables +@chapter OpenMP Environment Variables =20 The environment variables which beginning with @env{OMP_} are defined by section 4 of the OpenMP specification in version 4.0, while those @@ -1701,6 +2302,7 @@ presented by libgomp. Only maintainers should need t= hem. * Implementing ORDERED construct:: * Implementing SECTIONS construct:: * Implementing SINGLE construct:: +* Implementing OpenACC's PARALLEL construct:: @end menu =20 =20 @@ -2065,15 +2667,24 @@ becomes =20 =20 =20 +@node Implementing OpenACC's PARALLEL construct +@section Implementing OpenACC's PARALLEL construct + +@smallexample + void GOACC_parallel () +@end smallexample + + + @c --------------------------------------------------------------------- -@c=20 +@c Reporting Bugs @c --------------------------------------------------------------------- =20 @node Reporting Bugs @chapter Reporting Bugs =20 -Bugs in the GNU OpenMP implementation should be reported via=20 -@uref{http://gcc.gnu.org/bugzilla/, Bugzilla}. For all cases, please add= =20 +Bugs in the GNU OpenACC or OpenMP implementation should be reported via +@uref{http://gcc.gnu.org/bugzilla/, Bugzilla}. For OpenMP cases, please a= dd "openmp" to the keywords field in the bug report. =20 =20 --=20 1.7.10.4 --MP_/_NGz4nxS4FtWC45w2Ke2fv4--