fix gcc-plugin conflict

gcc-plugin/.dir-locals.el

@@ -0,0 +1,3 @@
+;; Hey Emacs, this is much nicer down here.
+
+((c-mode . ((c-file-style . "gnu"))))

gcc-plugin/COPYING

@@ -0,0 +1,674 @@
+                    GNU GENERAL PUBLIC LICENSE
+                       Version 3, 29 June 2007
+
+ Copyright (C) 2007 Free Software Foundation, Inc. <http://fsf.org/>
+ Everyone is permitted to copy and distribute verbatim copies
+ of this license document, but changing it is not allowed.
+
+                            Preamble
+
+  The GNU General Public License is a free, copyleft license for
+software and other kinds of works.
+
+  The licenses for most software and other practical works are designed
+to take away your freedom to share and change the works.  By contrast,
+the GNU General Public License is intended to guarantee your freedom to
+share and change all versions of a program--to make sure it remains free
+software for all its users.  We, the Free Software Foundation, use the
+GNU General Public License for most of our software; it applies also to
+any other work released this way by its authors.  You can apply it to
+your programs, too.
+
+  When we speak of free software, we are referring to freedom, not
+price.  Our General Public Licenses are designed to make sure that you
+have the freedom to distribute copies of free software (and charge for
+them if you wish), that you receive source code or can get it if you
+want it, that you can change the software or use pieces of it in new
+free programs, and that you know you can do these things.
+
+  To protect your rights, we need to prevent others from denying you
+these rights or asking you to surrender the rights.  Therefore, you have
+certain responsibilities if you distribute copies of the software, or if
+you modify it: responsibilities to respect the freedom of others.
+
+  For example, if you distribute copies of such a program, whether
+gratis or for a fee, you must pass on to the recipients the same
+freedoms that you received.  You must make sure that they, too, receive
+or can get the source code.  And you must show them these terms so they
+know their rights.
+
+  Developers that use the GNU GPL protect your rights with two steps:
+(1) assert copyright on the software, and (2) offer you this License
+giving you legal permission to copy, distribute and/or modify it.
+
+  For the developers' and authors' protection, the GPL clearly explains
+that there is no warranty for this free software.  For both users' and
+authors' sake, the GPL requires that modified versions be marked as
+changed, so that their problems will not be attributed erroneously to
+authors of previous versions.
+
+  Some devices are designed to deny users access to install or run
+modified versions of the software inside them, although the manufacturer
+can do so.  This is fundamentally incompatible with the aim of
+protecting users' freedom to change the software.  The systematic
+pattern of such abuse occurs in the area of products for individuals to
+use, which is precisely where it is most unacceptable.  Therefore, we
+have designed this version of the GPL to prohibit the practice for those
+products.  If such problems arise substantially in other domains, we
+stand ready to extend this provision to those domains in future versions
+of the GPL, as needed to protect the freedom of users.
+
+  Finally, every program is threatened constantly by software patents.
+States should not allow patents to restrict development and use of
+software on general-purpose computers, but in those that do, we wish to
+avoid the special danger that patents applied to a free program could
+make it effectively proprietary.  To prevent this, the GPL assures that
+patents cannot be used to render the program non-free.
+
+  The precise terms and conditions for copying, distribution and
+modification follow.
+
+                       TERMS AND CONDITIONS
+
+  0. Definitions.
+
+  "This License" refers to version 3 of the GNU General Public License.
+
+  "Copyright" also means copyright-like laws that apply to other kinds of
+works, such as semiconductor masks.
+
+  "The Program" refers to any copyrightable work licensed under this
+License.  Each licensee is addressed as "you".  "Licensees" and
+"recipients" may be individuals or organizations.
+
+  To "modify" a work means to copy from or adapt all or part of the work
+in a fashion requiring copyright permission, other than the making of an
+exact copy.  The resulting work is called a "modified version" of the
+earlier work or a work "based on" the earlier work.
+
+  A "covered work" means either the unmodified Program or a work based
+on the Program.
+
+  To "propagate" a work means to do anything with it that, without
+permission, would make you directly or secondarily liable for
+infringement under applicable copyright law, except executing it on a
+computer or modifying a private copy.  Propagation includes copying,
+distribution (with or without modification), making available to the
+public, and in some countries other activities as well.
+
+  To "convey" a work means any kind of propagation that enables other
+parties to make or receive copies.  Mere interaction with a user through
+a computer network, with no transfer of a copy, is not conveying.
+
+  An interactive user interface displays "Appropriate Legal Notices"
+to the extent that it includes a convenient and prominently visible
+feature that (1) displays an appropriate copyright notice, and (2)
+tells the user that there is no warranty for the work (except to the
+extent that warranties are provided), that licensees may convey the
+work under this License, and how to view a copy of this License.  If
+the interface presents a list of user commands or options, such as a
+menu, a prominent item in the list meets this criterion.
+
+  1. Source Code.
+
+  The "source code" for a work means the preferred form of the work
+for making modifications to it.  "Object code" means any non-source
+form of a work.
+
+  A "Standard Interface" means an interface that either is an official
+standard defined by a recognized standards body, or, in the case of
+interfaces specified for a particular programming language, one that
+is widely used among developers working in that language.
+
+  The "System Libraries" of an executable work include anything, other
+than the work as a whole, that (a) is included in the normal form of
+packaging a Major Component, but which is not part of that Major
+Component, and (b) serves only to enable use of the work with that
+Major Component, or to implement a Standard Interface for which an
+implementation is available to the public in source code form.  A
+"Major Component", in this context, means a major essential component
+(kernel, window system, and so on) of the specific operating system
+(if any) on which the executable work runs, or a compiler used to
+produce the work, or an object code interpreter used to run it.
+
+  The "Corresponding Source" for a work in object code form means all
+the source code needed to generate, install, and (for an executable
+work) run the object code and to modify the work, including scripts to
+control those activities.  However, it does not include the work's
+System Libraries, or general-purpose tools or generally available free
+programs which are used unmodified in performing those activities but
+which are not part of the work.  For example, Corresponding Source
+includes interface definition files associated with source files for
+the work, and the source code for shared libraries and dynamically
+linked subprograms that the work is specifically designed to require,
+such as by intimate data communication or control flow between those
+subprograms and other parts of the work.
+
+  The Corresponding Source need not include anything that users
+can regenerate automatically from other parts of the Corresponding
+Source.
+
+  The Corresponding Source for a work in source code form is that
+same work.
+
+  2. Basic Permissions.
+
+  All rights granted under this License are granted for the term of
+copyright on the Program, and are irrevocable provided the stated
+conditions are met.  This License explicitly affirms your unlimited
+permission to run the unmodified Program.  The output from running a
+covered work is covered by this License only if the output, given its
+content, constitutes a covered work.  This License acknowledges your
+rights of fair use or other equivalent, as provided by copyright law.
+
+  You may make, run and propagate covered works that you do not
+convey, without conditions so long as your license otherwise remains
+in force.  You may convey covered works to others for the sole purpose
+of having them make modifications exclusively for you, or provide you
+with facilities for running those works, provided that you comply with
+the terms of this License in conveying all material for which you do
+not control copyright.  Those thus making or running the covered works
+for you must do so exclusively on your behalf, under your direction
+and control, on terms that prohibit them from making any copies of
+your copyrighted material outside their relationship with you.
+
+  Conveying under any other circumstances is permitted solely under
+the conditions stated below.  Sublicensing is not allowed; section 10
+makes it unnecessary.
+
+  3. Protecting Users' Legal Rights From Anti-Circumvention Law.
+
+  No covered work shall be deemed part of an effective technological
+measure under any applicable law fulfilling obligations under article
+11 of the WIPO copyright treaty adopted on 20 December 1996, or
+similar laws prohibiting or restricting circumvention of such
+measures.
+
+  When you convey a covered work, you waive any legal power to forbid
+circumvention of technological measures to the extent such circumvention
+is effected by exercising rights under this License with respect to
+the covered work, and you disclaim any intention to limit operation or
+modification of the work as a means of enforcing, against the work's
+users, your or third parties' legal rights to forbid circumvention of
+technological measures.
+
+  4. Conveying Verbatim Copies.
+
+  You may convey verbatim copies of the Program's source code as you
+receive it, in any medium, provided that you conspicuously and
+appropriately publish on each copy an appropriate copyright notice;
+keep intact all notices stating that this License and any
+non-permissive terms added in accord with section 7 apply to the code;
+keep intact all notices of the absence of any warranty; and give all
+recipients a copy of this License along with the Program.
+
+  You may charge any price or no price for each copy that you convey,
+and you may offer support or warranty protection for a fee.
+
+  5. Conveying Modified Source Versions.
+
+  You may convey a work based on the Program, or the modifications to
+produce it from the Program, in the form of source code under the
+terms of section 4, provided that you also meet all of these conditions:
+
+    a) The work must carry prominent notices stating that you modified
+    it, and giving a relevant date.
+
+    b) The work must carry prominent notices stating that it is
+    released under this License and any conditions added under section
+    7.  This requirement modifies the requirement in section 4 to
+    "keep intact all notices".
+
+    c) You must license the entire work, as a whole, under this
+    License to anyone who comes into possession of a copy.  This
+    License will therefore apply, along with any applicable section 7
+    additional terms, to the whole of the work, and all its parts,
+    regardless of how they are packaged.  This License gives no
+    permission to license the work in any other way, but it does not
+    invalidate such permission if you have separately received it.
+
+    d) If the work has interactive user interfaces, each must display
+    Appropriate Legal Notices; however, if the Program has interactive
+    interfaces that do not display Appropriate Legal Notices, your
+    work need not make them do so.
+
+  A compilation of a covered work with other separate and independent
+works, which are not by their nature extensions of the covered work,
+and which are not combined with it such as to form a larger program,
+in or on a volume of a storage or distribution medium, is called an
+"aggregate" if the compilation and its resulting copyright are not
+used to limit the access or legal rights of the compilation's users
+beyond what the individual works permit.  Inclusion of a covered work
+in an aggregate does not cause this License to apply to the other
+parts of the aggregate.
+
+  6. Conveying Non-Source Forms.
+
+  You may convey a covered work in object code form under the terms
+of sections 4 and 5, provided that you also convey the
+machine-readable Corresponding Source under the terms of this License,
+in one of these ways:
+
+    a) Convey the object code in, or embodied in, a physical product
+    (including a physical distribution medium), accompanied by the
+    Corresponding Source fixed on a durable physical medium
+    customarily used for software interchange.
+
+    b) Convey the object code in, or embodied in, a physical product
+    (including a physical distribution medium), accompanied by a
+    written offer, valid for at least three years and valid for as
+    long as you offer spare parts or customer support for that product
+    model, to give anyone who possesses the object code either (1) a
+    copy of the Corresponding Source for all the software in the
+    product that is covered by this License, on a durable physical
+    medium customarily used for software interchange, for a price no
+    more than your reasonable cost of physically performing this
+    conveying of source, or (2) access to copy the
+    Corresponding Source from a network server at no charge.
+
+    c) Convey individual copies of the object code with a copy of the
+    written offer to provide the Corresponding Source.  This
+    alternative is allowed only occasionally and noncommercially, and
+    only if you received the object code with such an offer, in accord
+    with subsection 6b.
+
+    d) Convey the object code by offering access from a designated
+    place (gratis or for a charge), and offer equivalent access to the
+    Corresponding Source in the same way through the same place at no
+    further charge.  You need not require recipients to copy the
+    Corresponding Source along with the object code.  If the place to
+    copy the object code is a network server, the Corresponding Source
+    may be on a different server (operated by you or a third party)
+    that supports equivalent copying facilities, provided you maintain
+    clear directions next to the object code saying where to find the
+    Corresponding Source.  Regardless of what server hosts the
+    Corresponding Source, you remain obligated to ensure that it is
+    available for as long as needed to satisfy these requirements.
+
+    e) Convey the object code using peer-to-peer transmission, provided
+    you inform other peers where the object code and Corresponding
+    Source of the work are being offered to the general public at no
+    charge under subsection 6d.
+
+  A separable portion of the object code, whose source code is excluded
+from the Corresponding Source as a System Library, need not be
+included in conveying the object code work.
+
+  A "User Product" is either (1) a "consumer product", which means any
+tangible personal property which is normally used for personal, family,
+or household purposes, or (2) anything designed or sold for incorporation
+into a dwelling.  In determining whether a product is a consumer product,
+doubtful cases shall be resolved in favor of coverage.  For a particular
+product received by a particular user, "normally used" refers to a
+typical or common use of that class of product, regardless of the status
+of the particular user or of the way in which the particular user
+actually uses, or expects or is expected to use, the product.  A product
+is a consumer product regardless of whether the product has substantial
+commercial, industrial or non-consumer uses, unless such uses represent
+the only significant mode of use of the product.
+
+  "Installation Information" for a User Product means any methods,
+procedures, authorization keys, or other information required to install
+and execute modified versions of a covered work in that User Product from
+a modified version of its Corresponding Source.  The information must
+suffice to ensure that the continued functioning of the modified object
+code is in no case prevented or interfered with solely because
+modification has been made.
+
+  If you convey an object code work under this section in, or with, or
+specifically for use in, a User Product, and the conveying occurs as
+part of a transaction in which the right of possession and use of the
+User Product is transferred to the recipient in perpetuity or for a
+fixed term (regardless of how the transaction is characterized), the
+Corresponding Source conveyed under this section must be accompanied
+by the Installation Information.  But this requirement does not apply
+if neither you nor any third party retains the ability to install
+modified object code on the User Product (for example, the work has
+been installed in ROM).
+
+  The requirement to provide Installation Information does not include a
+requirement to continue to provide support service, warranty, or updates
+for a work that has been modified or installed by the recipient, or for
+the User Product in which it has been modified or installed.  Access to a
+network may be denied when the modification itself materially and
+adversely affects the operation of the network or violates the rules and
+protocols for communication across the network.
+
+  Corresponding Source conveyed, and Installation Information provided,
+in accord with this section must be in a format that is publicly
+documented (and with an implementation available to the public in
+source code form), and must require no special password or key for
+unpacking, reading or copying.
+
+  7. Additional Terms.
+
+  "Additional permissions" are terms that supplement the terms of this
+License by making exceptions from one or more of its conditions.
+Additional permissions that are applicable to the entire Program shall
+be treated as though they were included in this License, to the extent
+that they are valid under applicable law.  If additional permissions
+apply only to part of the Program, that part may be used separately
+under those permissions, but the entire Program remains governed by
+this License without regard to the additional permissions.
+
+  When you convey a copy of a covered work, you may at your option
+remove any additional permissions from that copy, or from any part of
+it.  (Additional permissions may be written to require their own
+removal in certain cases when you modify the work.)  You may place
+additional permissions on material, added by you to a covered work,
+for which you have or can give appropriate copyright permission.
+
+  Notwithstanding any other provision of this License, for material you
+add to a covered work, you may (if authorized by the copyright holders of
+that material) supplement the terms of this License with terms:
+
+    a) Disclaiming warranty or limiting liability differently from the
+    terms of sections 15 and 16 of this License; or
+
+    b) Requiring preservation of specified reasonable legal notices or
+    author attributions in that material or in the Appropriate Legal
+    Notices displayed by works containing it; or
+
+    c) Prohibiting misrepresentation of the origin of that material, or
+    requiring that modified versions of such material be marked in
+    reasonable ways as different from the original version; or
+
+    d) Limiting the use for publicity purposes of names of licensors or
+    authors of the material; or
+
+    e) Declining to grant rights under trademark law for use of some
+    trade names, trademarks, or service marks; or
+
+    f) Requiring indemnification of licensors and authors of that
+    material by anyone who conveys the material (or modified versions of
+    it) with contractual assumptions of liability to the recipient, for
+    any liability that these contractual assumptions directly impose on
+    those licensors and authors.
+
+  All other non-permissive additional terms are considered "further
+restrictions" within the meaning of section 10.  If the Program as you
+received it, or any part of it, contains a notice stating that it is
+governed by this License along with a term that is a further
+restriction, you may remove that term.  If a license document contains
+a further restriction but permits relicensing or conveying under this
+License, you may add to a covered work material governed by the terms
+of that license document, provided that the further restriction does
+not survive such relicensing or conveying.
+
+  If you add terms to a covered work in accord with this section, you
+must place, in the relevant source files, a statement of the
+additional terms that apply to those files, or a notice indicating
+where to find the applicable terms.
+
+  Additional terms, permissive or non-permissive, may be stated in the
+form of a separately written license, or stated as exceptions;
+the above requirements apply either way.
+
+  8. Termination.
+
+  You may not propagate or modify a covered work except as expressly
+provided under this License.  Any attempt otherwise to propagate or
+modify it is void, and will automatically terminate your rights under
+this License (including any patent licenses granted under the third
+paragraph of section 11).
+
+  However, if you cease all violation of this License, then your
+license from a particular copyright holder is reinstated (a)
+provisionally, unless and until the copyright holder explicitly and
+finally terminates your license, and (b) permanently, if the copyright
+holder fails to notify you of the violation by some reasonable means
+prior to 60 days after the cessation.
+
+  Moreover, your license from a particular copyright holder is
+reinstated permanently if the copyright holder notifies you of the
+violation by some reasonable means, this is the first time you have
+received notice of violation of this License (for any work) from that
+copyright holder, and you cure the violation prior to 30 days after
+your receipt of the notice.
+
+  Termination of your rights under this section does not terminate the
+licenses of parties who have received copies or rights from you under
+this License.  If your rights have been terminated and not permanently
+reinstated, you do not qualify to receive new licenses for the same
+material under section 10.
+
+  9. Acceptance Not Required for Having Copies.
+
+  You are not required to accept this License in order to receive or
+run a copy of the Program.  Ancillary propagation of a covered work
+occurring solely as a consequence of using peer-to-peer transmission
+to receive a copy likewise does not require acceptance.  However,
+nothing other than this License grants you permission to propagate or
+modify any covered work.  These actions infringe copyright if you do
+not accept this License.  Therefore, by modifying or propagating a
+covered work, you indicate your acceptance of this License to do so.
+
+  10. Automatic Licensing of Downstream Recipients.
+
+  Each time you convey a covered work, the recipient automatically
+receives a license from the original licensors, to run, modify and
+propagate that work, subject to this License.  You are not responsible
+for enforcing compliance by third parties with this License.
+
+  An "entity transaction" is a transaction transferring control of an
+organization, or substantially all assets of one, or subdividing an
+organization, or merging organizations.  If propagation of a covered
+work results from an entity transaction, each party to that
+transaction who receives a copy of the work also receives whatever
+licenses to the work the party's predecessor in interest had or could
+give under the previous paragraph, plus a right to possession of the
+Corresponding Source of the work from the predecessor in interest, if
+the predecessor has it or can get it with reasonable efforts.
+
+  You may not impose any further restrictions on the exercise of the
+rights granted or affirmed under this License.  For example, you may
+not impose a license fee, royalty, or other charge for exercise of
+rights granted under this License, and you may not initiate litigation
+(including a cross-claim or counterclaim in a lawsuit) alleging that
+any patent claim is infringed by making, using, selling, offering for
+sale, or importing the Program or any portion of it.
+
+  11. Patents.
+
+  A "contributor" is a copyright holder who authorizes use under this
+License of the Program or a work on which the Program is based.  The
+work thus licensed is called the contributor's "contributor version".
+
+  A contributor's "essential patent claims" are all patent claims
+owned or controlled by the contributor, whether already acquired or
+hereafter acquired, that would be infringed by some manner, permitted
+by this License, of making, using, or selling its contributor version,
+but do not include claims that would be infringed only as a
+consequence of further modification of the contributor version.  For
+purposes of this definition, "control" includes the right to grant
+patent sublicenses in a manner consistent with the requirements of
+this License.
+
+  Each contributor grants you a non-exclusive, worldwide, royalty-free
+patent license under the contributor's essential patent claims, to
+make, use, sell, offer for sale, import and otherwise run, modify and
+propagate the contents of its contributor version.
+
+  In the following three paragraphs, a "patent license" is any express
+agreement or commitment, however denominated, not to enforce a patent
+(such as an express permission to practice a patent or covenant not to
+sue for patent infringement).  To "grant" such a patent license to a
+party means to make such an agreement or commitment not to enforce a
+patent against the party.
+
+  If you convey a covered work, knowingly relying on a patent license,
+and the Corresponding Source of the work is not available for anyone
+to copy, free of charge and under the terms of this License, through a
+publicly available network server or other readily accessible means,
+then you must either (1) cause the Corresponding Source to be so
+available, or (2) arrange to deprive yourself of the benefit of the
+patent license for this particular work, or (3) arrange, in a manner
+consistent with the requirements of this License, to extend the patent
+license to downstream recipients.  "Knowingly relying" means you have
+actual knowledge that, but for the patent license, your conveying the
+covered work in a country, or your recipient's use of the covered work
+in a country, would infringe one or more identifiable patents in that
+country that you have reason to believe are valid.
+
+  If, pursuant to or in connection with a single transaction or
+arrangement, you convey, or propagate by procuring conveyance of, a
+covered work, and grant a patent license to some of the parties
+receiving the covered work authorizing them to use, propagate, modify
+or convey a specific copy of the covered work, then the patent license
+you grant is automatically extended to all recipients of the covered
+work and works based on it.
+
+  A patent license is "discriminatory" if it does not include within
+the scope of its coverage, prohibits the exercise of, or is
+conditioned on the non-exercise of one or more of the rights that are
+specifically granted under this License.  You may not convey a covered
+work if you are a party to an arrangement with a third party that is
+in the business of distributing software, under which you make payment
+to the third party based on the extent of your activity of conveying
+the work, and under which the third party grants, to any of the
+parties who would receive the covered work from you, a discriminatory
+patent license (a) in connection with copies of the covered work
+conveyed by you (or copies made from those copies), or (b) primarily
+for and in connection with specific products or compilations that
+contain the covered work, unless you entered into that arrangement,
+or that patent license was granted, prior to 28 March 2007.
+
+  Nothing in this License shall be construed as excluding or limiting
+any implied license or other defenses to infringement that may
+otherwise be available to you under applicable patent law.
+
+  12. No Surrender of Others' Freedom.
+
+  If conditions are imposed on you (whether by court order, agreement or
+otherwise) that contradict the conditions of this License, they do not
+excuse you from the conditions of this License.  If you cannot convey a
+covered work so as to satisfy simultaneously your obligations under this
+License and any other pertinent obligations, then as a consequence you may
+not convey it at all.  For example, if you agree to terms that obligate you
+to collect a royalty for further conveying from those to whom you convey
+the Program, the only way you could satisfy both those terms and this
+License would be to refrain entirely from conveying the Program.
+
+  13. Use with the GNU Affero General Public License.
+
+  Notwithstanding any other provision of this License, you have
+permission to link or combine any covered work with a work licensed
+under version 3 of the GNU Affero General Public License into a single
+combined work, and to convey the resulting work.  The terms of this
+License will continue to apply to the part which is the covered work,
+but the special requirements of the GNU Affero General Public License,
+section 13, concerning interaction through a network will apply to the
+combination as such.
+
+  14. Revised Versions of this License.
+
+  The Free Software Foundation may publish revised and/or new versions of
+the GNU General Public License from time to time.  Such new versions will
+be similar in spirit to the present version, but may differ in detail to
+address new problems or concerns.
+
+  Each version is given a distinguishing version number.  If the
+Program specifies that a certain numbered version of the GNU General
+Public License "or any later version" applies to it, you have the
+option of following the terms and conditions either of that numbered
+version or of any later version published by the Free Software
+Foundation.  If the Program does not specify a version number of the
+GNU General Public License, you may choose any version ever published
+by the Free Software Foundation.
+
+  If the Program specifies that a proxy can decide which future
+versions of the GNU General Public License can be used, that proxy's
+public statement of acceptance of a version permanently authorizes you
+to choose that version for the Program.
+
+  Later license versions may give you additional or different
+permissions.  However, no additional obligations are imposed on any
+author or copyright holder as a result of your choosing to follow a
+later version.
+
+  15. Disclaimer of Warranty.
+
+  THERE IS NO WARRANTY FOR THE PROGRAM, TO THE EXTENT PERMITTED BY
+APPLICABLE LAW.  EXCEPT WHEN OTHERWISE STATED IN WRITING THE COPYRIGHT
+HOLDERS AND/OR OTHER PARTIES PROVIDE THE PROGRAM "AS IS" WITHOUT WARRANTY
+OF ANY KIND, EITHER EXPRESSED OR IMPLIED, INCLUDING, BUT NOT LIMITED TO,
+THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+PURPOSE.  THE ENTIRE RISK AS TO THE QUALITY AND PERFORMANCE OF THE PROGRAM
+IS WITH YOU.  SHOULD THE PROGRAM PROVE DEFECTIVE, YOU ASSUME THE COST OF
+ALL NECESSARY SERVICING, REPAIR OR CORRECTION.
+
+  16. Limitation of Liability.
+
+  IN NO EVENT UNLESS REQUIRED BY APPLICABLE LAW OR AGREED TO IN WRITING
+WILL ANY COPYRIGHT HOLDER, OR ANY OTHER PARTY WHO MODIFIES AND/OR CONVEYS
+THE PROGRAM AS PERMITTED ABOVE, BE LIABLE TO YOU FOR DAMAGES, INCLUDING ANY
+GENERAL, SPECIAL, INCIDENTAL OR CONSEQUENTIAL DAMAGES ARISING OUT OF THE
+USE OR INABILITY TO USE THE PROGRAM (INCLUDING BUT NOT LIMITED TO LOSS OF
+DATA OR DATA BEING RENDERED INACCURATE OR LOSSES SUSTAINED BY YOU OR THIRD
+PARTIES OR A FAILURE OF THE PROGRAM TO OPERATE WITH ANY OTHER PROGRAMS),
+EVEN IF SUCH HOLDER OR OTHER PARTY HAS BEEN ADVISED OF THE POSSIBILITY OF
+SUCH DAMAGES.
+
+  17. Interpretation of Sections 15 and 16.
+
+  If the disclaimer of warranty and limitation of liability provided
+above cannot be given local legal effect according to their terms,
+reviewing courts shall apply local law that most closely approximates
+an absolute waiver of all civil liability in connection with the
+Program, unless a warranty or assumption of liability accompanies a
+copy of the Program in return for a fee.
+
+                     END OF TERMS AND CONDITIONS
+
+            How to Apply These Terms to Your New Programs
+
+  If you develop a new program, and you want it to be of the greatest
+possible use to the public, the best way to achieve this is to make it
+free software which everyone can redistribute and change under these terms.
+
+  To do so, attach the following notices to the program.  It is safest
+to attach them to the start of each source file to most effectively
+state the exclusion of warranty; and each file should have at least
+the "copyright" line and a pointer to where the full notice is found.
+
+    <one line to give the program's name and a brief idea of what it does.>
+    Copyright (C) <year>  <name of author>
+
+    This program is free software: you can redistribute it and/or modify
+    it under the terms of the GNU General Public License as published by
+    the Free Software Foundation, either version 3 of the License, or
+    (at your option) any later version.
+
+    This program is distributed in the hope that it will be useful,
+    but WITHOUT ANY WARRANTY; without even the implied warranty of
+    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+    GNU General Public License for more details.
+
+    You should have received a copy of the GNU General Public License
+    along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+Also add information on how to contact you by electronic and paper mail.
+
+  If the program does terminal interaction, make it output a short
+notice like this when it starts in an interactive mode:
+
+    <program>  Copyright (C) <year>  <name of author>
+    This program comes with ABSOLUTELY NO WARRANTY; for details type `show w'.
+    This is free software, and you are welcome to redistribute it
+    under certain conditions; type `show c' for details.
+
+The hypothetical commands `show w' and `show c' should show the appropriate
+parts of the General Public License.  Of course, your program's commands
+might be different; for a GUI interface, you would use an "about box".
+
+  You should also get your employer (if you work as a programmer) or school,
+if any, to sign a "copyright disclaimer" for the program, if necessary.
+For more information on this, and how to apply and follow the GNU GPL, see
+<http://www.gnu.org/licenses/>.
+
+  The GNU General Public License does not permit incorporating your program
+into proprietary programs.  If your program is a subroutine library, you
+may consider it more useful to permit linking proprietary applications with
+the library.  If this is what you want to do, use the GNU Lesser General
+Public License instead of this License.  But first, please read
+<http://www.gnu.org/philosophy/why-not-lgpl.html>.

gcc-plugin/ChangeLog

@@ -0,0 +1,2 @@
+This file should be generated when making a tarball from a Subversion
+checkout.  If it's not, please refer to the output "svn log".

gcc-plugin/Makefile.am

@@ -0,0 +1,58 @@
+# StarPU --- Runtime system for heterogeneous multicore architectures.
+#
+# Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+#
+# StarPU is free software; you can redistribute it and/or modify
+# it under the terms of the GNU Lesser General Public License as published by
+# the Free Software Foundation; either version 2.1 of the License, or (at
+# your option) any later version.
+#
+# StarPU is distributed in the hope that it will be useful, but
+# WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+#
+# See the GNU Lesser General Public License in COPYING.LGPL for more details.
+
+SUBDIRS = src tests
+if BUILD_EXAMPLES
+SUBDIRS += examples
+endif
+
+EXTRA_DIST =					\
+  COPYING					\
+  README					\
+  ChangeLog
+
+noinst_HEADERS =				\
+  include/starpu-gcc/opencl.h			\
+  include/starpu-gcc/tasks.h			\
+  include/starpu-gcc/utils.h			\
+  include/starpu-gcc/warn-unregistered.h
+
+
+showcheck:
+	for i in $(SUBDIRS) ; do \
+		make -C $$i showcheck ; \
+	done
+
+# Generate a GNU-style ChangeLog for inclusion in the tarball.
+# It requires network access and may be slow.
+gen-ChangeLog:
+	if test "x$$CHANGELOG" = "xno"; then					\
+	   echo "ChangeLog not built, per user request" >&2;			\
+	elif ! xsltproc --version > /dev/null 2>&1; then			\
+	   echo "xsltproc not found, ChangeLog not generated" >&2;		\
+	elif ! test -d "$(srcdir)/.svn"; then					\
+	   echo "Subversion meta-data not found, ChangeLog not generated" >&2;	\
+	elif ! svn --version > /dev/null 2>&1; then				\
+	   echo "Subversion not found, ChangeLog not generated" >&2;		\
+	else									\
+	   ( cd "$(srcdir)";							\
+	     svn log --xml --verbose ) |					\
+	   xsltproc "$(top_srcdir)/build-aux/svn2cl.xsl" - > "ChangeLog.tmp";	\
+	   mv "ChangeLog.tmp" "$(distdir)/ChangeLog";				\
+	fi
+
+dist-hook: gen-ChangeLog
+
+.PHONY: showcheck dist-hook gen-ChangeLog

gcc-plugin/README

@@ -0,0 +1,12 @@
+StarPU's GCC Plug-In
+=====================
+
+This directory contains an (experimental) GCC plug-in that provides C
+language extensions to make it easier to define and invoke StarPU
+tasks.
+
+Plug-ins are supported starting from GCC 4.5.
+
+To run the test suite, GNU Guile 1.8.x or 2.0.x is needed.
+
+When building from SVN, GNU Bison 2.5+ is required.

gcc-plugin/examples/Makefile.am

@@ -0,0 +1,78 @@
+# StarPU --- Runtime system for heterogeneous multicore architectures.
+#
+# Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+#
+# StarPU is free software; you can redistribute it and/or modify
+# it under the terms of the GNU Lesser General Public License as published by
+# the Free Software Foundation; either version 2.1 of the License, or (at
+# your option) any later version.
+#
+# StarPU is distributed in the hope that it will be useful, but
+# WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+#
+# See the GNU Lesser General Public License in COPYING.LGPL for more details.
+
+if STARPU_USE_CPU
+
+noinst_PROGRAMS =				\
+  matrix-mult stencil5 vector_scal/vector_scal
+
+if !NO_BLAS_LIB
+noinst_PROGRAMS +=				\
+  cholesky/cholesky
+endif !NO_BLAS_LIB
+endif STARPU_USE_CPU
+
+AM_LDFLAGS = $(top_builddir)/src/@LIBSTARPU_LINK@
+AM_LDFLAGS += $(STARPU_OPENCL_LDFLAGS) $(STARPU_CUDA_LDFLAGS) 
+
+AM_CPPFLAGS =						\
+  -I$(top_srcdir)/include				\
+  -I$(top_srcdir)/examples				\
+  $(STARPU_OPENCL_CPPFLAGS) $(STARPU_CUDA_CPPFLAGS)
+
+AM_CFLAGS =							\
+  -fplugin="$(builddir)/../src/.libs/starpu.so"			\
+  -fplugin-arg-starpu-include-dir="$(top_srcdir)/include"	\
+  -fplugin-arg-starpu-verbose					\
+  -Wall
+
+noinst_HEADERS =				\
+  cholesky/cholesky.h				\
+  cholesky/cholesky_kernels.h
+
+EXTRA_DIST = vector_scal/vector_scal_opencl_kernel.cl
+
+if !NO_BLAS_LIB
+cholesky_cholesky_SOURCES	=		\
+	cholesky/cholesky.c		\
+	cholesky/cholesky_models.c	\
+	cholesky/cholesky_kernels.c	\
+	$(top_srcdir)/examples/common/blas.c
+
+cholesky_cholesky_LDADD	=	\
+	$(STARPU_BLAS_LDFLAGS)
+endif
+
+vector_scal_vector_scal_SOURCES = vector_scal/vector_scal.c
+
+if STARPU_USE_CUDA
+
+vector_scal_vector_scal_SOURCES += vector_scal/vector_scal_cuda.cu
+
+.cu.o:
+	$(NVCC) $< -c -o $@ $(NVCCFLAGS)			\
+	  -I$(top_builddir)/include -I$(top_srcdir)/include
+
+else !STARPU_USE_CUDA
+
+EXTRA_DIST += vector_scal/vector_scal_cuda.cu
+
+endif
+
+showcheck:
+	-cat $(TEST_LOGS) /dev/null
+	for i in $(SUBDIRS) ; do \
+		make -C $$i showcheck ; \
+	done

gcc-plugin/examples/cholesky/cholesky.c

@@ -0,0 +1,255 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2009-2011  Université de Bordeaux 1
+ * Copyright (C) 2010  Mehdi Juhoor <mjuhoor@gmail.com>
+ * Copyright (C) 2010, 2011, 2012  Centre National de la Recherche Scientifique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include "cholesky.h"
+#include "cholesky_kernels.h"
+
+#define __heap  __attribute__ ((heap_allocated))
+
+/*
+ *	code to bootstrap the factorization
+ *	and construct the DAG
+ */
+static void dw_cholesky(unsigned nblocks, unsigned size, unsigned ld,
+			float matA[nblocks][nblocks][size/nblocks * size/nblocks])
+{
+	struct timeval start;
+	struct timeval end;
+        int x, y;
+
+	/* create all the DAG nodes */
+	unsigned i,j,k;
+
+        for(x = 0; x < nblocks ;  x++) {
+                for (y = 0; y < nblocks; y++) {
+#pragma starpu register matA[x][y]
+		}
+        }
+
+	gettimeofday(&start, NULL);
+
+	for (k = 0; k < nblocks; k++)
+        {
+#ifdef STARPU_DEVEL
+#  warning deal with prio and models
+#endif
+//                int prio = STARPU_DEFAULT_PRIO;
+//                if (!noprio) prio = STARPU_MAX_PRIO;
+
+		chol_codelet_update_u11(matA[k][k], size/nblocks, ld);
+
+		for (j = k+1; j<nblocks; j++)
+		{
+//                        prio = STARPU_DEFAULT_PRIO;
+//                        if (!noprio&& (j == k+1)) prio = STARPU_MAX_PRIO;
+
+			chol_codelet_update_u21(matA[k][k], matA[k][j], ld, ld, size/nblocks, size/nblocks);
+
+			for (i = k+1; i<nblocks; i++)
+			{
+				if (i <= j)
+                                {
+//                                        prio = STARPU_DEFAULT_PRIO;
+//                                        if (!noprio && (i == k + 1) && (j == k +1) ) prio = STARPU_MAX_PRIO;
+
+					chol_codelet_update_u22(matA[k][i],
+								matA[k][j],
+								matA[i][j],
+								size/nblocks, size/nblocks, size/nblocks, ld, ld, ld);
+                                }
+			}
+		}
+        }
+
+#pragma starpu wait
+
+        for(x = 0; x < nblocks ;  x++) {
+                for (y = 0; y < nblocks; y++) {
+#pragma starpu unregister matA[x][y]
+                }
+        }
+
+	gettimeofday(&end, NULL);
+
+	double timing = (double)((end.tv_sec - start.tv_sec)*1000000 + (end.tv_usec - start.tv_usec));
+	fprintf(stderr, "Computation took (in ms)\n");
+	fprintf(stdout, "%2.2f\n", timing/1000);
+	double flop = (1.0f*size*size*size)/3.0f;
+	fprintf(stderr, "Synthetic GFlops : %2.2f\n", (flop/timing/1000.0f));
+}
+
+int main(int argc, char **argv)
+{
+	/* create a simple definite positive symetric matrix example
+	 *
+	 *	Hilbert matrix : h(i,j) = 1/(i+j+1)
+	 * */
+
+	parse_args(argc, argv);
+
+#ifdef STARPU_DEVEL
+#  warning todo
+#endif
+//	struct starpu_conf conf;
+//	starpu_conf_init(&conf);
+//	conf.sched_policy_name = "heft";
+//	conf.calibrate = 1;
+#pragma starpu initialize
+
+        starpu_helper_cublas_init();
+
+	float bmat[nblocks][nblocks][BLOCKSIZE * BLOCKSIZE] __heap;
+
+	unsigned i,j,x,y;
+        for(x=0 ; x<nblocks ; x++)
+	{
+                for(y=0 ; y<nblocks ; y++)
+		{
+			for (i = 0; i < BLOCKSIZE; i++)
+			{
+				for (j = 0; j < BLOCKSIZE; j++)
+				{
+                                        bmat[x][y][j +i*BLOCKSIZE] = (1.0f/(1.0f+(i+(x*BLOCKSIZE)+j+(y*BLOCKSIZE)))) + ((i+(x*BLOCKSIZE) == j+(y*BLOCKSIZE))?1.0f*size:0.0f);
+					//mat[j +i*size] = ((i == j)?1.0f*size:0.0f);
+				}
+			}
+		}
+	}
+
+
+        if (display) {
+		for(y=0 ; y<nblocks ; y++)
+		{
+			for(x=0 ; x<nblocks ; x++)
+			{
+                                printf("Block %d,%d :\n", x, y);
+				for (j = 0; j < BLOCKSIZE; j++)
+				{
+					for (i = 0; i < BLOCKSIZE; i++)
+					{
+						if (i <= j) {
+							printf("%2.2f\t", bmat[y][x][j +i*BLOCKSIZE]);
+						}
+						else {
+							printf(".\t");
+						}
+					}
+					printf("\n");
+				}
+			}
+		}
+	}
+
+	dw_cholesky(nblocks, size, size/nblocks, bmat);
+
+        if (display) {
+                printf("Results:\n");
+		for(y=0 ; y<nblocks ; y++)
+		{
+			for(x=0 ; x<nblocks ; x++)
+			{
+                                printf("Block %d,%d :\n", x, y);
+				for (j = 0; j < BLOCKSIZE; j++)
+				{
+					for (i = 0; i < BLOCKSIZE; i++)
+					{
+						if (i <= j) {
+							printf("%2.2f\t", bmat[y][x][j +i*BLOCKSIZE]);
+						}
+						else {
+							printf(".\t");
+						}
+					}
+					printf("\n");
+				}
+			}
+		}
+	}
+
+	float rmat[size * size] __heap;
+        for(x=0 ; x<nblocks ; x++) {
+                for(y=0 ; y<nblocks ; y++) {
+                        for (i = 0; i < BLOCKSIZE; i++) {
+                                for (j = 0; j < BLOCKSIZE; j++) {
+                                        rmat[j+(y*BLOCKSIZE)+(i+(x*BLOCKSIZE))*size] = bmat[x][y][j +i*BLOCKSIZE];
+                                }
+                        }
+                }
+        }
+
+	fprintf(stderr, "compute explicit LLt ...\n");
+	for (j = 0; j < size; j++)
+	{
+		for (i = 0; i < size; i++)
+		{
+			if (i > j) {
+				rmat[j+i*size] = 0.0f; // debug
+			}
+		}
+	}
+	float test_mat[size * size] __heap;
+	SSYRK("L", "N", size, size, 1.0f,
+	      rmat, size, 0.0f, test_mat, size);
+
+	fprintf(stderr, "comparing results ...\n");
+        if (display) {
+                for (j = 0; j < size; j++)
+		{
+                        for (i = 0; i < size; i++)
+			{
+                                if (i <= j) {
+                                        printf("%2.2f\t", test_mat[j +i*size]);
+                                }
+                                else {
+                                        printf(".\t");
+                                }
+                        }
+                        printf("\n");
+                }
+        }
+
+	int correctness = 1;
+        for(x = 0; x < nblocks ;  x++)
+	{
+                for (y = 0; y < nblocks; y++)
+		{
+			for (i = (size/nblocks)*x ; i < (size/nblocks)*x+(size/nblocks); i++)
+                                {
+                                        for (j = (size/nblocks)*y ; j < (size/nblocks)*y+(size/nblocks); j++)
+						{
+							if (i <= j)
+								{
+									float orig = (1.0f/(1.0f+i+j)) + ((i == j)?1.0f*size:0.0f);
+									float err = abs(test_mat[j +i*size] - orig);
+									if (err > 0.00001) {
+										fprintf(stderr, "Error[%d, %d] --> %2.2f != %2.2f (err %2.2f)\n", i, j, test_mat[j +i*size], orig, err);
+										correctness = 0;
+										break;
+									}
+								}
+						}
+                                }
+		}
+        }
+
+        starpu_helper_cublas_shutdown();
+#pragma starpu shutdown
+
+	assert(correctness);
+	return 0;
+}

gcc-plugin/examples/cholesky/cholesky.h

@@ -0,0 +1,93 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2009, 2010  Université de Bordeaux 1
+ * Copyright (C) 2010, 2011  Centre National de la Recherche Scientifique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef __DW_CHOLESKY_H__
+#define __DW_CHOLESKY_H__
+
+#include <semaphore.h>
+#include <string.h>
+#include <math.h>
+#include <sys/time.h>
+#include <limits.h>
+#ifdef STARPU_USE_CUDA
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <cublas.h>
+#endif
+
+#include <common/blas.h>
+#include <starpu.h>
+
+#define NMAXBLOCKS	32
+
+#define BLOCKSIZE	(size/nblocks)
+
+static unsigned size = 4*1024;
+static unsigned nblocks = 16;
+static unsigned nbigblocks = 8;
+static unsigned noprio = 0;
+static unsigned display = 0;
+
+static void __attribute__((unused)) parse_args(int argc, char **argv)
+{
+	int i;
+	for (i = 1; i < argc; i++) {
+		if (strcmp(argv[i], "-size") == 0) {
+		        char *argptr;
+			long int ret;
+			ret = strtol(argv[++i], &argptr, 10);
+			if (*argptr != '\0') {
+				fprintf(stderr, "Invalid size %s, aborting \n", argv[i]);
+				exit(EXIT_FAILURE);
+			}
+			if (ret == LONG_MIN || ret == LONG_MAX) {
+				perror("strtol");
+				exit(EXIT_FAILURE);
+			}
+			if (ret == 0) {
+				fprintf(stderr, "0 is not a valid size\n");
+				exit(EXIT_FAILURE);
+			}
+			size = ret;
+		}
+
+		if (strcmp(argv[i], "-nblocks") == 0) {
+		        char *argptr;
+			nblocks = strtol(argv[++i], &argptr, 10);
+		}
+
+		if (strcmp(argv[i], "-nbigblocks") == 0) {
+		        char *argptr;
+			nbigblocks = strtol(argv[++i], &argptr, 10);
+		}
+
+		if (strcmp(argv[i], "-no-prio") == 0) {
+			noprio = 1;
+		}
+
+		if (strcmp(argv[i], "-display") == 0) {
+			display = 1;
+		}
+
+		if (strcmp(argv[i], "-h") == 0) {
+			printf("usage : %s [-display] [-size size] [-nblocks nblocks]\n", argv[0]);
+		}
+	}
+	if (nblocks > size) nblocks = size;
+}
+
+#endif // __DW_CHOLESKY_H__

gcc-plugin/examples/cholesky/cholesky_kernels.c

@@ -0,0 +1,225 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2009, 2010, 2012  Université de Bordeaux 1
+ * Copyright (C) 2010, 2011  Centre National de la Recherche Scientifique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include "cholesky.h"
+#include "cholesky_kernels.h"
+#include "common/blas.h"
+#ifdef STARPU_USE_CUDA
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <cublas.h>
+#ifdef STARPU_HAVE_MAGMA
+#include "magma.h"
+#include "magma_lapack.h"
+#endif
+#endif
+
+/*
+ *   U22
+ */
+
+static inline void chol_common_cpu_codelet_update_u22(const float *left, const float *right, float *center, unsigned dx, unsigned dy, unsigned dz,
+						      unsigned ld21, unsigned ld12, unsigned ld22, int s)
+{
+	//printf("22\n");
+#ifdef STARPU_USE_CUDA
+	cublasStatus st;
+#endif
+
+	switch (s) {
+		case 0:
+			SGEMM("N", "T", dy, dx, dz, -1.0f, left, ld21,
+				right, ld12, 1.0f, center, ld22);
+			break;
+#ifdef STARPU_USE_CUDA
+		case 1:
+			cublasSgemm('n', 't', dy, dx, dz,
+					-1.0f, left, ld21, right, ld12,
+					 1.0f, center, ld22);
+			st = cublasGetError();
+			if (STARPU_UNLIKELY(st != CUBLAS_STATUS_SUCCESS))
+				STARPU_CUBLAS_REPORT_ERROR(st);
+
+			cudaStreamSynchronize(starpu_cuda_get_local_stream());
+
+			break;
+#endif
+		default:
+			STARPU_ABORT();
+			break;
+	}
+}
+
+void chol_cpu_codelet_update_u22(const float *left, const float *right, float *center, unsigned dx, unsigned dy, unsigned dz,
+				 unsigned ld21, unsigned ld12, unsigned ld22)
+	__attribute__ ((task_implementation ("cpu", chol_codelet_update_u22)));
+#ifdef STARPU_USE_CUDA
+void chol_cublas_codelet_update_u22(const float *left, const float *right, float *center, unsigned dx, unsigned dy, unsigned dz,
+                                    unsigned ld21, unsigned ld12, unsigned ld22)
+	__attribute__ ((task_implementation ("cuda", chol_codelet_update_u22)));
+#endif
+
+void chol_cpu_codelet_update_u22(const float *left, const float *right, float *center, unsigned dx, unsigned dy, unsigned dz,
+				 unsigned ld21, unsigned ld12, unsigned ld22)
+{
+	chol_common_cpu_codelet_update_u22(left, right, center, dx, dx, dz, ld21, ld12, ld22, 0);
+}
+
+#ifdef STARPU_USE_CUDA
+void chol_cublas_codelet_update_u22(const float *left, const float *right, float *center, unsigned dx, unsigned dy, unsigned dz,
+                                    unsigned ld21, unsigned ld12, unsigned ld22)
+{
+	chol_common_cpu_codelet_update_u22(left, right, center, dx, dx, dz, ld21, ld12, ld22, 1);
+}
+#endif// STARPU_USE_CUDA
+
+/*
+ * U21
+ */
+
+static inline void chol_common_codelet_update_u21(const float *sub11, float *sub21, unsigned ld11, unsigned ld21, unsigned nx21, unsigned ny21, int s)
+{
+	switch (s) {
+		case 0:
+			STRSM("R", "L", "T", "N", nx21, ny21, 1.0f, sub11, ld11, sub21, ld21);
+			break;
+#ifdef STARPU_USE_CUDA
+		case 1:
+			cublasStrsm('R', 'L', 'T', 'N', nx21, ny21, 1.0f, sub11, ld11, sub21, ld21);
+			cudaStreamSynchronize(starpu_cuda_get_local_stream());
+			break;
+#endif
+		default:
+			STARPU_ABORT();
+			break;
+	}
+}
+
+void chol_cpu_codelet_update_u21(const float *sub11, float *sub21, unsigned ld11, unsigned ld21, unsigned nx21, unsigned ny21)
+	__attribute__ ((task_implementation ("cpu", chol_codelet_update_u21)));
+#ifdef STARPU_USE_CUDA
+void chol_cublas_codelet_update_u21(const float *sub11, float *sub21, unsigned ld11, unsigned ld21, unsigned nx21, unsigned ny21)
+	__attribute__ ((task_implementation ("cuda", chol_codelet_update_u21)));
+#endif
+
+void chol_cpu_codelet_update_u21(const float *sub11, float *sub21, unsigned ld11, unsigned ld21, unsigned nx21, unsigned ny21)
+{
+	chol_common_codelet_update_u21(sub11, sub21, ld11, ld21, nx21, ny21, 0);
+}
+
+#ifdef STARPU_USE_CUDA
+void chol_cublas_codelet_update_u21(const float *sub11, float *sub21, unsigned ld11, unsigned ld21, unsigned nx21, unsigned ny21)
+{
+        chol_common_codelet_update_u21(sub11, sub21, ld11, ld21, nx21, ny21, 1);
+}
+#endif
+
+/*
+ *	U11
+ */
+
+static inline void chol_common_codelet_update_u11(float *sub11, unsigned nx, unsigned ld, int s)
+{
+	unsigned z;
+
+	switch (s) {
+		case 0:
+
+			/*
+			 *	- alpha 11 <- lambda 11 = sqrt(alpha11)
+			 *	- alpha 21 <- l 21	= alpha 21 / lambda 11
+			 *	- A22 <- A22 - l21 trans(l21)
+			 */
+
+			for (z = 0; z < nx; z++)
+			{
+				float lambda11;
+				lambda11 = sqrt(sub11[z+z*ld]);
+				sub11[z+z*ld] = lambda11;
+
+				STARPU_ASSERT(lambda11 != 0.0f);
+
+				SSCAL(nx - z - 1, 1.0f/lambda11, &sub11[(z+1)+z*ld], 1);
+
+				SSYR("L", nx - z - 1, -1.0f,
+							&sub11[(z+1)+z*ld], 1,
+							&sub11[(z+1)+(z+1)*ld], ld);
+			}
+			break;
+#ifdef STARPU_USE_CUDA
+		case 1:
+#ifdef STARPU_HAVE_MAGMA
+			{
+				int ret;
+				int info;
+				ret = magma_spotrf_gpu('L', nx, sub11, ld, &info);
+				if (ret != MAGMA_SUCCESS) {
+					fprintf(stderr, "Error in Magma: %d\n", ret);
+					STARPU_ABORT();
+				}
+				cudaError_t cures = cudaStreamSynchronize(starpu_cuda_get_local_stream());
+				STARPU_ASSERT(!cures);
+			}
+#else
+			for (z = 0; z < nx; z++)
+			{
+				float lambda11;
+				cudaMemcpyAsync(&lambda11, &sub11[z+z*ld], sizeof(float), cudaMemcpyDeviceToHost, starpu_cuda_get_local_stream());
+				cudaStreamSynchronize(starpu_cuda_get_local_stream());
+
+				STARPU_ASSERT(lambda11 != 0.0f);
+
+				lambda11 = sqrt(lambda11);
+
+				cublasSetVector(1, sizeof(float), &lambda11, sizeof(float), &sub11[z+z*ld], sizeof(float));
+
+				cublasSscal(nx - z - 1, 1.0f/lambda11, &sub11[(z+1)+z*ld], 1);
+
+				cublasSsyr('U', nx - z - 1, -1.0f,
+							&sub11[(z+1)+z*ld], 1,
+							&sub11[(z+1)+(z+1)*ld], ld);
+			}
+
+			cudaStreamSynchronize(starpu_cuda_get_local_stream());
+#endif
+			break;
+#endif
+		default:
+			STARPU_ABORT();
+			break;
+	}
+}
+
+
+void chol_cpu_codelet_update_u11(float *mat, unsigned nx, unsigned ld)
+	__attribute__ ((task_implementation ("cpu", chol_codelet_update_u11)));
+#ifdef STARPU_USE_CUDA
+void chol_cublas_codelet_update_u11(float *mat, unsigned nx, unsigned ld)
+	__attribute__ ((task_implementation ("cuda", chol_codelet_update_u11)));
+#endif
+
+void chol_cpu_codelet_update_u11(float *mat, unsigned nx, unsigned ld)
+{
+	chol_common_codelet_update_u11(mat, nx, ld, 0);
+}
+
+#ifdef STARPU_USE_CUDA
+void chol_cublas_codelet_update_u11(float *mat, unsigned nx, unsigned ld)
+{
+	chol_common_codelet_update_u11(mat, nx, ld, 1);
+}
+#endif// STARPU_USE_CUDA

gcc-plugin/examples/cholesky/cholesky_kernels.h

@@ -0,0 +1,33 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2009, 2010  Université de Bordeaux 1
+ * Copyright (C) 2010, 2011  Centre National de la Recherche Scientifique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#ifndef __DW_CHOLESKY_MODELS_H__
+#define __DW_CHOLESKY_MODELS_H__
+
+void chol_codelet_update_u11(float* mat, unsigned nx, unsigned ld)
+	__attribute__ ((task));
+
+void chol_codelet_update_u21(const float *sub11, float *sub21, unsigned ld11,
+			     unsigned ld21, unsigned nx21, unsigned ny21)
+	__attribute__ ((task));
+
+void chol_codelet_update_u22(const float *left, const float *right, float *center,
+			     unsigned dx, unsigned dy, unsigned dz,
+			     unsigned ld21, unsigned ld12, unsigned ld22)
+	__attribute__ ((task));
+
+#endif // __DW_CHOLESKY_MODELS_H__

gcc-plugin/examples/cholesky/cholesky_models.c

@@ -0,0 +1,40 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2009, 2010  Université de Bordeaux 1
+ * Copyright (C) 2010, 2011  Centre National de la Recherche Scientifique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+/*
+ * As a convention, in that file, descr[0] is represented by A,
+ * 				  descr[1] is B ...
+ */
+
+/*
+ *	Number of flops of Gemm 
+ */
+
+struct starpu_perfmodel chol_model_11 = {
+	.type = STARPU_HISTORY_BASED,
+	.symbol = "chol_model_11"
+};
+
+struct starpu_perfmodel chol_model_21 = {
+	.type = STARPU_HISTORY_BASED,
+	.symbol = "chol_model_21"
+};
+
+struct starpu_perfmodel chol_model_22 = {
+	.type = STARPU_HISTORY_BASED,
+	.symbol = "chol_model_22"
+};

gcc-plugin/examples/matrix-mult.c

@@ -0,0 +1,277 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+   Copyright (C) 2011, 2012 Inria
+   Copyright (C) 2010 Sylvain Gault
+
+   StarPU is free software; you can redistribute it and/or modify
+   it under the terms of the GNU Lesser General Public License as published by
+   the Free Software Foundation; either version 2.1 of the License, or (at
+   your option) any later version.
+
+   StarPU is distributed in the hope that it will be useful, but
+   WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  */
+
+/* The classical matrix multiplication example, implemented as a StarPU
+   task invoked for different slices of the matrices.  Each invocation may
+   of course execute in parallel.  */
+
+#ifndef STARPU_GCC_PLUGIN
+# error must be compiled with the StarPU GCC plug-in
+#endif
+
+/* Convenience macro.  */
+#define __heap __attribute__ ((__heap_allocated__))
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <time.h>
+#include <stdint.h>
+#include <sys/time.h>
+
+
+/* Definition of the StarPU task and its CPU implementation.  */
+
+static void matmul (const float *A, const float *B, float *C,
+		    size_t nx, size_t ny, size_t nz)
+  __attribute__ ((task));
+
+static void matmul_cpu (const float *A, const float *B, float *C,
+			size_t nx, size_t ny, size_t nz)
+  __attribute__ ((task_implementation ("cpu", matmul)));
+
+
+static void
+matmul_cpu (const float *A, const float *B, float *C,
+	    size_t nx, size_t ny, size_t nz)
+{
+  size_t i, j, k;
+
+  for (j = 0; j < ny; j++)
+    for (i = 0; i < nx; i++)
+      {
+	for (k = 0; k < nz; k++)
+	  C[j * nx + i] += A[j * nz + k] * B[k * nx + i];
+      }
+}
+
+
+
+static void print_matrix (const float *M, size_t nslicesx,
+			  size_t nslicesy, size_t bxdim,
+			  size_t bydim)
+  __attribute__ ((unused));
+
+static void
+print_matrix (const float *M, size_t nslicesx, size_t nslicesy, size_t bxdim,
+	      size_t bydim)
+{
+  size_t a, b, i, j;
+
+  for (b = 0; b < nslicesy; b++)
+    for (j = 0; j < bydim; j++)
+      {
+	for (a = 0; a < nslicesx; a++)
+	  {
+	    for (i = 0; i < bxdim; i++)
+	      printf ("%f ",
+		      M[b * nslicesx * bxdim * bydim + j * bxdim +
+			a * bxdim * bydim + i]);
+	  }
+	printf ("\b\n");
+      }
+
+  printf ("\n");
+}
+
+
+static float
+my_rand (void)
+{
+  return (float) rand () / (float) RAND_MAX;
+}
+
+static double
+mean (const double *v, size_t size)
+{
+  double sum = 0;
+  size_t i;
+
+  for (i = 0; i < size; i++)
+    sum += v[i];
+
+  return sum / size;
+}
+
+
+static double
+stddev (const double *v, size_t size)
+{
+  double m = mean (v, size);
+  double sqsum = 0;
+  size_t i;
+
+  for (i = 0; i < size; i++)
+    sqsum += (v[i] - m) * (v[i] - m);
+
+  return sqrt (sqsum / size);
+}
+
+
+int
+main (int argc, char **argv)
+{
+  int mloop, nloop = 0;
+  size_t i, j, k;
+  size_t nslicesx;
+  size_t nslicesy;
+  size_t nslicesz;
+  size_t xdim, ydim, zdim;
+  size_t bxdim, bydim, bzdim;
+  struct timeval start_all, end_all;
+  struct timeval start_register, end_register;
+  struct timeval start_tasks, end_tasks;
+  struct timeval start_unregister, end_unregister;
+  struct timeval start_compute, end_compute;
+
+
+  if (argc < 4)
+    {
+      fprintf (stderr, "Usage: %s NLOOPS MATRIX-SIZE NSLICES\n", argv[0]);
+      return EXIT_FAILURE;
+    }
+
+  mloop = nloop = atoi (argv[1]);
+  zdim = ydim = xdim = atoi (argv[2]);
+  nslicesz = nslicesy = nslicesx = atoi (argv[3]);
+  bxdim = xdim / nslicesx;
+  bydim = ydim / nslicesy;
+  bzdim = zdim / nslicesz;
+
+  if (xdim % nslicesx)
+    {
+      fprintf (stderr, "MATRIX-SIZE must be a multiple of NSLICES\n");
+      return EXIT_FAILURE;
+    }
+
+  fprintf (stderr, "running %d loops with %lux%lux%lu matrices and %lux%lux%lu blocks...\n",
+	   nloop,
+	   xdim, ydim, zdim,
+	   bxdim, bydim, bzdim);
+
+  double computetime[nloop];
+  double starttaskstime[nloop];
+
+#pragma starpu initialize
+
+  gettimeofday (&start_all, NULL);
+
+  float A[zdim * ydim] __heap;
+  float B[xdim * zdim] __heap;
+  float C[xdim * ydim] __heap;
+
+  srand (time (NULL));
+  for (i = 0; i < zdim * ydim; i++)
+    A[i] = my_rand () * 100;
+
+  for (i = 0; i < xdim * zdim; i++)
+    B[i] = my_rand () * 100;
+
+#if 0
+  print_matrix (A, nslicesz, nslicesy, bzdim, bydim);
+  print_matrix (B, nslicesx, nslicesz, bxdim, bzdim);
+#endif
+
+  for (i = 0; i < xdim * ydim; i++)
+    C[i] = 0;
+
+  gettimeofday (&start_register, NULL);
+  for (i = 0; i < nslicesy; i++)
+    for (j = 0; j < nslicesz; j++)
+#pragma starpu register &A[i*zdim*bydim + j*bzdim*bydim] (bzdim * bydim)
+
+  for (i = 0; i < nslicesz; i++)
+    for (j = 0; j < nslicesx; j++)
+#pragma starpu register &B[i*xdim*bzdim + j*bxdim*bzdim] (bxdim * bzdim)
+
+  for (i = 0; i < nslicesy; i++)
+    for (j = 0; j < nslicesx; j++)
+#pragma starpu register &C[i*xdim*bydim + j*bxdim*bydim] (bxdim * bydim)
+
+
+  gettimeofday (&end_register, NULL);
+
+  while (nloop--)
+    {
+      gettimeofday (&start_tasks, NULL);
+      gettimeofday (&start_compute, NULL);
+      for (i = 0; i < nslicesy; i++)
+	for (j = 0; j < nslicesx; j++)
+	  for (k = 0; k < nslicesz; k++)
+	    /* Make an asynchronous call to `matmul', leading to the
+	       instantiation of a StarPU task executing in parallel.  */
+	    matmul (&A[i * zdim * bydim + k * bzdim * bydim],
+		    &B[k * xdim * bzdim + j * bxdim * bzdim],
+		    &C[i * xdim * bydim + j * bxdim * bydim], bxdim,
+		    bydim, bzdim);
+
+      gettimeofday (&end_tasks, NULL);
+      starttaskstime[nloop] =
+	(end_tasks.tv_sec - start_tasks.tv_sec) + (end_tasks.tv_usec -
+						   start_tasks.tv_usec) /
+	1000000.0;
+
+      /* Wait for the asynchronous calls to complete.  */
+#pragma starpu wait
+
+      gettimeofday (&end_compute, NULL);
+      computetime[nloop] =
+	(end_compute.tv_sec - start_compute.tv_sec) + (end_compute.tv_usec -
+						       start_compute.
+						       tv_usec) / 1000000.0;
+    }
+
+#if 0
+  print_matrix (C, nslicesx, nslicesy, bxdim, bydim);
+#endif
+
+  gettimeofday (&start_unregister, NULL);
+  for (i = 0; i < nslicesy; i++)
+    for (j = 0; j < nslicesz; j++)
+#pragma starpu unregister &A[i*zdim*bydim + j*bzdim*bydim]
+
+
+  for (i = 0; i < nslicesz; i++)
+    for (j = 0; j < nslicesx; j++)
+#pragma starpu unregister &B[i*xdim*bzdim + j*bxdim*bzdim]
+
+  for (i = 0; i < nslicesy; i++)
+    for (j = 0; j < nslicesx; j++)
+#pragma starpu unregister &C[i*xdim*bydim + j*bxdim*bydim]
+
+  gettimeofday (&end_unregister, NULL);
+  gettimeofday (&end_all, NULL);
+
+#pragma starpu shutdown
+
+  printf ("total: %f\n",
+	  (end_all.tv_sec - start_all.tv_sec) + (end_all.tv_usec -
+						 start_all.tv_usec) /
+	  1000000.0);
+  printf ("register: %f\n",
+	  (end_register.tv_sec - start_register.tv_sec) +
+	  (end_register.tv_usec - start_register.tv_usec) / 1000000.0);
+  printf ("unregister: %f\n",
+	  (end_unregister.tv_sec - start_unregister.tv_sec) +
+	  (end_unregister.tv_usec - start_unregister.tv_usec) / 1000000.0);
+
+  printf ("mean task launch : %f\n", mean (starttaskstime, mloop));
+  printf ("std task launch: %f\n", stddev (starttaskstime, mloop));
+
+  printf ("mean compute: %f\n", mean (computetime, mloop));
+  printf ("std compute: %f\n", stddev (computetime, mloop));
+  printf ("Compute performance : %f GFLOPS\n",
+	  .002 * xdim * ydim * zdim / (mean (computetime, mloop) * 1000000));
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/examples/stencil5.c

@@ -0,0 +1,128 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Centre National de la Recherche Scientifique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+
+#ifndef STARPU_GCC_PLUGIN
+# error must be compiled with the StarPU GCC plug-in
+#endif
+
+/* Definition of the StarPU task and its CPU implementation.  */
+static void stencil5(float *xy, const float *xm1y, const float *xp1y, const float *xym1, const float *xyp1)
+	__attribute__ ((task));
+
+static void stencil5_cpu(float *xy, const float *xm1y, const float *xp1y, const float *xym1, const float *xyp1)
+	__attribute__ ((task_implementation ("cpu", stencil5)));
+
+static void stencil5_cpu(float *xy, const float *xm1y, const float *xp1y, const float *xym1, const float *xyp1)
+{
+	*xy = (*xy + *xm1y + *xp1y + *xym1 + *xyp1) / 5;
+}
+
+#define NITER_DEF 20000
+#define X         10
+#define Y         10
+
+int display = 0;
+int niter = NITER_DEF;
+
+static void parse_args(int argc, char **argv)
+{
+	int i;
+	for (i = 1; i < argc; i++) {
+		if (strcmp(argv[i], "-iter") == 0) {
+			char *argptr;
+			niter = strtol(argv[++i], &argptr, 10);
+		}
+		if (strcmp(argv[i], "-display") == 0) {
+			display = 1;
+		}
+	}
+}
+
+static float my_rand (void)
+{
+	return (float) rand () / (float) RAND_MAX;
+}
+
+int main(int argc, char **argv)
+{
+        int x, y;
+        float mean=0;
+        float matrix[X][Y];
+
+        parse_args(argc, argv);
+
+	srand (time (NULL));
+        for(x = 0; x < X; x++) {
+                for (y = 0; y < Y; y++) {
+                        matrix[x][y] = my_rand () * 100;
+                        mean += matrix[x][y];
+                }
+        }
+        mean /= (x*y);
+
+        if (display) {
+                fprintf(stdout, "mean=%f\n", mean);
+                for(x = 0; x < X; x++) {
+                        for (y = 0; y < Y; y++) {
+                                fprintf(stdout, "%3f ", matrix[x][y]);
+                        }
+                        fprintf(stdout, "\n");
+                }
+        }
+
+#pragma starpu initialize
+
+        for(x = 0; x < X; x++) {
+		for (y = 0; y < Y; y++) {
+#pragma starpu register &matrix[x][y] 1
+		}
+	}
+
+	while(niter--) {
+                for (x = 1; x < X-1; x++) {
+                        for (y = 1; y < Y-1; y++) {
+                                stencil5(&matrix[x][y], &matrix[x-1][y], &matrix[x+1][y],
+					 &matrix[x][y-1], &matrix[x][y+1]);
+                        }
+                }
+        }
+
+#pragma starpu wait
+
+        for(x = 0; x < X; x++) {
+                for (y = 0; y < Y; y++) {
+#pragma starpu unregister &matrix[x][y]
+                }
+        }
+
+#pragma starpu shutdown
+
+        if (display) {
+                fprintf(stdout, "mean=%f\n", mean);
+                for(x = 0; x < X; x++) {
+                        for (y = 0; y < Y; y++) {
+                                fprintf(stdout, "%3f ", matrix[x][y]);
+                        }
+                        fprintf(stdout, "\n");
+                }
+        }
+
+	return EXIT_SUCCESS;
+}

gcc-plugin/examples/vector_scal/vector_scal.c

@@ -0,0 +1,192 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+/* This example showcases features of the StarPU GCC plug-in.  It defines a
+   "vector scaling" task with multiple CPU implementations, an OpenCL
+   implementation, and a CUDA implementation.
+
+   Compiling it without `-fplugin=starpu.so' yields valid sequential code.  */
+
+#include <math.h>
+#include <stdbool.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+
+/* Declare and define the standard CPU implementation.  */
+
+static void vector_scal (unsigned int size, float vector[size], float factor)
+  __attribute__ ((task));
+
+/* The CPU implementation.  */
+static void
+vector_scal (unsigned int size, float vector[size], float factor)
+{
+  unsigned int i;
+  for (i = 0; i < size; i++)
+    vector[i] *= factor;
+}
+
+
+#if defined STARPU_GCC_PLUGIN && defined __SSE__
+/* The SSE-capable CPU implementation.  */
+
+#include <xmmintrin.h>
+
+static void vector_scal_sse (unsigned int size, float vector[size], float factor)
+  __attribute__ ((task_implementation ("cpu", vector_scal)));
+
+static void
+vector_scal_sse (unsigned int size, float vector[size], float factor)
+{
+  unsigned int n_iterations = size / 4;
+
+  __m128 *VECTOR = (__m128 *) vector;
+  __m128 _FACTOR __attribute__ ((aligned (16)));
+  _FACTOR = _mm_set1_ps (factor);
+
+  unsigned int i;
+  for (i = 0; i < n_iterations; i++)
+    VECTOR[i] = _mm_mul_ps (_FACTOR, VECTOR[i]);
+
+  unsigned int remainder = size % 4;
+  if (remainder != 0)
+    {
+      unsigned int start = 4 * n_iterations;
+      for (i = start; i < start + remainder; ++i)
+	vector[i] = factor * vector[i];
+    }
+}
+#endif /* __SSE__ */
+
+
+/* Declaration and definition of the OpenCL implementation.  */
+
+#if defined STARPU_GCC_PLUGIN && defined STARPU_USE_OPENCL
+
+#include <starpu_opencl.h>
+
+/* The OpenCL programs, loaded from `main'.  */
+static struct starpu_opencl_program cl_programs;
+
+static void vector_scal_opencl (unsigned int size, float vector[size], float factor)
+  __attribute__ ((task_implementation ("opencl", vector_scal)));
+
+static void
+vector_scal_opencl (unsigned int size, float vector[size], float factor)
+{
+  int id, devid, err;
+  cl_kernel kernel;
+  cl_command_queue queue;
+  cl_event event;
+
+  cl_mem val = (cl_mem) vector;
+
+  id = starpu_worker_get_id ();
+  devid = starpu_worker_get_devid (id);
+
+  /* Prepare to invoke the kernel.  In the future, this will be largely
+     automated.  */
+  err = starpu_opencl_load_kernel (&kernel, &queue, &cl_programs,
+				   "vector_mult_opencl", devid);
+  if (err != CL_SUCCESS)
+    STARPU_OPENCL_REPORT_ERROR (err);
+
+  err = clSetKernelArg (kernel, 0, sizeof (val), &val);
+  err |= clSetKernelArg (kernel, 1, sizeof (size), &size);
+  err |= clSetKernelArg (kernel, 2, sizeof (factor), &factor);
+  if (err)
+    STARPU_OPENCL_REPORT_ERROR (err);
+
+  size_t global = size, local = 1;
+  err = clEnqueueNDRangeKernel (queue, kernel, 1, NULL, &global, &local, 0,
+				NULL, &event);
+  if (err != CL_SUCCESS)
+    STARPU_OPENCL_REPORT_ERROR (err);
+
+  clFinish (queue);
+  starpu_opencl_collect_stats (event);
+  clReleaseEvent (event);
+
+  starpu_opencl_release_kernel (kernel);
+}
+
+#endif
+
+
+#ifdef STARPU_USE_CUDA
+
+/* Declaration of the CUDA implementation.  The definition itself is in the
+   `.cu' file itself.  */
+
+extern void vector_scal_cuda (unsigned int size, float vector[size], float factor)
+  __attribute__ ((task_implementation ("cuda", vector_scal)));
+
+#endif
+
+
+#define EPSILON 1e-3
+static bool
+check (size_t size, float vector[size], float factor)
+{
+  size_t i;
+
+  for (i = 0; i < size; i++)
+    {
+      if (fabs(vector[i] - i * factor) > i*factor*EPSILON)
+        {
+          fprintf(stderr, "%.2f != %.2f\n", vector[i], i*factor);
+          return false;
+        }
+    }
+  return true;
+}
+
+
+int
+main (void)
+{
+  bool valid;
+
+#pragma starpu initialize
+
+#if defined STARPU_GCC_PLUGIN && defined STARPU_USE_OPENCL
+  starpu_opencl_load_opencl_from_file ("vector_scal_opencl_kernel.cl",
+				       &cl_programs, "");
+#endif
+
+#define NX     0x100000
+#define FACTOR 3.14
+
+  {
+    float vector[NX] __attribute__ ((heap_allocated, registered));
+
+    size_t i;
+    for (i = 0; i < NX; i++)
+      vector[i] = (float) i;
+
+    vector_scal (NX, vector, FACTOR);
+
+#pragma starpu wait
+
+    valid = check (NX, vector, FACTOR);
+
+  } /* VECTOR is automatically freed here.  */
+
+#pragma starpu shutdown
+
+  return valid ? EXIT_SUCCESS : EXIT_FAILURE;
+}

gcc-plugin/examples/vector_scal/vector_scal_cuda.cu

@@ -0,0 +1,44 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+ * Copyright (C) 2010, 2011  Centre National de la Recherche Scientifique
+ * Copyright (C) 2010  Université de Bordeaux 1
+ *
+ * StarPU is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU Lesser General Public License as published by
+ * the Free Software Foundation; either version 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * StarPU is distributed in the hope that it will be useful, but
+ * WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+ *
+ * See the GNU Lesser General Public License in COPYING.LGPL for more details.
+ */
+
+/* CUDA implementation of the `vector_scal' task.  */
+
+#include <starpu.h>
+#include <starpu_cuda.h>
+#include <stdlib.h>
+
+static __global__ void
+vector_mult_cuda (float *val, unsigned n, float factor)
+{
+  unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+
+  if (i < n)
+    val[i] *= factor;
+}
+
+extern "C" void
+vector_scal_cuda (size_t size, float vector[], float factor)
+{
+  unsigned threads_per_block = 64;
+  unsigned nblocks = (size + threads_per_block - 1) / threads_per_block;
+
+  vector_mult_cuda <<< nblocks, threads_per_block, 0,
+    starpu_cuda_get_local_stream () >>> (vector, size, factor);
+
+  cudaStreamSynchronize (starpu_cuda_get_local_stream ());
+}

gcc-plugin/examples/vector_scal/vector_scal_opencl_kernel.cl

@@ -0,0 +1,38 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010, 2011  Centre National de la Recherche Scientifique
+ * Copyright (C) 2010, 2011  Université de Bordeaux 1
+ *
+ * Redistribution  and  use  in  source and binary forms, with or without
+ * modification,  are  permitted  provided  that the following conditions
+ * are met:
+ *
+ * * Redistributions  of  source  code  must  retain  the above copyright
+ *   notice,  this  list  of  conditions  and  the  following  disclaimer.
+ * * Redistributions  in  binary  form must reproduce the above copyright
+ *   notice,  this list of conditions and the following disclaimer in the
+ *   documentation  and/or other materials provided with the distribution.
+ * * The name of the author may not be used to endorse or promote products
+ *   derived from this software without specific prior written permission.
+ *
+ * THIS  SOFTWARE  IS  PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * ``AS IS''  AND  ANY  EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
+ * LIMITED  TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
+ * A  PARTICULAR  PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+ * HOLDERS OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL
+ * SPECIAL,  EXEMPLARY,  OR  CONSEQUENTIAL  DAMAGES  (INCLUDING,  BUT NOT
+ * LIMITED  TO,  PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE
+ * DATA,  OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+ * THEORY  OF  LIABILITY,  WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING  NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF  THIS  SOFTWARE,  EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+__kernel void vector_mult_opencl(__global float* val, unsigned int nx, float factor)
+{
+        const int i = get_global_id(0);
+        if (i < nx) {
+                val[i] *= factor;
+        }
+}
+

gcc-plugin/include/starpu-gcc/config.h.in

@@ -0,0 +1,41 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Hand-written config.h template for GCC-StarPU.  Autoheader's generated
+   template cannot be used because it defines PACKAGE_NAME & co., which GCC's
+   <auto-host.h> header shamelessly defines.  */
+
+#undef HAVE_DECL_BUILD_CALL_EXPR_LOC_ARRAY
+
+#undef HAVE_DECL_BUILD_CALL_EXPR_LOC_VEC
+
+#undef HAVE_DECL_BUILD_ARRAY_REF
+
+#undef HAVE_DECL_BUILD_ZERO_CST
+
+#undef HAVE_DECL_BUILTIN_DECL_EXPLICIT
+
+#undef HAVE_ATTRIBUTE_SPEC_AFFECTS_TYPE_IDENTITY
+
+#undef HAVE_C_FAMILY_C_COMMON_H
+#undef HAVE_C_COMMON_H
+
+#undef HAVE_C_FAMILY_C_PRAGMA_H
+#undef HAVE_C_PRAGMA_H
+
+#undef STARPU_INCLUDE_DIR
+
+#undef HAVE_DECL_PTR_DEREFS_MAY_ALIAS_P

gcc-plugin/include/starpu-gcc/opencl.h

@@ -0,0 +1,28 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Inria
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#pragma once
+
+#include <gcc-plugin.h>
+#include <tree.h>
+#include <cpplib.h>
+
+#include <starpu-gcc/utils.h>
+
+extern tree opencl_include_dirs;
+
+extern void handle_pragma_opencl (struct cpp_reader *reader);
+extern void validate_opencl_argument_type (location_t loc, const_tree type);

gcc-plugin/include/starpu-gcc/tasks.h

@@ -0,0 +1,58 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Inria
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Various utilities.  */
+
+#pragma once
+
+#include <starpu-gcc/config.h>
+
+/* Note: Users of this file must first include <starpu.h>, but we don't do
+   that here because it has to be done very early, to avoid the dreaded "use
+   of poisoned malloc" in xmmintrin.h.  */
+
+#include <starpu-gcc/utils.h>
+
+
+extern const char task_attribute_name[];
+extern const char task_implementation_attribute_name[];
+extern const char output_attribute_name[];
+
+extern const char task_implementation_wrapper_attribute_name[];
+extern const char task_implementation_list_attribute_name[];
+
+extern bool task_p (const_tree decl);
+extern bool task_implementation_p (const_tree decl);
+extern int task_implementation_where (const_tree task_impl);
+extern int task_implementation_target_to_int (const_tree target);
+extern tree task_implementation_task (const_tree task_impl);
+extern tree task_codelet_declaration (const_tree task_decl);
+extern tree task_implementation_list (const_tree task_decl);
+extern tree task_pointer_parameter_types (const_tree task_decl);
+extern int task_where (const_tree task_decl);
+extern tree task_implementation_wrapper (const_tree task_impl);
+extern enum starpu_access_mode access_mode (const_tree type);
+extern bool output_type_p (const_tree type);
+
+extern tree codelet_type (void);
+extern void taskify_function (tree fn);
+extern tree build_codelet_identifier (tree task_decl);
+extern tree build_codelet_declaration (tree task_decl);
+extern tree build_codelet_initializer (tree task_decl);
+extern tree declare_codelet (tree task_decl);
+extern void define_task (tree task_decl);
+extern void add_task_implementation (tree task_decl, tree fn,
+				     const_tree where);

gcc-plugin/include/starpu-gcc/utils.h

@@ -0,0 +1,152 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Inria
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Various utilities.  */
+
+#pragma once
+
+#include <starpu-gcc/config.h>
+
+#include <unistd.h>
+
+/* GCC 4.7 requires compilation with `g++', and C++ lacks a number of GNU C
+   features, so work around that.  */
+
+#ifdef __cplusplus
+
+/* G++ doesn't implement nested functions, so use C++11 lambdas instead.  */
+
+# include <functional>
+
+# define local_define(ret, name, parms)     auto name = [=]parms
+# define function_parm(ret, name, parms)    std::function<ret parms> name
+
+/* G++ lacks designated initializers.  */
+# define designated_field_init(name, value) value /* XXX: cross fingers */
+
+#else  /* !__cplusplus */
+
+/* GNU C nested functions.  */
+
+# define local_define(ret, name, parms)	    ret name parms
+# define function_parm(ret, name, parms)    ret (*name) parms
+
+/* Designated field initializer.  */
+
+# define designated_field_init(name, value) .name = value
+
+#endif /* !__cplusplus */
+
+
+/* List and vector utilities, à la SRFI-1.  */
+
+extern tree chain_trees (tree t, ...)
+  __attribute__ ((sentinel));
+
+extern tree filter (function_parm (bool, pred, (const_tree)), tree t);
+extern tree list_remove (function_parm (bool, pred, (const_tree)), tree t);
+extern tree map (function_parm (tree, func, (const_tree)), tree t);
+extern void for_each (function_parm (void, func, (tree)), tree t);
+extern size_t count (function_parm (bool, pred, (const_tree)), const_tree t);
+
+
+/* Compatibility tricks & workarounds.  */
+
+#include <tree.h>
+#include <vec.h>
+
+/* This declaration is from `c-tree.h', but that header doesn't get
+   installed.  */
+
+extern tree xref_tag (enum tree_code, tree);
+
+#if !HAVE_DECL_BUILTIN_DECL_EXPLICIT
+
+/* This function was introduced in GCC 4.7 as a replacement for the
+   `built_in_decls' array.  */
+
+static inline tree
+builtin_decl_explicit (enum built_in_function fncode)
+{
+  return built_in_decls[fncode];
+}
+
+#endif
+
+#if !HAVE_DECL_BUILD_CALL_EXPR_LOC_ARRAY
+
+extern tree build_call_expr_loc_array (location_t loc, tree fndecl, int n,
+				       tree *argarray);
+
+#endif
+
+#if !HAVE_DECL_BUILD_CALL_EXPR_LOC_VEC
+
+extern tree build_call_expr_loc_vec (location_t loc, tree fndecl,
+				     VEC(tree,gc) *vec);
+
+#endif
+
+#if !HAVE_DECL_BUILD_ZERO_CST
+
+extern tree build_zero_cst (tree type);
+
+#endif
+
+#ifndef VEC_qsort
+
+/* This macro is missing in GCC 4.5.  */
+
+# define VEC_qsort(T,V,CMP) qsort(VEC_address (T,V), VEC_length(T,V),	\
+				  sizeof (T), CMP)
+
+#endif
+
+
+/* Helpers.  */
+
+extern bool verbose_output_p;
+
+extern tree build_pointer_lookup (tree pointer);
+extern tree build_starpu_error_string (tree error_var);
+extern tree build_constructor_from_unsorted_list (tree type, tree vals);
+extern tree read_pragma_expressions (const char *pragma, location_t loc);
+extern tree type_decl_for_struct_tag (const char *tag);
+extern tree build_function_arguments (tree fn);
+extern tree build_error_statements (location_t, tree,
+				    function_parm (tree, f, (tree)),
+				    const char *, ...)
+  __attribute__ ((format (printf, 4, 5)));
+
+extern bool void_type_p (const_tree lst);
+extern bool pointer_type_p (const_tree lst);
+
+/* Lookup the StarPU function NAME in the global scope and store the result
+   in VAR (this can't be done from `lower_starpu'.)  */
+
+#define LOOKUP_STARPU_FUNCTION(var, name)				\
+  if ((var) == NULL_TREE)						\
+    {									\
+      (var) = lookup_name (get_identifier (name));			\
+      gcc_assert ((var) != NULL_TREE && TREE_CODE (var) == FUNCTION_DECL); \
+    }
+
+
+/* Don't warn about the unused `gcc_version' variable, from
+   <plugin-version.h>.  */
+
+static const struct plugin_gcc_version *starpu_gcc_version
+  __attribute__ ((__unused__)) = &gcc_version;

gcc-plugin/include/starpu-gcc/warn-unregistered.h

@@ -0,0 +1,25 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Inria
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Buffer registration warning pass.  */
+
+#pragma once
+
+#include <tree-pass.h>
+
+/* A pass that warns about use of possibly unregistered buffers.  */
+
+extern struct opt_pass pass_warn_starpu_unregistered;

gcc-plugin/src/Makefile.am

@@ -0,0 +1,53 @@
+# StarPU --- Runtime system for heterogeneous multicore architectures.
+#
+# Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+#
+# StarPU is free software; you can redistribute it and/or modify
+# it under the terms of the GNU Lesser General Public License as published by
+# the Free Software Foundation; either version 2.1 of the License, or (at
+# your option) any later version.
+#
+# StarPU is distributed in the hope that it will be useful, but
+# WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+#
+# See the GNU Lesser General Public License in COPYING.LGPL for more details.
+
+# `LIBRARIES' would be more appropriate than `LTLIBRARIES' but it
+# requires a name prefixed by `lib'.
+gccplugin_LTLIBRARIES = starpu.la
+
+starpu_la_SOURCES =				\
+  c-expr.y					\
+  opencl.c					\
+  starpu.c					\
+  tasks.c					\
+  utils.c
+
+if HAVE_PTR_DEREFS_MAY_ALIAS_P
+
+# Only for GCC >= 4.6.
+starpu_la_SOURCES += warn-unregistered.c
+
+endif
+
+# Use the Yacc-compatibility mode so that Bison doesn't error out upon
+# reduce/reduce conflicts.
+AM_YFLAGS = -y
+
+AM_CPPFLAGS =						\
+  -I$(top_builddir)/gcc-plugin/include			\
+  -I$(top_srcdir)/gcc-plugin/include			\
+  -I$(top_srcdir)/include				\
+  -I$(GCC_PLUGIN_INCLUDE_DIR) -Wall -DYYERROR_VERBOSE=1	\
+  $(STARPU_CUDA_CPPFLAGS) $(STARPU_OPENCL_CPPFLAGS)
+
+AM_LDFLAGS = -module
+
+# Use either `gcc' or `g++', whichever is appropriate to build
+# plug-ins for this version of GCC.
+AM_LIBTOOLFLAGS = --tag="$(GCC_FOR_PLUGIN_LIBTOOL_TAG)"
+CC = $(GCC_FOR_PLUGIN)
+
+showcheck:
+	-cat $(TEST_LOGS) /dev/null

gcc-plugin/src/c-expr.y

@@ -0,0 +1,304 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Parser for simple C expressions in pragmas.  */
+
+%define api.pure
+%parse-param { location_t loc }
+%parse-param { const char *pragma }
+%parse-param { tree *seq }
+%debug
+
+%{
+  #include <starpu-gcc/config.h>
+
+  #include <gcc-plugin.h>
+  #include <plugin.h>
+  #include <tree.h>
+  #include <cpplib.h>
+
+  #ifdef HAVE_C_FAMILY_C_COMMON_H
+  # include <c-family/c-common.h>
+  #elif HAVE_C_COMMON_H
+  # include <c-common.h>
+  #endif
+
+  #ifdef HAVE_C_FAMILY_C_PRAGMA_H
+  # include <c-family/c-pragma.h>
+  #elif HAVE_C_PRAGMA_H
+  # include <c-pragma.h>
+  #endif
+
+  #include <diagnostic.h>
+
+  #if !HAVE_DECL_BUILD_ARRAY_REF
+  /* This declaration is missing in GCC 4.6.1.  */
+  extern tree build_array_ref (location_t loc, tree array, tree index);
+  #endif
+
+
+  #define YYSTYPE tree
+  #define YYLTYPE location_t
+
+  static void
+  yyerror (location_t loc, const char *pragma, tree *seq,
+	   char const *message)
+  {
+    error_at (loc, "parse error in pragma %qs: %s", pragma, message);
+  }
+
+  /* Return SOMETHING if it's a VAR_DECL, an identifier bound to a VAR_DECL,
+     or another object; raise an error otherwise.  */
+
+  static tree
+  ensure_bound (location_t loc, tree something)
+  {
+    gcc_assert (something != NULL_TREE);
+
+    if (DECL_P (something))
+      return something;
+    else if (TREE_CODE (something) == IDENTIFIER_NODE)
+      {
+	tree var = lookup_name (something);
+	if (var == NULL_TREE)
+	  {
+	    error_at (loc, "unbound variable %qE", something);
+	    return error_mark_node;
+	  }
+	else
+	  return var;
+      }
+
+    return something;
+  }
+
+  static tree
+  build_component_ref (location_t loc, tree what, tree field)
+  {
+    sorry ("struct field access not implemented yet"); /* XXX */
+    return error_mark_node;
+  }
+
+  /* Interpret the string beneath CST, and return a new string constant.  */
+  static tree
+  interpret_string (const_tree cst)
+  {
+    gcc_assert (TREE_CODE (cst) == STRING_CST);
+
+    cpp_string input, interpreted;
+    input.text = (unsigned char *) TREE_STRING_POINTER (cst);
+    input.len = TREE_STRING_LENGTH (cst);
+
+    bool success;
+    success = cpp_interpret_string (parse_in, &input, 1, &interpreted,
+				    CPP_STRING);
+    gcc_assert (success);
+
+    return build_string (interpreted.len, (char *) interpreted.text);
+  }
+%}
+
+%code {
+  /* Mapping of libcpp token names to Bison-generated token names.  This is
+     not ideal but Bison cannot be told to use the `enum cpp_ttype'
+     values.  */
+
+#define STARPU_CPP_TOKENS			\
+  TK (CPP_NAME)					\
+  TK (CPP_NUMBER)				\
+  TK (CPP_AND)					\
+  TK (CPP_OPEN_SQUARE)				\
+  TK (CPP_CLOSE_SQUARE)				\
+  TK (CPP_OPEN_PAREN)				\
+  TK (CPP_CLOSE_PAREN)				\
+  TK (CPP_PLUS)					\
+  TK (CPP_MINUS)				\
+  TK (CPP_MULT)					\
+  TK (CPP_DIV)					\
+  TK (CPP_DOT)					\
+  TK (CPP_DEREF)				\
+  TK (CPP_STRING)
+
+#ifndef __cplusplus
+
+  static const int cpplib_bison_token_map[] =
+    {
+# define TK(x) [x] = Y ## x,
+      STARPU_CPP_TOKENS
+# undef TK
+    };
+
+#else /* __cplusplus */
+
+  /* No designated initializers in C++.  */
+  static int cpplib_bison_token_map[CPP_PADDING];
+
+#endif	/* __cplusplus */
+
+  static int
+  yylex (YYSTYPE *lvalp)
+  {
+    int ret;
+    enum cpp_ttype type;
+    location_t loc;
+
+#ifdef __cplusplus
+    if (cpplib_bison_token_map[CPP_NAME] != YCPP_NAME)
+      {
+	/* Initialize the table.  */
+# define TK(x) cpplib_bison_token_map[x] = Y ## x;
+	STARPU_CPP_TOKENS
+# undef TK
+      }
+#endif
+
+    /* First check whether EOL is reached, because the EOL token needs to be
+       left to the C parser.  */
+    type = cpp_peek_token (parse_in, 0)->type;
+    if (type == CPP_PRAGMA_EOL)
+      ret = -1;
+    else
+      {
+	/* Tell the lexer to not concatenate adjacent strings like cpp and
+	   `pragma_lex' normally do, because we want to be able to
+	   distinguish adjacent STRING_CST.  */
+	type = c_lex_with_flags (lvalp, &loc, NULL, C_LEX_STRING_NO_JOIN);
+
+	if (type == CPP_STRING)
+	  /* XXX: When using `C_LEX_STRING_NO_JOIN', `c_lex_with_flags'
+	     doesn't call `cpp_interpret_string', leaving us with an
+	     uninterpreted string (with quotes, etc.)  This hack works around
+	     that.  */
+	  *lvalp = interpret_string (*lvalp);
+
+	if (type < sizeof cpplib_bison_token_map / sizeof cpplib_bison_token_map[0])
+	  ret = cpplib_bison_token_map[type];
+	else
+	  ret = -1;
+      }
+
+    return ret;
+  }
+}
+
+%token YCPP_NAME "identifier"
+%token YCPP_NUMBER "integer"
+%token YCPP_AND "&"
+%token YCPP_OPEN_SQUARE "["
+%token YCPP_CLOSE_SQUARE "]"
+%token YCPP_OPEN_PAREN "("
+%token YCPP_CLOSE_PAREN ")"
+%token YCPP_PLUS "+"
+%token YCPP_MINUS "-"
+%token YCPP_MULT "*"
+%token YCPP_DIV "/"
+%token YCPP_DOT "."
+%token YCPP_DEREF "->"
+%token YCPP_STRING "string"
+
+%% /* Grammar rules.  */
+
+ /* Always return a TREE_LIST rather than a raw chain, because the elements
+    of that list may be already chained for other purposes---e.g., PARM_DECLs
+    of a function are chained together.  */
+
+sequence: expression {
+          gcc_assert (*seq == NULL_TREE);
+	  *seq = tree_cons (NULL_TREE, $1, NULL_TREE);
+	  $$ = *seq;
+      }
+      | expression sequence {
+	  gcc_assert ($2 == *seq);
+	  *seq = tree_cons (NULL_TREE, $1, $2);
+	  $$ = *seq;
+      }
+;
+
+expression: binary_expression
+;
+
+/* XXX: `ensure_bound' below leads to errors raised even for non-significant
+   arguments---e.g., junk after pragma.  */
+identifier: YCPP_NAME  { $$ = ensure_bound (loc, $1); }
+;
+
+binary_expression: additive_expression
+;
+
+multiplicative_expression: multiplicative_expression YCPP_MULT cast_expression {
+       $$ = build_binary_op (UNKNOWN_LOCATION, MULT_EXPR, $1, $3, 0);
+     }
+     | multiplicative_expression YCPP_DIV cast_expression {
+       $$ = build_binary_op (UNKNOWN_LOCATION, TRUNC_DIV_EXPR, $1, $3, 0);
+     }
+     | cast_expression
+;
+
+additive_expression: multiplicative_expression
+     | additive_expression YCPP_PLUS multiplicative_expression {
+       $$ = build_binary_op (UNKNOWN_LOCATION, PLUS_EXPR, $1, $3, 0);
+     }
+     | additive_expression YCPP_MINUS multiplicative_expression {
+       $$ = build_binary_op (UNKNOWN_LOCATION, MINUS_EXPR, $1, $3, 0);
+     }
+;
+
+cast_expression: unary_expression
+		 /* XXX: No support for '(' TYPE-NAME ')' UNARY-EXPRESSION.  */
+;
+
+unary_expression: postfix_expression
+     | YCPP_AND cast_expression {
+       $$ = build_addr (ensure_bound (loc, $2), current_function_decl);
+     }
+;
+
+postfix_expression:
+       primary_expression
+     | postfix_expression YCPP_OPEN_SQUARE expression YCPP_CLOSE_SQUARE {
+#if 1
+	 /* Build the array ref with proper error checking.  */
+	 $$ = build_array_ref (loc, ensure_bound (loc, $1),
+			       ensure_bound (loc, $3));
+#else /* TIMTOWTDI */
+	 $$ = build_indirect_ref (loc,
+	       build_binary_op (loc, PLUS_EXPR, ensure_bound (loc, $1), ensure_bound (loc, $3), 0),
+		RO_ARRAY_INDEXING);
+#endif
+     }
+     | postfix_expression YCPP_DOT identifier {
+        $$ = build_component_ref (loc, ensure_bound (loc, $1), $2);
+     }
+     | postfix_expression YCPP_DEREF identifier {
+        $$ = build_component_ref (loc,
+               build_indirect_ref (loc, ensure_bound (loc, $1), RO_ARRAY_INDEXING),
+               $2);
+     }
+;
+
+primary_expression: identifier
+     | constant
+     | string_literal
+     | YCPP_OPEN_PAREN expression YCPP_CLOSE_PAREN { $$ = $2; }
+;
+
+constant: YCPP_NUMBER { $$ = $1; }
+;
+
+string_literal: YCPP_STRING { $$ = $1; }
+;
+
+%%

gcc-plugin/src/opencl.c

@@ -0,0 +1,698 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Inria
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#include <starpu-gcc/config.h>
+
+/* We must include starpu.h here, otherwise gcc will complain about a poisoned
+   malloc in xmmintrin.h.  */
+#include <starpu.h>
+
+#include <gcc-plugin.h>
+#include <plugin-version.h>
+#include <plugin.h>
+#include <tree.h>
+#include <tree-iterator.h>
+#include <gimple.h>
+#include <cgraph.h>
+#include <toplev.h>
+#include <langhooks.h>
+
+#ifdef HAVE_C_FAMILY_C_COMMON_H
+# include <c-family/c-common.h>
+#elif HAVE_C_COMMON_H
+# include <c-common.h>
+#endif
+
+#include <stdlib.h>
+#include <unistd.h>
+#include <sys/mman.h>
+
+#include <starpu-gcc/utils.h>
+#include <starpu-gcc/tasks.h>
+
+
+/* Search path for OpenCL source files for the `opencl' pragma, as a
+   `TREE_LIST'.  */
+tree opencl_include_dirs = NULL_TREE;
+
+/* Names of data structures defined in <starpu.h>.  */
+static const char opencl_program_struct_tag[] = "starpu_opencl_program";
+
+
+/* Return the type corresponding to OPENCL_PROGRAM_STRUCT_TAG.  */
+
+static tree
+opencl_program_type (void)
+{
+  tree t = TREE_TYPE (type_decl_for_struct_tag (opencl_program_struct_tag));
+
+  if (TYPE_SIZE (t) == NULL_TREE)
+    {
+      /* Incomplete type definition, for instance because <starpu_opencl.h>
+	 wasn't included.  */
+      error_at (UNKNOWN_LOCATION, "StarPU OpenCL support is lacking");
+      t = error_mark_node;
+    }
+
+  return t;
+}
+
+static tree
+opencl_kernel_type (void)
+{
+  tree t = lookup_name (get_identifier ("cl_kernel"));
+  gcc_assert (t != NULL_TREE);
+  if (TREE_CODE (t) == TYPE_DECL)
+    t = TREE_TYPE (t);
+  gcc_assert (TYPE_P (t));
+  return t;
+}
+
+static tree
+opencl_command_queue_type (void)
+{
+  tree t = lookup_name (get_identifier ("cl_command_queue"));
+  gcc_assert (t != NULL_TREE);
+  if (TREE_CODE (t) == TYPE_DECL)
+    t = TREE_TYPE (t);
+  gcc_assert (TYPE_P (t));
+  return t;
+}
+
+static tree
+opencl_event_type (void)
+{
+  tree t = lookup_name (get_identifier ("cl_event"));
+  gcc_assert (t != NULL_TREE);
+  if (TREE_CODE (t) == TYPE_DECL)
+    t = TREE_TYPE (t);
+  gcc_assert (TYPE_P (t));
+  return t;
+}
+
+
+
+/* Return a private global string literal VAR_DECL, whose contents are the
+   LEN bytes at CONTENTS.  */
+
+static tree
+build_string_variable (location_t loc, const char *name_seed,
+		       const char *contents, size_t len)
+{
+  tree decl;
+
+  decl = build_decl (loc, VAR_DECL, create_tmp_var_name (name_seed),
+		     string_type_node);
+  TREE_PUBLIC (decl) = false;
+  TREE_STATIC (decl) = true;
+  TREE_USED (decl) = true;
+
+  DECL_INITIAL (decl) =				  /* XXX: off-by-one? */
+    build_string_literal (len + 1, contents);
+
+  DECL_ARTIFICIAL (decl) = true;
+
+  return decl;
+}
+
+/* Return a VAR_DECL for a string variable containing the contents of FILE,
+   which is looked for in each of the directories listed in SEARCH_PATH.  If
+   FILE could not be found, return NULL_TREE.  */
+
+static tree
+build_variable_from_file_contents (location_t loc,
+				   const char *name_seed,
+				   const char *file,
+				   const_tree search_path)
+{
+  gcc_assert (search_path != NULL_TREE
+	      && TREE_CODE (search_path) == TREE_LIST);
+
+  int err, dir_fd;
+  struct stat st;
+  const_tree dirs;
+  tree var = NULL_TREE;
+
+  /* Look for FILE in each directory in SEARCH_PATH, and pick the first one
+     that matches.  */
+  for (err = ENOENT, dir_fd = -1, dirs = search_path;
+       (err != 0 || err == ENOENT) && dirs != NULL_TREE;
+       dirs = TREE_CHAIN (dirs))
+    {
+      gcc_assert (TREE_VALUE (dirs) != NULL_TREE
+		  && TREE_CODE (TREE_VALUE (dirs)) == STRING_CST);
+
+      dir_fd = open (TREE_STRING_POINTER (TREE_VALUE (dirs)),
+		     O_DIRECTORY | O_RDONLY);
+      if (dir_fd < 0)
+	err = ENOENT;
+      else
+	{
+	  err = fstatat (dir_fd, file, &st, 0);
+	  if (err != 0)
+	    close (dir_fd);
+	  else
+	    /* Leave DIRS unchanged so it can be referred to in diagnostics
+	       below.  */
+	    break;
+	}
+    }
+
+  if (err != 0 || dir_fd < 0)
+    error_at (loc, "failed to access %qs: %m", file);
+  else if (st.st_size == 0)
+    {
+      error_at (loc, "source file %qs is empty", file);
+      close (dir_fd);
+    }
+  else
+    {
+      if (verbose_output_p)
+	inform (loc, "found file %qs in %qs",
+		file, TREE_STRING_POINTER (TREE_VALUE (dirs)));
+
+      int fd;
+
+      fd = openat (dir_fd, file, O_RDONLY);
+      close (dir_fd);
+
+      if (fd < 0)
+	error_at (loc, "failed to open %qs: %m", file);
+      else
+	{
+	  void *contents;
+
+	  contents = mmap (NULL, st.st_size, PROT_READ, MAP_SHARED, fd, 0);
+	  if (contents == NULL)
+	    error_at (loc, "failed to map contents of %qs: %m", file);
+	  else
+	    {
+	      var = build_string_variable (loc, name_seed,
+					   (char *) contents, st.st_size);
+	      pushdecl (var);
+	      munmap (contents, st.st_size);
+	    }
+
+	  close (fd);
+	}
+    }
+
+  return var;
+}
+
+/* Return an expression that, given the OpenCL error code in ERROR_VAR,
+   returns a string.  */
+
+static tree
+build_opencl_error_string (tree error_var)
+{
+  static tree clstrerror_fn;
+  LOOKUP_STARPU_FUNCTION (clstrerror_fn, "starpu_opencl_error_string");
+
+  return build_call_expr (clstrerror_fn, 1, error_var);
+}
+
+/* Return an error-checking `clSetKernelArg' call for argument ARG, at
+   index IDX, of KERNEL.  */
+
+static tree
+build_opencl_set_kernel_arg_call (location_t loc, tree fn,
+				  tree kernel, unsigned int idx,
+				  tree arg)
+{
+  gcc_assert (TREE_CODE (fn) == FUNCTION_DECL
+	      && TREE_TYPE (kernel) == opencl_kernel_type ());
+
+  static tree setkernarg_fn;
+  LOOKUP_STARPU_FUNCTION (setkernarg_fn, "clSetKernelArg");
+
+  tree call = build_call_expr (setkernarg_fn, 4, kernel,
+			       build_int_cst (integer_type_node, idx),
+			       size_in_bytes (TREE_TYPE (arg)),
+			       build_addr (arg, fn));
+  tree error_var = build_decl (loc, VAR_DECL,
+			       create_tmp_var_name ("setkernelarg_error"),
+			       integer_type_node);
+  DECL_ARTIFICIAL (error_var) = true;
+  DECL_CONTEXT (error_var) = fn;
+
+  tree assignment = build2 (INIT_EXPR, TREE_TYPE (error_var),
+			    error_var, call);
+
+  /* Build `if (ERROR_VAR != 0) error ();'.  */
+  tree cond;
+  cond = build3 (COND_EXPR, void_type_node,
+		 build2 (NE_EXPR, boolean_type_node,
+			 error_var, integer_zero_node),
+		 build_error_statements (loc, error_var,
+					 build_opencl_error_string,
+					 "failed to set OpenCL kernel "
+					 "argument %d", idx),
+		 NULL_TREE);
+
+  tree stmts = NULL_TREE;
+  append_to_statement_list (assignment, &stmts);
+  append_to_statement_list (cond, &stmts);
+
+  return build4 (TARGET_EXPR, void_type_node, error_var,
+		 stmts, NULL_TREE, NULL_TREE);
+}
+
+/* Return the sequence of `clSetKernelArg' calls for KERNEL.  */
+
+static tree
+build_opencl_set_kernel_arg_calls (location_t loc, tree task_impl,
+				   tree kernel)
+{
+  gcc_assert (task_implementation_p (task_impl));
+
+  size_t n;
+  tree arg, stmts = NULL_TREE;
+
+  for (arg = DECL_ARGUMENTS (task_impl), n = 0;
+       arg != NULL_TREE;
+       arg = TREE_CHAIN (arg), n++)
+    {
+      tree call = build_opencl_set_kernel_arg_call (loc, task_impl,
+						    kernel, n, arg);
+      append_to_statement_list (call, &stmts);
+    }
+
+  return stmts;
+}
+
+/* Define a body for TASK_IMPL that loads OpenCL source from FILE and calls
+   KERNEL.  */
+
+static void
+define_opencl_task_implementation (location_t loc, tree task_impl,
+				   const char *file, const_tree kernel,
+				   tree groupsize)
+{
+  gcc_assert (task_implementation_p (task_impl)
+	      && task_implementation_where (task_impl) == STARPU_OPENCL);
+  gcc_assert (TREE_CODE (kernel) == STRING_CST);
+  gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (groupsize)));
+
+  local_define (tree, local_var, (tree type))
+  {
+    tree var = build_decl (loc, VAR_DECL,
+			   create_tmp_var_name ("opencl_var"),
+			   type);
+    DECL_ARTIFICIAL (var) = true;
+    DECL_CONTEXT (var) = task_impl;
+    return var;
+  };
+
+  if (!verbose_output_p)
+    /* No further warnings for this node.  */
+    TREE_NO_WARNING (task_impl) = true;
+
+  static tree load_fn, load_kern_fn, enqueue_kern_fn, wid_fn, devid_fn, clfinish_fn,
+    collect_stats_fn, release_ev_fn;
+
+  if (load_fn == NULL_TREE)
+    {
+      load_fn =
+	lookup_name (get_identifier ("starpu_opencl_load_opencl_from_string"));
+      if (load_fn == NULL_TREE)
+	{
+	  inform (loc, "no OpenCL support, task implementation %qE "
+		  "not generated", DECL_NAME (task_impl));
+	  return;
+	}
+    }
+
+  LOOKUP_STARPU_FUNCTION (load_kern_fn, "starpu_opencl_load_kernel");
+  LOOKUP_STARPU_FUNCTION (wid_fn, "starpu_worker_get_id");
+  LOOKUP_STARPU_FUNCTION (devid_fn, "starpu_worker_get_devid");
+  LOOKUP_STARPU_FUNCTION (enqueue_kern_fn, "clEnqueueNDRangeKernel");
+  LOOKUP_STARPU_FUNCTION (clfinish_fn, "clFinish");
+  LOOKUP_STARPU_FUNCTION (collect_stats_fn, "starpu_opencl_collect_stats");
+  LOOKUP_STARPU_FUNCTION (release_ev_fn, "clReleaseEvent");
+
+  if (verbose_output_p)
+    inform (loc, "defining %qE, with OpenCL kernel %qs from file %qs",
+	    DECL_NAME (task_impl), TREE_STRING_POINTER (kernel), file);
+
+  tree source_var;
+  source_var = build_variable_from_file_contents (loc, "opencl_source",
+						  file, opencl_include_dirs);
+  if (source_var != NULL_TREE)
+    {
+      /* Give TASK_IMPL an actual argument list.  */
+      DECL_ARGUMENTS (task_impl) = build_function_arguments (task_impl);
+
+      tree prog_var, prog_loaded_var;
+
+      /* Global variable to hold the `starpu_opencl_program' object.  */
+
+      prog_var = build_decl (loc, VAR_DECL,
+			     create_tmp_var_name ("opencl_program"),
+			     opencl_program_type ());
+      TREE_PUBLIC (prog_var) = false;
+      TREE_STATIC (prog_var) = true;
+      TREE_USED (prog_var) = true;
+      DECL_ARTIFICIAL (prog_var) = true;
+      pushdecl (prog_var);
+
+      /* Global variable indicating whether the program has already been
+	 loaded.  */
+
+      prog_loaded_var = build_decl (loc, VAR_DECL,
+				    create_tmp_var_name ("opencl_prog_loaded"),
+				    boolean_type_node);
+      TREE_PUBLIC (prog_loaded_var) = false;
+      TREE_STATIC (prog_loaded_var) = true;
+      TREE_USED (prog_loaded_var) = true;
+      DECL_ARTIFICIAL (prog_loaded_var) = true;
+      DECL_INITIAL (prog_loaded_var) = build_zero_cst (boolean_type_node);
+      pushdecl (prog_loaded_var);
+
+      /* Build `starpu_opencl_load_opencl_from_string (SOURCE_VAR,
+	                                               &PROG_VAR, "")'.  */
+      tree load = build_call_expr (load_fn, 3, source_var,
+				   build_addr (prog_var, task_impl),
+				   build_string_literal (1, ""));
+
+      tree load_stmts = NULL_TREE;
+      append_to_statement_list (load, &load_stmts);
+      append_to_statement_list (build2 (MODIFY_EXPR, boolean_type_node,
+					prog_loaded_var,
+					build_int_cst (boolean_type_node, 1)),
+				&load_stmts);
+
+      /* Build `if (!PROG_LOADED_VAR) { ...; PROG_LOADED_VAR = true; }'.  */
+
+      tree load_cond = build3 (COND_EXPR, void_type_node,
+			       prog_loaded_var,
+			       NULL_TREE,
+			       load_stmts);
+
+      /* Local variables.  */
+      tree kernel_var, queue_var, event_var, group_size_var, ngroups_var,
+	error_var;
+
+      kernel_var = local_var (opencl_kernel_type ());
+      queue_var = local_var (opencl_command_queue_type ());
+      event_var = local_var (opencl_event_type ());
+      group_size_var = local_var (size_type_node);
+      ngroups_var = local_var (size_type_node);
+      error_var = local_var (integer_type_node);
+
+      /* Build `starpu_opencl_load_kernel (...)'.
+         TODO: Check return value.  */
+      tree devid =
+	build_call_expr (devid_fn, 1, build_call_expr (wid_fn, 0));
+
+      tree load_kern = build_call_expr (load_kern_fn, 5,
+					build_addr (kernel_var, task_impl),
+					build_addr (queue_var, task_impl),
+					build_addr (prog_var, task_impl),
+					build_string_literal
+					(TREE_STRING_LENGTH (kernel) + 1,
+					 TREE_STRING_POINTER (kernel)),
+					devid);
+
+      tree enqueue_kern =
+	build_call_expr (enqueue_kern_fn, 9,
+			 queue_var, kernel_var,
+			 build_int_cst (integer_type_node, 1),
+			 null_pointer_node,
+			 build_addr (group_size_var, task_impl),
+			 build_addr (ngroups_var, task_impl),
+			 integer_zero_node,
+			 null_pointer_node,
+			 build_addr (event_var, task_impl));
+      tree enqueue_err =
+	build2 (INIT_EXPR, TREE_TYPE (error_var), error_var, enqueue_kern);
+
+      tree enqueue_cond =
+	build3 (COND_EXPR, void_type_node,
+		build2 (NE_EXPR, boolean_type_node,
+			error_var, integer_zero_node),
+		build_error_statements (loc, error_var,
+					build_opencl_error_string,
+					"failed to enqueue kernel"),
+		NULL_TREE);
+
+      tree clfinish =
+	build_call_expr (clfinish_fn, 1, queue_var);
+
+      tree collect_stats =
+	build_call_expr (collect_stats_fn, 1, event_var);
+
+      tree release_ev =
+	build_call_expr (release_ev_fn, 1, event_var);
+
+      tree enqueue_stmts = NULL_TREE;
+      append_to_statement_list (enqueue_err, &enqueue_stmts);
+      append_to_statement_list (enqueue_cond, &enqueue_stmts);
+
+
+      /* TODO: Build `clFinish', `clReleaseEvent', & co.  */
+      /* Put it all together.  */
+      tree stmts = NULL_TREE;
+      append_to_statement_list (load_cond, &stmts);
+      append_to_statement_list (load_kern, &stmts);
+      append_to_statement_list (build_opencl_set_kernel_arg_calls (loc,
+								   task_impl,
+								   kernel_var),
+				&stmts);
+
+      /* TODO: Support user-provided values.  */
+      append_to_statement_list (build2 (INIT_EXPR, TREE_TYPE (group_size_var),
+					group_size_var,
+					fold_convert (TREE_TYPE (group_size_var),
+						      groupsize)),
+				&stmts);
+      append_to_statement_list (build2 (INIT_EXPR, TREE_TYPE (ngroups_var),
+					ngroups_var,
+					build_int_cst (TREE_TYPE (ngroups_var),
+						       1)),
+				&stmts);
+      append_to_statement_list (build4 (TARGET_EXPR, void_type_node,
+					error_var, enqueue_stmts,
+					NULL_TREE, NULL_TREE),
+				&stmts);
+      append_to_statement_list (clfinish, &stmts);
+      append_to_statement_list (collect_stats, &stmts);
+      append_to_statement_list (release_ev, &stmts);
+
+      /* Bind the local vars.  */
+      tree vars = chain_trees (kernel_var, queue_var, event_var,
+			       group_size_var, ngroups_var, NULL_TREE);
+      tree bind = build3 (BIND_EXPR, void_type_node, vars, stmts,
+			  build_block (vars, NULL_TREE, task_impl, NULL_TREE));
+
+      TREE_USED (task_impl) = true;
+      TREE_STATIC (task_impl) = true;
+      DECL_EXTERNAL (task_impl) = false;
+      DECL_ARTIFICIAL (task_impl) = true;
+      DECL_SAVED_TREE (task_impl) = bind;
+      DECL_INITIAL (task_impl) = BIND_EXPR_BLOCK (bind);
+      DECL_RESULT (task_impl) =
+	build_decl (loc, RESULT_DECL, NULL_TREE, void_type_node);
+
+      /* Compile TASK_IMPL.  */
+      rest_of_decl_compilation (task_impl, true, 0);
+      allocate_struct_function (task_impl, false);
+      cgraph_finalize_function (task_impl, false);
+      cgraph_mark_needed_node (cgraph_get_node (task_impl));
+
+      /* Generate a wrapper for TASK_IMPL, and possibly the body of its task.
+	 This needs to be done explicitly here, because otherwise
+	 `handle_pre_genericize' would never see TASK_IMPL's task.  */
+      tree task = task_implementation_task (task_impl);
+      if (!TREE_STATIC (task))
+	{
+	  declare_codelet (task);
+	  define_task (task);
+
+	  /* Compile TASK's body.  */
+	  rest_of_decl_compilation (task, true, 0);
+	  allocate_struct_function (task, false);
+	  cgraph_finalize_function (task, false);
+	  cgraph_mark_needed_node (cgraph_get_node (task));
+	}
+    }
+  else
+    DECL_SAVED_TREE (task_impl) = error_mark_node;
+
+  return;
+}
+
+/* Handle the `opencl' pragma, which defines an OpenCL task
+   implementation.  */
+
+void
+handle_pragma_opencl (struct cpp_reader *reader)
+{
+  tree args;
+  location_t loc;
+
+  loc = cpp_peek_token (reader, 0)->src_loc;
+
+  if (current_function_decl != NULL_TREE)
+    {
+      error_at (loc, "%<starpu opencl%> pragma can only be used "
+		"at the top-level");
+      return;
+    }
+
+  args = read_pragma_expressions ("opencl", loc);
+  if (args == NULL_TREE)
+    return;
+
+  /* TODO: Add "number of groups" arguments.  */
+  if (list_length (args) < 4)
+    {
+      error_at (loc, "wrong number of arguments for %<starpu opencl%> pragma");
+      return;
+    }
+
+  if (task_implementation_p (TREE_VALUE (args)))
+    {
+      tree task_impl = TREE_VALUE (args);
+      if (task_implementation_where (task_impl) == STARPU_OPENCL)
+  	{
+  	  args = TREE_CHAIN (args);
+  	  if (TREE_CODE (TREE_VALUE (args)) == STRING_CST)
+  	    {
+  	      tree file = TREE_VALUE (args);
+  	      args = TREE_CHAIN (args);
+  	      if (TREE_CODE (TREE_VALUE (args)) == STRING_CST)
+  		{
+  		  tree kernel = TREE_VALUE (args);
+		  args = TREE_CHAIN (args);
+
+		  if (TREE_TYPE (TREE_VALUE (args)) != NULL_TREE &&
+		      INTEGRAL_TYPE_P (TREE_TYPE (TREE_VALUE (args))))
+		    {
+		      tree groupsize = TREE_VALUE (args);
+		      if (TREE_CHAIN (args) == NULL_TREE)
+			define_opencl_task_implementation (loc, task_impl,
+							   TREE_STRING_POINTER (file),
+							   kernel, groupsize);
+		      else
+			error_at (loc, "junk after %<starpu opencl%> pragma");
+		    }
+		  else
+		    error_at (loc, "%<groupsize%> argument must be an integral type");
+  		}
+  	      else
+  		error_at (loc, "%<kernel%> argument must be a string constant");
+	    }
+	  else
+	    error_at (loc, "%<file%> argument must be a string constant");
+	}
+      else
+	error_at (loc, "%qE is not an OpenCL task implementation",
+		  DECL_NAME (task_impl));
+    }
+  else
+    error_at (loc, "%qE is not a task implementation", TREE_VALUE (args));
+}
+
+/* Diagnose use of C types that are either nonexistent or different in
+   OpenCL.  */
+
+void
+validate_opencl_argument_type (location_t loc, const_tree type)
+{
+  /* When TYPE is a pointer type, get to the base element type.  */
+  for (; POINTER_TYPE_P (type); type = TREE_TYPE (type));
+
+  if (!RECORD_OR_UNION_TYPE_P (type) && !VOID_TYPE_P (type))
+    {
+      tree decl = TYPE_NAME (type);
+
+      if (DECL_P (decl))
+	{
+	  static const struct { const char *c; const char *cl; }
+	  type_map[] =
+	    {
+	      /* Scalar types defined in OpenCL 1.2.  See
+		 <http://www.khronos.org/files/opencl-1-2-quick-reference-card.pdf>.  */
+	      { "char", "cl_char" },
+	      { "signed char", "cl_char" },
+	      { "unsigned char", "cl_uchar" },
+	      { "uchar", "cl_uchar" },
+	      { "short int", "cl_short" },
+	      { "unsigned short", "cl_ushort" },
+	      { "int", "cl_int" },
+	      { "unsigned int", "cl_uint" },
+	      { "uint", "cl_uint" },
+	      { "long int", "cl_long" },
+	      { "long unsigned int", "cl_ulong" },
+	      { "ulong", "cl_ulong" },
+	      { "float", "cl_float" },
+	      { "double", "cl_double" },
+	      { NULL, NULL }
+	    };
+
+	  const char *c_name = IDENTIFIER_POINTER (DECL_NAME (decl));
+	  const char *cl_name =
+	    ({
+	      size_t i;
+	      for (i = 0; type_map[i].c != NULL; i++)
+		{
+		  if (strcmp (type_map[i].c, c_name) == 0)
+		    break;
+		}
+	      type_map[i].cl;
+	    });
+
+	  if (cl_name != NULL)
+	    {
+	      tree cl_type = lookup_name (get_identifier (cl_name));
+
+	      if (cl_type != NULL_TREE)
+		{
+		  if (DECL_P (cl_type))
+		    cl_type = TREE_TYPE (cl_type);
+
+		  if (!lang_hooks.types_compatible_p ((tree) type, cl_type))
+		    {
+		      tree st, sclt;
+
+		      st = c_common_signed_type ((tree) type);
+		      sclt = c_common_signed_type (cl_type);
+
+		      if (st == sclt)
+			warning_at (loc, 0, "C type %qE differs in signedness "
+				    "from the same-named OpenCL type",
+				    DECL_NAME (decl));
+		      else
+			/* TYPE should be avoided because the it differs from
+			   CL_TYPE, and thus cannot be used safely in
+			   `clSetKernelArg'.  */
+			warning_at (loc, 0, "C type %qE differs from the "
+				    "same-named OpenCL type",
+				    DECL_NAME (decl));
+		    }
+		}
+
+	      /* Otherwise we can't conclude.  It could be that <CL/cl.h>
+		 wasn't included in the program, for instance.  */
+	    }
+	  else
+	    /* Recommend against use of `size_t', etc.  */
+	    warning_at (loc, 0, "%qE does not correspond to a known "
+			"OpenCL type", DECL_NAME (decl));
+	}
+    }
+}

gcc-plugin/src/tasks.c

@@ -0,0 +1,674 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Inria
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#include <starpu-gcc/config.h>
+
+/* We must include starpu.h here, otherwise gcc will complain about a poisoned
+   malloc in xmmintrin.h.  */
+#include <starpu.h>
+
+#include <gcc-plugin.h>
+#include <plugin-version.h>
+
+#include <plugin.h>
+#include <cpplib.h>
+#include <tree.h>
+#include <tree-iterator.h>
+#include <gimple.h>
+
+#ifdef HAVE_C_FAMILY_C_COMMON_H
+# include <c-family/c-common.h>
+#elif HAVE_C_COMMON_H
+# include <c-common.h>
+#endif
+
+#include <diagnostic.h>
+
+#include <starpu-gcc/tasks.h>
+#include <starpu-gcc/utils.h>
+#include <starpu-gcc/opencl.h>
+
+
+/* Task-related functions.  */
+
+/* Name of public attributes.  */
+const char task_attribute_name[] = "task";
+const char task_implementation_attribute_name[] = "task_implementation";
+const char output_attribute_name[] = "output";
+
+/* Names of attributes used internally.  */
+static const char task_codelet_attribute_name[] = ".codelet";
+const char task_implementation_list_attribute_name[] =
+  ".task_implementation_list";
+const char task_implementation_wrapper_attribute_name[] =
+  ".task_implementation_wrapper";
+
+/* Names of data structures defined in <starpu.h>.  */
+static const char codelet_struct_tag[] = "starpu_codelet";
+
+
+/* Return true if DECL is a task.  */
+
+bool
+task_p (const_tree decl)
+{
+  return (TREE_CODE (decl) == FUNCTION_DECL &&
+	  lookup_attribute (task_attribute_name,
+			    DECL_ATTRIBUTES (decl)) != NULL_TREE);
+}
+
+/* Return true if DECL is a task implementation.  */
+
+bool
+task_implementation_p (const_tree decl)
+{
+  return (TREE_CODE (decl) == FUNCTION_DECL &&
+	  lookup_attribute (task_implementation_attribute_name,
+			    DECL_ATTRIBUTES (decl)) != NULL_TREE);
+}
+
+/* Return a value indicating where TASK_IMPL should execute (`STARPU_CPU',
+   `STARPU_CUDA', etc.).  */
+
+int
+task_implementation_where (const_tree task_impl)
+{
+  tree impl_attr, args, where;
+
+  gcc_assert (TREE_CODE (task_impl) == FUNCTION_DECL);
+
+  impl_attr = lookup_attribute (task_implementation_attribute_name,
+				DECL_ATTRIBUTES (task_impl));
+  gcc_assert (impl_attr != NULL_TREE);
+
+  args = TREE_VALUE (impl_attr);
+  where = TREE_VALUE (args);
+
+  return task_implementation_target_to_int (where);
+}
+
+/* Return the StarPU integer constant corresponding to string TARGET.  */
+
+int
+task_implementation_target_to_int (const_tree target)
+{
+  gcc_assert (TREE_CODE (target) == STRING_CST);
+
+  int where_int;
+
+  if (!strncmp (TREE_STRING_POINTER (target), "cpu",
+		TREE_STRING_LENGTH (target)))
+    where_int = STARPU_CPU;
+  else if (!strncmp (TREE_STRING_POINTER (target), "opencl",
+		     TREE_STRING_LENGTH (target)))
+    where_int = STARPU_OPENCL;
+  else if (!strncmp (TREE_STRING_POINTER (target), "cuda",
+		     TREE_STRING_LENGTH (target)))
+    where_int = STARPU_CUDA;
+  else if (!strncmp (TREE_STRING_POINTER (target), "gordon",
+		     TREE_STRING_LENGTH (target)))
+    where_int = STARPU_GORDON;
+  else
+    where_int = 0;
+
+  return where_int;
+}
+
+/* Return the task implemented by TASK_IMPL.  */
+
+tree
+task_implementation_task (const_tree task_impl)
+{
+  tree impl_attr, args, task;
+
+  gcc_assert (TREE_CODE (task_impl) == FUNCTION_DECL);
+
+  impl_attr = lookup_attribute (task_implementation_attribute_name,
+				DECL_ATTRIBUTES (task_impl));
+  gcc_assert (impl_attr != NULL_TREE);
+
+  args = TREE_VALUE (impl_attr);
+
+  task = TREE_VALUE (TREE_CHAIN (args));
+  if (task_implementation_p (task))
+    /* TASK is an implicit CPU task implementation, so return its real
+       task.  */
+    return task_implementation_task (task);
+
+  return task;
+}
+
+/* Return the declaration of the `struct starpu_codelet' variable associated with
+   TASK_DECL.  */
+
+tree
+task_codelet_declaration (const_tree task_decl)
+{
+  tree cl_attr;
+
+  cl_attr = lookup_attribute (task_codelet_attribute_name,
+			      DECL_ATTRIBUTES (task_decl));
+  gcc_assert (cl_attr != NULL_TREE);
+
+  return TREE_VALUE (cl_attr);
+}
+
+/* Return the list of implementations of TASK_DECL.  */
+
+tree
+task_implementation_list (const_tree task_decl)
+{
+  tree attr;
+
+  attr = lookup_attribute (task_implementation_list_attribute_name,
+			   DECL_ATTRIBUTES (task_decl));
+  return TREE_VALUE (attr);
+}
+
+/* Return the list of pointer parameter types of TASK_DECL.  */
+
+tree
+task_pointer_parameter_types (const_tree task_decl)
+{
+  return filter (pointer_type_p, TYPE_ARG_TYPES (TREE_TYPE (task_decl)));
+}
+
+/* Return a bitwise-or of the supported targets of TASK_DECL.  */
+
+int
+task_where (const_tree task_decl)
+{
+  gcc_assert (task_p (task_decl));
+
+  int where;
+  const_tree impl;
+
+  for (impl = task_implementation_list (task_decl), where = 0;
+       impl != NULL_TREE;
+       impl = TREE_CHAIN (impl))
+    where |= task_implementation_where (TREE_VALUE (impl));
+
+  return where;
+}
+
+/* Return the FUNCTION_DECL of the wrapper generated for TASK_IMPL.  */
+
+tree
+task_implementation_wrapper (const_tree task_impl)
+{
+  tree attr;
+
+  gcc_assert (TREE_CODE (task_impl) == FUNCTION_DECL);
+
+  attr = lookup_attribute (task_implementation_wrapper_attribute_name,
+			   DECL_ATTRIBUTES (task_impl));
+  gcc_assert (attr != NULL_TREE);
+
+  return TREE_VALUE (attr);
+}
+
+tree
+codelet_type (void)
+{
+  /* XXX: Hack to allow the type declaration to be accessible at lower
+     time.  */
+  static tree type_decl = NULL_TREE;
+
+  if (type_decl == NULL_TREE)
+    /* Lookup the `struct starpu_codelet' struct type.  This should succeed since
+       we push <starpu.h> early on.  */
+    type_decl = type_decl_for_struct_tag (codelet_struct_tag);
+
+  return TREE_TYPE (type_decl);
+}
+
+/* Return the access mode for POINTER, a PARM_DECL of a task.  */
+
+enum starpu_access_mode
+access_mode (const_tree type)
+{
+  gcc_assert (POINTER_TYPE_P (type));
+
+  /* If TYPE points to a const-qualified type, then mark the data as
+     read-only; if is has the `output' attribute, then mark it as write-only;
+     otherwise default to read-write.  */
+  return ((TYPE_QUALS (TREE_TYPE (type)) & TYPE_QUAL_CONST)
+	  ? STARPU_R
+	  : (output_type_p (type) ? STARPU_W : STARPU_RW));
+}
+
+/* Return true if TYPE is `output'-qualified.  */
+
+bool
+output_type_p (const_tree type)
+{
+  return (lookup_attribute (output_attribute_name,
+			    TYPE_ATTRIBUTES (type)) != NULL_TREE);
+}
+
+
+/* Code generation.  */
+
+/* Turn FN into a task, and push its associated codelet declaration.  */
+
+void
+taskify_function (tree fn)
+{
+  gcc_assert (TREE_CODE (fn) == FUNCTION_DECL);
+
+  /* Add a `task' attribute and an empty `task_implementation_list'
+     attribute.  */
+  DECL_ATTRIBUTES (fn) =
+    tree_cons (get_identifier (task_implementation_list_attribute_name),
+	       NULL_TREE,
+	       tree_cons (get_identifier (task_attribute_name), NULL_TREE,
+			  DECL_ATTRIBUTES (fn)));
+
+  /* Push a declaration for the corresponding `struct starpu_codelet' object and
+     add it as an attribute of FN.  */
+  tree cl = build_codelet_declaration (fn);
+  DECL_ATTRIBUTES (fn) =
+    tree_cons (get_identifier (task_codelet_attribute_name), cl,
+	       DECL_ATTRIBUTES (fn));
+
+  pushdecl (cl);
+}
+
+
+/* Return a NODE_IDENTIFIER for the variable holding the `struct starpu_codelet'
+   structure associated with TASK_DECL.  */
+
+tree
+build_codelet_identifier (tree task_decl)
+{
+  static const char suffix[] = ".codelet";
+
+  tree id;
+  char *cl_name;
+  const char *task_name;
+
+  id = DECL_NAME (task_decl);
+  task_name = IDENTIFIER_POINTER (id);
+
+  cl_name = (char *) alloca (IDENTIFIER_LENGTH (id) + strlen (suffix) + 1);
+  memcpy (cl_name, task_name, IDENTIFIER_LENGTH (id));
+  strcpy (&cl_name[IDENTIFIER_LENGTH (id)], suffix);
+
+  return get_identifier (cl_name);
+}
+
+/* Return a VAR_DECL that declares a `struct starpu_codelet' structure for
+   TASK_DECL.  */
+
+tree
+build_codelet_declaration (tree task_decl)
+{
+  tree name, cl_decl;
+
+  name = build_codelet_identifier (task_decl);
+
+  cl_decl = build_decl (DECL_SOURCE_LOCATION (task_decl),
+			VAR_DECL, name,
+			/* c_build_qualified_type (type, TYPE_QUAL_CONST) */
+			codelet_type ());
+
+  DECL_ARTIFICIAL (cl_decl) = true;
+  TREE_PUBLIC (cl_decl) = TREE_PUBLIC (task_decl);
+  TREE_STATIC (cl_decl) = false;
+  TREE_USED (cl_decl) = true;
+  DECL_EXTERNAL (cl_decl) = true;
+  DECL_CONTEXT (cl_decl) = NULL_TREE;
+
+  return cl_decl;
+}
+
+/* Return a `struct starpu_codelet' initializer for TASK_DECL.  */
+
+tree
+build_codelet_initializer (tree task_decl)
+{
+  tree fields;
+
+  fields = TYPE_FIELDS (codelet_type ());
+  gcc_assert (TREE_CODE (fields) == FIELD_DECL);
+
+  local_define (tree, lookup_field, (const char *name))
+  {
+    tree fdecl, fname;
+
+    fname = get_identifier (name);
+    for (fdecl = fields;
+	 fdecl != NULL_TREE;
+	 fdecl = TREE_CHAIN (fdecl))
+      {
+	if (DECL_NAME (fdecl) == fname)
+	  return fdecl;
+      }
+
+    /* Field NAME wasn't found.  */
+    gcc_assert (false);
+  };
+
+  local_define (tree, field_initializer, (const char *name, tree value))
+  {
+    tree field, init;
+
+    field = lookup_field (name);
+    init = make_node (TREE_LIST);
+    TREE_PURPOSE (init) = field;
+    TREE_CHAIN (init) = NULL_TREE;
+
+    if (TREE_CODE (TREE_TYPE (value)) != ARRAY_TYPE)
+      TREE_VALUE (init) = fold_convert (TREE_TYPE (field), value);
+    else
+      TREE_VALUE (init) = value;
+
+    return init;
+  };
+
+  local_define (tree, codelet_name, ())
+  {
+    const char *name = IDENTIFIER_POINTER (DECL_NAME (task_decl));
+    return build_string_literal (strlen (name) + 1, name);
+  };
+
+  local_define (tree, where_init, (tree impls))
+  {
+    tree impl;
+    int where_int = 0;
+
+    for (impl = impls;
+	 impl != NULL_TREE;
+	 impl = TREE_CHAIN (impl))
+      {
+	tree impl_decl;
+
+	impl_decl = TREE_VALUE (impl);
+	gcc_assert (TREE_CODE (impl_decl) == FUNCTION_DECL);
+
+	if (verbose_output_p)
+	  /* List the implementations of TASK_DECL.  */
+	  inform (DECL_SOURCE_LOCATION (impl_decl),
+		  "   %qE", DECL_NAME (impl_decl));
+
+	where_int |= task_implementation_where (impl_decl);
+      }
+
+    return build_int_cstu (integer_type_node, where_int);
+  };
+
+  local_define (tree, implementation_pointers, (tree impls, int where))
+  {
+    size_t len;
+    tree impl, pointers;
+
+    for (impl = impls, pointers = NULL_TREE, len = 0;
+	 impl != NULL_TREE;
+	 impl = TREE_CHAIN (impl))
+      {
+	tree impl_decl;
+
+	impl_decl = TREE_VALUE (impl);
+	if (task_implementation_where (impl_decl) == where)
+	  {
+	    /* Return a pointer to the wrapper of IMPL_DECL.  */
+	    tree addr = build_addr (task_implementation_wrapper (impl_decl),
+				    NULL_TREE);
+	    pointers = tree_cons (size_int (len), addr, pointers);
+	    len++;
+
+	    if (len > STARPU_MAXIMPLEMENTATIONS)
+	      error_at (DECL_SOURCE_LOCATION (impl_decl),
+			"maximum number of per-target task implementations "
+			"exceeded");
+	  }
+      }
+
+    /* POINTERS must be null-terminated.  */
+    pointers = tree_cons (size_int (len), build_zero_cst (ptr_type_node),
+			  pointers);
+    len++;
+
+    /* Return an array initializer.  */
+    tree index_type = build_index_type (size_int (list_length (pointers)));
+
+    return build_constructor_from_list (build_array_type (ptr_type_node,
+							  index_type),
+					nreverse (pointers));
+  };
+
+  local_define (tree, pointer_arg_count, (void))
+  {
+    size_t len;
+
+    len = list_length (task_pointer_parameter_types (task_decl));
+    return build_int_cstu (integer_type_node, len);
+  };
+
+  local_define (tree, access_mode_array, (void))
+  {
+    const_tree type;
+    tree modes;
+    size_t index;
+
+    for (type = task_pointer_parameter_types (task_decl),
+	   modes = NULL_TREE, index = 0;
+	 type != NULL_TREE && index < STARPU_NMAXBUFS;
+	 type = TREE_CHAIN (type), index++)
+      {
+	tree value = build_int_cst (integer_type_node,
+				    access_mode (TREE_VALUE (type)));
+
+	modes = tree_cons (size_int (index), value, modes);
+      }
+
+    tree index_type = build_index_type (size_int (list_length (modes)));
+
+    return build_constructor_from_list (build_array_type (integer_type_node,
+							  index_type),
+					nreverse (modes));
+  };
+
+  if (verbose_output_p)
+    inform (DECL_SOURCE_LOCATION (task_decl),
+	    "implementations for task %qE:", DECL_NAME (task_decl));
+
+  tree impls, inits;
+
+  impls = task_implementation_list (task_decl);
+
+  inits =
+    chain_trees (field_initializer ("name", codelet_name ()),
+		 field_initializer ("where", where_init (impls)),
+		 field_initializer ("nbuffers", pointer_arg_count ()),
+		 field_initializer ("modes", access_mode_array ()),
+		 field_initializer ("cpu_funcs",
+				    implementation_pointers (impls,
+							     STARPU_CPU)),
+		 field_initializer ("opencl_funcs",
+		 		    implementation_pointers (impls,
+							     STARPU_OPENCL)),
+		 field_initializer ("cuda_funcs",
+		 		    implementation_pointers (impls,
+							     STARPU_CUDA)),
+		 NULL_TREE);
+
+  return build_constructor_from_unsorted_list (codelet_type (), inits);
+}
+
+/* Return the VAR_DECL that defines a `struct starpu_codelet' structure for
+   TASK_DECL.  The VAR_DECL is assumed to already exists, so it must not be
+   pushed again.  */
+
+tree
+declare_codelet (tree task_decl)
+{
+  /* Retrieve the declaration of the `struct starpu_codelet' object.  */
+  tree cl_decl;
+  cl_decl = lookup_name (build_codelet_identifier (task_decl));
+  gcc_assert (cl_decl != NULL_TREE && TREE_CODE (cl_decl) == VAR_DECL);
+
+  /* Turn the codelet declaration into a definition.  */
+  TREE_TYPE (cl_decl) = codelet_type ();
+  TREE_PUBLIC (cl_decl) = TREE_PUBLIC (task_decl);
+
+  return cl_decl;
+}
+
+/* Build the body of TASK_DECL, which will call `starpu_insert_task'.  */
+
+void
+define_task (tree task_decl)
+{
+  /* First of all, give TASK_DECL an argument list.  */
+  DECL_ARGUMENTS (task_decl) = build_function_arguments (task_decl);
+
+  VEC(tree, gc) *args = NULL;
+  location_t loc = DECL_SOURCE_LOCATION (task_decl);
+  tree p, params = DECL_ARGUMENTS (task_decl);
+
+  /* The first argument will be a pointer to the codelet.  */
+
+  VEC_safe_push (tree, gc, args,
+		 build_addr (task_codelet_declaration (task_decl),
+			     current_function_decl));
+
+  for (p = params; p != NULL_TREE; p = TREE_CHAIN (p))
+    {
+      gcc_assert (TREE_CODE (p) == PARM_DECL);
+
+      tree type = TREE_TYPE (p);
+
+      if (POINTER_TYPE_P (type))
+	{
+	  /* A pointer: the arguments will be:
+	     `STARPU_RW, ptr' or similar.  */
+
+	  VEC_safe_push (tree, gc, args,
+			 build_int_cst (integer_type_node,
+					access_mode (type)));
+	  VEC_safe_push (tree, gc, args, build_pointer_lookup (p));
+	}
+      else
+	{
+	  /* A scalar: the arguments will be:
+	     `STARPU_VALUE, &scalar, sizeof (scalar)'.  */
+
+	  mark_addressable (p);
+
+	  VEC_safe_push (tree, gc, args,
+			 build_int_cst (integer_type_node, STARPU_VALUE));
+	  VEC_safe_push (tree, gc, args,
+			 build_addr (p, current_function_decl));
+	  VEC_safe_push (tree, gc, args,
+			 size_in_bytes (type));
+	}
+    }
+
+  /* Push the terminating zero.  */
+
+  VEC_safe_push (tree, gc, args,
+		 build_int_cst (integer_type_node, 0));
+
+  /* Introduce a local variable to hold the error code.  */
+
+  tree error_var = build_decl (loc, VAR_DECL,
+  			       create_tmp_var_name (".insert_task_error"),
+  			       integer_type_node);
+  DECL_CONTEXT (error_var) = task_decl;
+  DECL_ARTIFICIAL (error_var) = true;
+
+  /* Build this:
+
+       err = starpu_insert_task (...);
+       if (err != 0)
+         { printf ...; abort (); }
+   */
+
+  static tree insert_task_fn;
+  LOOKUP_STARPU_FUNCTION (insert_task_fn, "starpu_insert_task");
+
+  tree call = build_call_expr_loc_vec (loc, insert_task_fn, args);
+
+  tree assignment = build2 (INIT_EXPR, TREE_TYPE (error_var),
+  			    error_var, call);
+
+  tree name = DECL_NAME (task_decl);
+  tree cond = build3 (COND_EXPR, void_type_node,
+		      build2 (NE_EXPR, boolean_type_node,
+			      error_var, integer_zero_node),
+		      build_error_statements (loc, error_var,
+					      build_starpu_error_string,
+					      "failed to insert task `%s'",
+					      IDENTIFIER_POINTER (name)),
+		      NULL_TREE);
+
+  tree stmts = NULL;
+  append_to_statement_list (assignment, &stmts);
+  append_to_statement_list (cond, &stmts);
+
+  tree bind = build3 (BIND_EXPR, void_type_node, error_var, stmts,
+  		      NULL_TREE);
+
+  /* Put it all together.  */
+
+  DECL_SAVED_TREE (task_decl) = bind;
+  TREE_STATIC (task_decl) = true;
+  DECL_EXTERNAL (task_decl) = false;
+  DECL_ARTIFICIAL (task_decl) = true;
+  DECL_INITIAL (task_decl) =
+    build_block (error_var, NULL_TREE, task_decl, NULL_TREE);
+  DECL_RESULT (task_decl) =
+    build_decl (loc, RESULT_DECL, NULL_TREE, void_type_node);
+  DECL_CONTEXT (DECL_RESULT (task_decl)) = task_decl;
+}
+
+/* Add FN to the list of implementations of TASK_DECL.  */
+
+void
+add_task_implementation (tree task_decl, tree fn, const_tree where)
+{
+  location_t loc;
+  tree attr, impls;
+
+  attr = lookup_attribute (task_implementation_list_attribute_name,
+			   DECL_ATTRIBUTES (task_decl));
+  gcc_assert (attr != NULL_TREE);
+
+  gcc_assert (TREE_CODE (where) == STRING_CST);
+
+  loc = DECL_SOURCE_LOCATION (fn);
+
+  impls = tree_cons (NULL_TREE, fn, TREE_VALUE (attr));
+  TREE_VALUE (attr) = impls;
+
+  TREE_USED (fn) = true;
+
+  /* Check the `where' argument to raise a warning if needed.  */
+  if (task_implementation_target_to_int (where) == 0)
+    warning_at (loc, 0,
+		"unsupported target %E; task implementation won't be used",
+		where);
+  else if (task_implementation_target_to_int (where) == STARPU_OPENCL)
+    {
+      local_define (void, validate, (tree t))
+	{
+	  validate_opencl_argument_type (loc, t);
+	};
+
+      for_each (validate, TYPE_ARG_TYPES (TREE_TYPE (fn)));
+    }
+}

gcc-plugin/src/utils.c

@@ -0,0 +1,432 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Inria
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#include <starpu-gcc/config.h>
+
+/* We must include starpu.h here, otherwise gcc will complain about a poisoned
+   malloc in xmmintrin.h.  */
+#include <starpu.h>
+
+#include <gcc-plugin.h>
+#include <plugin-version.h>
+
+#include <plugin.h>
+#include <cpplib.h>
+#include <tree.h>
+#include <tree-iterator.h>
+#include <gimple.h>
+
+#ifdef HAVE_C_FAMILY_C_COMMON_H
+# include <c-family/c-common.h>
+#elif HAVE_C_COMMON_H
+# include <c-common.h>
+#endif
+
+#include <starpu-gcc/utils.h>
+
+/* Whether to enable verbose output.  */
+bool verbose_output_p = false;
+
+
+/* Various helpers.  */
+
+/* Return a TYPE_DECL for the RECORD_TYPE with tag name TAG.  */
+
+tree
+type_decl_for_struct_tag (const char *tag)
+{
+  tree type_decl = xref_tag (RECORD_TYPE, get_identifier (tag));
+  gcc_assert (type_decl != NULL_TREE
+	      && TREE_CODE (type_decl) == RECORD_TYPE);
+
+  /* `build_decl' expects a TYPE_DECL, so give it what it wants.  */
+
+  type_decl = TYPE_STUB_DECL (type_decl);
+  gcc_assert (type_decl != NULL && TREE_CODE (type_decl) == TYPE_DECL);
+
+  return type_decl;
+}
+
+/* Given ERROR_VAR, an integer variable holding a StarPU error code, return
+   statements that print out the error message returned by
+   BUILD_ERROR_MESSAGE (ERROR_VAR) and abort.  */
+
+tree
+build_error_statements (location_t loc, tree error_var,
+			function_parm (tree, build_error_message, (tree)),
+			const char *fmt, ...)
+{
+  expanded_location xloc = expand_location (loc);
+
+  tree print;
+  char *str, *fmt_long;
+  va_list args;
+
+  va_start (args, fmt);
+
+  /* Build a longer format.  Since FMT itself contains % escapes, this needs
+     to be done in two steps.  */
+
+  vasprintf (&str, fmt, args);
+
+  if (error_var != NULL_TREE)
+    {
+      /* ERROR_VAR is an error code.  */
+      gcc_assert (TREE_CODE (error_var) == VAR_DECL
+		  && TREE_TYPE (error_var) == integer_type_node);
+
+      asprintf (&fmt_long, "%s:%d: error: %s: %%s\n",
+		xloc.file, xloc.line, str);
+
+      print =
+	build_call_expr (builtin_decl_explicit (BUILT_IN_PRINTF), 2,
+			 build_string_literal (strlen (fmt_long) + 1,
+					       fmt_long),
+			 build_error_message (error_var));
+    }
+  else
+    {
+      /* No error code provided.  */
+
+      asprintf (&fmt_long, "%s:%d: error: %s\n",
+		xloc.file, xloc.line, str);
+
+      print =
+	build_call_expr (builtin_decl_explicit (BUILT_IN_PUTS), 1,
+			 build_string_literal (strlen (fmt_long) + 1,
+					       fmt_long));
+    }
+
+  free (fmt_long);
+  free (str);
+  va_end (args);
+
+  tree stmts = NULL;
+  append_to_statement_list (print, &stmts);
+  append_to_statement_list (build_call_expr
+			    (builtin_decl_explicit (BUILT_IN_ABORT), 0),
+			    &stmts);
+
+  return stmts;
+}
+
+/* Return a fresh argument list for FN.  */
+
+tree
+build_function_arguments (tree fn)
+{
+  gcc_assert (TREE_CODE (fn) == FUNCTION_DECL
+	      && DECL_ARGUMENTS (fn) == NULL_TREE);
+
+  local_define (tree, build_argument, (const_tree lst))
+    {
+      tree param, type;
+
+      type = TREE_VALUE (lst);
+      param = build_decl (DECL_SOURCE_LOCATION (fn), PARM_DECL,
+			  create_tmp_var_name ("argument"),
+			  type);
+      DECL_ARG_TYPE (param) = type;
+      DECL_CONTEXT (param) = fn;
+
+      return param;
+    };
+
+  return map (build_argument,
+	      list_remove (void_type_p,
+			   TYPE_ARG_TYPES (TREE_TYPE (fn))));
+}
+
+/* Return true if LST holds the void type.  */
+
+bool
+void_type_p (const_tree lst)
+{
+  gcc_assert (TREE_CODE (lst) == TREE_LIST);
+  return VOID_TYPE_P (TREE_VALUE (lst));
+}
+
+/* Return true if LST holds a pointer type.  */
+
+bool
+pointer_type_p (const_tree lst)
+{
+  gcc_assert (TREE_CODE (lst) == TREE_LIST);
+  return POINTER_TYPE_P (TREE_VALUE (lst));
+}
+
+
+/* C expression parser, possibly with C++ linkage.  */
+
+extern int yyparse (location_t, const char *, tree *);
+extern int yydebug;
+
+/* Parse expressions from the CPP reader for PRAGMA, which is located at LOC.
+   Return a TREE_LIST of C expressions.  */
+
+tree
+read_pragma_expressions (const char *pragma, location_t loc)
+{
+  tree expr = NULL_TREE;
+
+  if (yyparse (loc, pragma, &expr))
+    /* Parse error or memory exhaustion.  */
+    expr = NULL_TREE;
+
+  return expr;
+}
+
+
+/* List and vector utilities, à la SRFI-1.  */
+
+tree
+chain_trees (tree t, ...)
+{
+  va_list args;
+
+  va_start (args, t);
+
+  tree next, prev = t;
+  for (prev = t, next = va_arg (args, tree);
+       next != NULL_TREE;
+       prev = next, next = va_arg (args, tree))
+    TREE_CHAIN (prev) = next;
+
+  va_end (args);
+
+  return t;
+}
+
+tree
+filter (function_parm (bool, pred, (const_tree)), tree t)
+{
+  tree result, lst;
+
+  gcc_assert (TREE_CODE (t) == TREE_LIST);
+
+  result = NULL_TREE;
+  for (lst = t; lst != NULL_TREE; lst = TREE_CHAIN (lst))
+    {
+      if (pred (lst))
+	result = tree_cons (TREE_PURPOSE (lst), TREE_VALUE (lst),
+			    result);
+    }
+
+  return nreverse (result);
+}
+
+tree
+list_remove (function_parm (bool, pred, (const_tree)), tree t)
+{
+  local_define (bool, opposite, (const_tree t))
+  {
+    return !pred (t);
+  };
+
+  return filter (opposite, t);
+}
+
+/* Map FUNC over chain T.  T does not have to be `TREE_LIST'; it can be a
+   chain of arbitrary tree objects.  */
+
+tree
+map (function_parm (tree, func, (const_tree)), tree t)
+{
+  tree result, tail, lst;
+
+  result = tail = NULL_TREE;
+  for (lst = t; lst != NULL_TREE; lst = TREE_CHAIN (lst))
+    {
+      tree r = func (lst);
+      if (tail != NULL_TREE)
+	TREE_CHAIN (tail) = r;
+      else
+	result = r;
+
+      tail = r;
+    }
+
+  return result;
+}
+
+void
+for_each (function_parm (void, func, (tree)), tree t)
+{
+  tree lst;
+
+  gcc_assert (TREE_CODE (t) == TREE_LIST);
+
+  for (lst = t; lst != NULL_TREE; lst = TREE_CHAIN (lst))
+    func (TREE_VALUE (lst));
+}
+
+size_t
+count (function_parm (bool, pred, (const_tree)), const_tree t)
+{
+  size_t result;
+  const_tree lst;
+
+  for (lst = t, result = 0; lst != NULL_TREE; lst = TREE_CHAIN (lst))
+    if (pred (lst))
+      result++;
+
+  return result;
+}
+
+
+/* Useful code backported from GCC 4.6.  */
+
+#if !HAVE_DECL_BUILD_CALL_EXPR_LOC_ARRAY
+
+tree
+build_call_expr_loc_array (location_t loc, tree fndecl, int n, tree *argarray)
+{
+  tree fntype = TREE_TYPE (fndecl);
+  tree fn = build1 (ADDR_EXPR, build_pointer_type (fntype), fndecl);
+
+  return fold_builtin_call_array (loc, TREE_TYPE (fntype), fn, n, argarray);
+}
+
+#endif
+
+#if !HAVE_DECL_BUILD_CALL_EXPR_LOC_VEC
+
+tree
+build_call_expr_loc_vec (location_t loc, tree fndecl, VEC(tree,gc) *vec)
+{
+  return build_call_expr_loc_array (loc, fndecl, VEC_length (tree, vec),
+				    VEC_address (tree, vec));
+}
+
+#endif
+
+#if !HAVE_DECL_BUILD_ZERO_CST
+
+tree
+build_zero_cst (tree type)
+{
+  switch (TREE_CODE (type))
+    {
+    case INTEGER_TYPE: case ENUMERAL_TYPE: case BOOLEAN_TYPE:
+    case POINTER_TYPE: case REFERENCE_TYPE:
+    case OFFSET_TYPE:
+      return build_int_cst (type, 0);
+
+    default:
+      abort ();
+    }
+}
+
+#endif
+
+/* Build a "conversion" from a raw C pointer to its data handle.  The
+   assumption is that the programmer should have already registered the
+   pointer by themselves.  */
+
+tree
+build_pointer_lookup (tree pointer)
+{
+  static tree data_lookup_fn;
+
+  /* Make sure DATA_LOOKUP_FN is valid.  */
+  LOOKUP_STARPU_FUNCTION (data_lookup_fn, "starpu_data_lookup");
+
+  location_t loc;
+
+  if (DECL_P (pointer))
+    loc = DECL_SOURCE_LOCATION (pointer);
+  else
+    loc = UNKNOWN_LOCATION;
+
+  /* Introduce a local variable to hold the handle.  */
+
+  tree result_var = build_decl (loc, VAR_DECL,
+  				create_tmp_var_name (".data_lookup_result"),
+  				ptr_type_node);
+  DECL_CONTEXT (result_var) = current_function_decl;
+  DECL_ARTIFICIAL (result_var) = true;
+  DECL_SOURCE_LOCATION (result_var) = loc;
+
+  tree call = build_call_expr (data_lookup_fn, 1, pointer);
+  tree assignment = build2 (INIT_EXPR, TREE_TYPE (result_var),
+  			    result_var, call);
+
+  /* Build `if (RESULT_VAR == NULL) error ();'.  */
+
+  tree cond = build3 (COND_EXPR, void_type_node,
+		      build2 (EQ_EXPR, boolean_type_node,
+			      result_var, null_pointer_node),
+		      build_error_statements (loc, NULL_TREE,
+					      build_starpu_error_string,
+					      "attempt to use unregistered "
+					      "pointer"),
+		      NULL_TREE);
+
+  tree stmts = NULL;
+  append_to_statement_list (assignment, &stmts);
+  append_to_statement_list (cond, &stmts);
+  append_to_statement_list (result_var, &stmts);
+
+  return build4 (TARGET_EXPR, ptr_type_node, result_var, stmts, NULL_TREE, NULL_TREE);
+}
+
+/* Build an error string for the StarPU return value in ERROR_VAR.  */
+
+tree
+build_starpu_error_string (tree error_var)
+{
+  static tree strerror_fn;
+  LOOKUP_STARPU_FUNCTION (strerror_fn, "strerror");
+
+  tree error_code =
+    build1 (NEGATE_EXPR, TREE_TYPE (error_var), error_var);
+
+  return build_call_expr (strerror_fn, 1, error_code);
+}
+
+/* Like `build_constructor_from_list', but sort VALS according to their
+   offset in struct TYPE.  Inspired by `gnat_build_constructor'.  */
+
+tree
+build_constructor_from_unsorted_list (tree type, tree vals)
+{
+  local_define (int, compare_elmt_bitpos, (const void *rt1, const void *rt2))
+  {
+    const constructor_elt *elmt1 = (constructor_elt *) rt1;
+    const constructor_elt *elmt2 = (constructor_elt *) rt2;
+    const_tree field1 = elmt1->index;
+    const_tree field2 = elmt2->index;
+    int ret
+      = tree_int_cst_compare (bit_position (field1), bit_position (field2));
+
+    return ret ? ret : (int) (DECL_UID (field1) - DECL_UID (field2));
+  };
+
+  tree t;
+  VEC(constructor_elt,gc) *v = NULL;
+
+  if (vals)
+    {
+      v = VEC_alloc (constructor_elt, gc, list_length (vals));
+      for (t = vals; t; t = TREE_CHAIN (t))
+	CONSTRUCTOR_APPEND_ELT (v, TREE_PURPOSE (t), TREE_VALUE (t));
+    }
+
+  /* Sort field initializers by field offset.  */
+  VEC_qsort (constructor_elt, v, compare_elmt_bitpos);
+
+  return build_constructor (type, v);
+}

gcc-plugin/src/warn-unregistered.c

@@ -0,0 +1,264 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Inria
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Use extensions of the GNU C Library.  */
+#define _GNU_SOURCE 1
+
+#include <starpu-gcc/config.h>
+
+/* We must include starpu.h here, otherwise gcc will complain about a poisoned
+   malloc in xmmintrin.h.  */
+#include <starpu.h>
+
+#include <gcc-plugin.h>
+#include <plugin-version.h>
+
+#include <plugin.h>
+#include <cpplib.h>
+#include <tree.h>
+#include <tree-pass.h>
+#include <gimple.h>
+#include <diagnostic.h>
+#include <cgraph.h>
+
+#include <starpu-gcc/utils.h>
+#include <starpu-gcc/tasks.h>
+
+/* Return true if there exists a `starpu_vector_data_register' call for VAR
+   before GSI in its basic block.  */
+
+static bool
+registration_in_bb_p (gimple_stmt_iterator gsi, tree var)
+{
+  gcc_assert (SSA_VAR_P (var));
+
+  tree register_fn_name;
+
+  register_fn_name = get_identifier ("starpu_vector_data_register");
+
+  local_define (bool, registration_function_p, (const_tree obj))
+  {
+    /* TODO: Compare against the real fndecl.  */
+    return (obj != NULL_TREE
+	    && TREE_CODE (obj) == FUNCTION_DECL
+	    && DECL_NAME (obj) == register_fn_name);
+  };
+
+  bool found;
+
+  for (found = false;
+       !gsi_end_p (gsi) && !found;
+       gsi_prev (&gsi))
+    {
+      gimple stmt;
+
+      stmt = gsi_stmt (gsi);
+      if (is_gimple_call (stmt))
+	{
+	  tree fn = gimple_call_fndecl (stmt);
+	  if (registration_function_p (fn))
+	    {
+	      tree arg = gimple_call_arg (stmt, 2);
+	      if (is_gimple_address (arg))
+	      	arg = TREE_OPERAND (arg, 0);
+
+	      if (((TREE_CODE (arg) == VAR_DECL
+		    || TREE_CODE (arg) == VAR_DECL)
+		   && refs_may_alias_p (arg, var))
+
+		  /* Both VAR and ARG should be SSA names, otherwise, if ARG
+		     is a VAR_DECL, `ptr_derefs_may_alias_p' will
+		     conservatively assume that they may alias.  */
+		  || (TREE_CODE (var) == SSA_NAME
+		      && TREE_CODE (arg) != VAR_DECL
+		      && ptr_derefs_may_alias_p (arg, var)))
+		{
+		  if (verbose_output_p)
+		    {
+		      var = TREE_CODE (var) == SSA_NAME ? SSA_NAME_VAR (var) : var;
+		      inform (gimple_location (stmt),
+			      "found registration of variable %qE",
+			      DECL_NAME (var));
+		    }
+		  found = true;
+		}
+	    }
+	}
+    }
+
+  return found;
+}
+
+/* Return true if BB is dominated by a registration of VAR.  */
+
+static bool
+dominated_by_registration (gimple_stmt_iterator gsi, tree var)
+{
+  /* Is there a registration call for VAR in GSI's basic block?  */
+  if (registration_in_bb_p (gsi, var))
+    return true;
+
+  edge e;
+  edge_iterator ei;
+  bool found = false;
+
+  /* If every incoming edge is dominated by a registration, then we're
+     fine.
+
+     FIXME: This triggers false positives when registration is done in a
+     loop, because there's always an edge through which no registration
+     happens--the edge corresponding to the case where the loop is not
+     entered.  */
+
+  FOR_EACH_EDGE (e, ei, gsi_bb (gsi)->preds)
+    {
+      if (!dominated_by_registration (gsi_last_bb (e->src), var))
+	return false;
+      else
+	found = true;
+    }
+
+  return found;
+}
+
+/* Return true if NAME aliases a global variable or a PARM_DECL.  Note that,
+   for the former, `ptr_deref_may_alias_global_p' is way too conservative,
+   hence this approach.  */
+
+static bool
+ssa_name_aliases_global_or_parm_p (const_tree name)
+{
+  gcc_assert (TREE_CODE (name) == SSA_NAME);
+
+  if (TREE_CODE (SSA_NAME_VAR (name)) == PARM_DECL)
+    return true;
+  else
+    {
+      gimple def_stmt;
+
+      def_stmt = SSA_NAME_DEF_STMT (name);
+      if (is_gimple_assign (def_stmt))
+	{
+	  tree rhs = gimple_assign_rhs1 (def_stmt);
+
+	  if (TREE_CODE (rhs) == VAR_DECL
+	      && (DECL_EXTERNAL (rhs) || TREE_STATIC (rhs)))
+	    return true;
+	}
+    }
+
+  return false;
+}
+
+
+/* Validate the arguments passed to tasks in FN's body.  */
+
+static void
+validate_task_invocations (tree fn)
+{
+  gcc_assert (TREE_CODE (fn) == FUNCTION_DECL);
+
+  const struct cgraph_node *cgraph;
+  const struct cgraph_edge *callee;
+
+  cgraph = cgraph_get_node (fn);
+
+  /* When a definition of IMPL is available, check its callees.  */
+  if (cgraph != NULL)
+    for (callee = cgraph->callees;
+	 callee != NULL;
+	 callee = callee->next_callee)
+      {
+	if (task_p (callee->callee->decl))
+	  {
+	    unsigned i;
+	    gimple call_stmt = callee->call_stmt;
+
+	    for (i = 0; i < gimple_call_num_args (call_stmt); i++)
+	      {
+		tree arg = gimple_call_arg (call_stmt, i);
+
+		if (TREE_CODE (arg) == ADDR_EXPR
+		    && TREE_CODE (TREE_OPERAND (arg, 0)) == VAR_DECL
+		    && (TREE_CODE (TREE_TYPE (TREE_OPERAND (arg, 0)))
+			== ARRAY_TYPE))
+		  /* This is a "pointer-to-array" of a variable, so what we
+		     really care about is the variable itself.  */
+		  arg = TREE_OPERAND (arg, 0);
+
+		if ((POINTER_TYPE_P (TREE_TYPE (arg))
+		     || (TREE_CODE (TREE_TYPE (arg)) == ARRAY_TYPE))
+		    && ((TREE_CODE (arg) == VAR_DECL
+			 && !TREE_STATIC (arg)
+			 && !DECL_EXTERNAL (arg)
+			 && !TREE_NO_WARNING (arg))
+			|| (TREE_CODE (arg) == SSA_NAME
+			    && !ssa_name_aliases_global_or_parm_p (arg))))
+		  {
+		    if (!dominated_by_registration (gsi_for_stmt (call_stmt),
+						    arg))
+		      {
+			if (TREE_CODE (arg) == SSA_NAME)
+			  {
+			    tree var = SSA_NAME_VAR (arg);
+			    if (DECL_NAME (var) != NULL)
+			      arg = var;
+
+			    /* TODO: Check whether we can get the original
+			       variable name via ARG's DEF_STMT.  */
+			  }
+
+			if (TREE_CODE (arg) == VAR_DECL
+			    && DECL_NAME (arg) != NULL_TREE)
+			  warning_at (gimple_location (call_stmt), 0,
+				      "variable %qE may be used unregistered",
+				      DECL_NAME (arg));
+			else
+			  warning_at (gimple_location (call_stmt), 0,
+				      "argument %i may be used unregistered",
+				      i);
+		      }
+		  }
+	      }
+	  }
+      }
+}
+
+/* A pass to warn about possibly unregistered task arguments.  */
+
+static unsigned int
+warn_starpu_unregistered (void)
+{
+  tree fndecl;
+
+  fndecl = current_function_decl;
+  gcc_assert (TREE_CODE (fndecl) == FUNCTION_DECL);
+
+  if (!task_p (fndecl))
+    validate_task_invocations (fndecl);
+
+  return 0;
+}
+
+struct opt_pass pass_warn_starpu_unregistered =
+  {
+    designated_field_init (type, GIMPLE_PASS),
+    designated_field_init (name, "warn_starpu_unregistered"),
+    designated_field_init (gate, NULL),
+    designated_field_init (execute, warn_starpu_unregistered),
+
+    /* The rest is zeroed.  */
+  };

gcc-plugin/tests/Makefile.am

@@ -0,0 +1,118 @@
+# StarPU --- Runtime system for heterogeneous multicore architectures.
+#
+# Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+#
+# StarPU is free software; you can redistribute it and/or modify
+# it under the terms of the GNU Lesser General Public License as published by
+# the Free Software Foundation; either version 2.1 of the License, or (at
+# your option) any later version.
+#
+# StarPU is distributed in the hope that it will be useful, but
+# WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
+#
+# See the GNU Lesser General Public License in COPYING.LGPL for more details.
+
+
+gcc_tests =					\
+  base.c					\
+  pointers.c					\
+  output-pointer.c				\
+  output-pointer-errors.c			\
+  register.c					\
+  register-errors.c				\
+  registered.c					\
+  registered-errors.c				\
+  acquire.c					\
+  acquire-errors.c				\
+  release.c					\
+  release-errors.c				\
+  unregister.c					\
+  unregister-errors.c				\
+  task-errors.c					\
+  scalar-tasks.c				\
+  pointer-tasks.c				\
+  external-task-impl.c				\
+  no-initialize.c				\
+  lib-user.c					\
+  wait-errors.c					\
+  heap-allocated.c				\
+  heap-allocated-errors.c			\
+  verbose.c					\
+  debug-tree.c					\
+  opencl.c					\
+  opencl-errors.c				\
+  shutdown-errors.c
+
+EXTRA_DIST =
+
+if HAVE_PTR_DEREFS_MAY_ALIAS_P
+
+gcc_tests += warn-unregistered.c
+
+else !HAVE_PTR_DEREFS_MAY_ALIAS_P
+
+EXTRA_DIST += warn-unregistered.c
+
+endif !HAVE_PTR_DEREFS_MAY_ALIAS_P
+
+if !STARPU_USE_OPENCL
+
+# XXX: This test simulates a buggy OpenCL implementation, and thus
+# cannot be run then a real <cl_platform.h> is included.
+gcc_tests += opencl-types.c
+
+# This test simulates errors when lacking an OpenCL implementation.
+gcc_tests += opencl-lacking.c
+
+else STARPU_USE_OPENCL
+
+EXTRA_DIST +=					\
+  opencl-types.c				\
+  opencl-lacking.c
+
+endif STARPU_USE_OPENCL
+
+
+dist_noinst_HEADERS = mocks.h
+
+CLEANFILES = *.gimple *.o			\
+  base						\
+  pointers					\
+  register					\
+  registered					\
+  release					\
+  scalar-tasks					\
+  pointer-tasks					\
+  lib-user					\
+  output-pointer				\
+  unregister					\
+  heap-allocated				\
+  acquire					\
+  opencl
+
+
+EXTRA_DIST += ./run-test.in			\
+  my-lib.h my-lib.c				\
+  test.cl					\
+  $(gcc_tests)
+
+# The test suite assumes that the CPU back-end is available.
+if RUN_GCC_PLUGIN_TESTS
+
+TESTS = $(gcc_tests)
+if STARPU_HAVE_AM111
+LOG_COMPILER = ./run-test
+else
+TESTS_ENVIRONMENT = ./run-test
+endif
+
+else !RUN_GCC_PLUGIN_TESTS
+
+check-hook:
+	-@echo "GNU Guile or CPU back-end not available, test suite not run."
+
+endif !RUN_GCC_PLUGIN_TESTS
+
+showcheck:
+	-cat $(TEST_LOGS) /dev/null

gcc-plugin/tests/acquire-errors.c

@@ -0,0 +1,49 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+extern float *y;
+
+static const double a[123];
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  int x[123] __attribute__ ((unused));
+  static char z[345] __attribute__ ((unused));
+
+#pragma starpu register x
+
+#pragma starpu acquire /* (error "parse error") */
+#pragma starpu acquire 123 /* (error "neither a pointer nor an array") */
+#pragma starpu acquire does_not_exit /* (error "unbound variable") */
+
+#pragma starpu acquire argc /* (error "neither a pointer nor an array") */
+#pragma starpu acquire y
+#pragma starpu acquire x
+
+#pragma starpu acquire x z			  /* (error "junk after") */
+
+  /* XXX: Uncomment below when this is supported.  */
+#if 0
+#pragma starpu acquire z /* error "not registered" */
+#pragma starpu acquire a /* error "not registered" */
+#pragma starpu acquire argv /* error "not registered" */
+#endif
+
+  return 1;
+}

gcc-plugin/tests/acquire.c

@@ -0,0 +1,61 @@
+/* GCC-StarPU
+   Copyright (C) 2011 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Test whether `#pragma starpu acquire ...' generates the right code.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+
+static void
+foo (char *x, int foo)
+{
+  expected_acquire_arguments.pointer = x;
+#pragma starpu acquire x
+}
+
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  int x[123];
+  static char z[345];
+
+  expected_register_arguments.pointer = x;
+  expected_register_arguments.elements = 123;
+  expected_register_arguments.element_size = sizeof x[0];
+#pragma starpu register x
+
+  expected_acquire_arguments.pointer = x;
+#pragma starpu acquire x
+
+  expected_register_arguments.pointer = z;
+  expected_register_arguments.elements = sizeof z;
+  expected_register_arguments.element_size = sizeof z[0];
+#pragma starpu register z
+
+  expected_acquire_arguments.pointer = z;
+#pragma starpu acquire z
+
+  foo (z, 345);
+
+  assert (data_register_calls == 2);
+  assert (data_acquire_calls == 3);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/base.c

@@ -0,0 +1,157 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+
+
+/* The task under test.  */
+
+static void my_scalar_task (int x, unsigned char y, int z) __attribute__ ((task));
+
+static void my_scalar_task_cpu (int, unsigned char, int)
+  __attribute__ ((task_implementation ("cpu", my_scalar_task)));
+static void my_scalar_task_opencl (int, unsigned char, int)
+  __attribute__ ((task_implementation ("opencl", my_scalar_task)));
+
+static void
+my_scalar_task_cpu (int x, unsigned char y, int z)
+{
+  printf ("%s: x = %i, y = %i, z = %i\n", __func__, x, (int) y, z);
+}
+
+static void
+my_scalar_task_opencl (int x, unsigned char y, int z)
+{
+  printf ("%s: x = %i, y = %i, z = %i\n", __func__, x, (int) y, z);
+}
+
+
+/* Another task, where task implementation declarations are interleaved with
+   task definitions.  */
+
+static void my_other_task (int x) __attribute__ ((task));
+
+static void my_other_task_cpu (int)
+  __attribute__ ((task_implementation ("cpu", my_other_task)));
+static void my_other_task_cpu_bis (int)
+  __attribute__ ((task_implementation ("cpu", my_other_task)));
+
+static void
+my_other_task_cpu (int x)
+{
+  printf ("cpu\n");
+}
+
+static void
+my_other_task_cpu_bis (int x)
+{
+  printf ("second cpu implementation\n");
+}
+
+static void my_other_task_opencl (int)
+  __attribute__ ((task_implementation ("opencl", my_other_task)));
+
+static void
+my_other_task_opencl (int x)
+{
+  printf ("opencl\n");
+}
+
+
+/* Task with a body.  */
+
+static void my_task_with_body (int x) __attribute__ ((task));
+
+static void
+my_task_with_body (int x)
+{
+  /* This body is implicitly the "cpu" implementation of the task.  */
+  int y = x + 2;
+  printf ("body: %i\n", y - 2);
+}
+
+static void my_task_with_body_opencl (int x)
+  __attribute__ ((task_implementation ("opencl", my_task_with_body)));
+
+static void
+my_task_with_body_opencl (int x)
+{
+}
+
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+  assert (initialized == 1);
+
+#pragma starpu hello
+
+  int x = 42, z = 99;
+  unsigned char y = 77;
+  long y_as_long_int = 77;
+
+  struct insert_task_argument expected[] =
+    {
+      { STARPU_VALUE, &x, sizeof x },
+      { STARPU_VALUE, &y, sizeof y },
+      { STARPU_VALUE, &z, sizeof z },
+      { 0, 0, 0 }
+    };
+
+  expected_insert_task_arguments = expected;
+
+  /* Invoke the task, which should make sure it gets called with
+     EXPECTED.  */
+  my_scalar_task (x, y, z);
+
+  /* Invoke the task using literal constants instead of variables.  */
+  my_scalar_task (42, y, z);
+  my_scalar_task (x, 77, z);
+  my_scalar_task (x, y, 99);
+
+  my_scalar_task (42, 77, z);
+  my_scalar_task (x, 77, 99);
+  my_scalar_task (42, y, 99);
+
+  my_scalar_task (42, 77, 99);
+
+  my_scalar_task (42, y_as_long_int, 99);
+
+  assert (tasks_submitted == 9);
+
+  struct insert_task_argument expected2[] =
+    {
+      { STARPU_VALUE, &x, sizeof x },
+      { 0, 0, 0 }
+    };
+
+  tasks_submitted = 0;
+  expected_insert_task_arguments = expected2;
+
+  my_other_task (42);
+  assert (tasks_submitted == 1);
+
+  my_task_with_body (42);
+  assert (tasks_submitted == 2);
+
+#pragma starpu shutdown
+  assert (initialized == 0);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/debug-tree.c

@@ -0,0 +1,32 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions compile) */
+
+#pragma starpu debug_tree int			  /* (note "debug_tree") */
+
+int
+foo (void)
+{
+#pragma starpu debug_tree foo			  /* (note "debug_tree")  */
+
+  static int x = 2;
+#pragma starpu debug_tree x			  /* (note "debug_tree") */
+
+  return x;
+}
+
+#pragma starpu debug_tree void int foo /* (note "debug_tree") *//* (warning "extraneous") */

gcc-plugin/tests/external-task-impl.c

@@ -0,0 +1,33 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions compile) */
+
+void the_task (int foo, float bar[foo])
+  __attribute__ ((task));
+
+static void the_task_cpu (int foo, float bar[foo])
+  __attribute__ ((task_implementation ("cpu", the_task)));
+
+/* Make sure GCC doesn't barf on this one.  */
+extern void the_task_cuda (int foo, float bar[foo])
+  __attribute__ ((task_implementation ("cuda", the_task)));
+
+static void
+the_task_cpu (int foo, float bar[foo])
+{
+  /* Nothingness.  */
+}

gcc-plugin/tests/heap-allocated-errors.c

@@ -0,0 +1,36 @@
+/* GCC-StarPU
+   Copyright (C) 2011 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions compile (cflags "-Wno-unused-variable")) */
+
+static int global[123]              /* (error "cannot be used") */
+  __attribute__ ((heap_allocated, used));
+
+extern int external[123]            /* (error "cannot be used") */
+  __attribute__ ((heap_allocated));
+
+void
+foo (size_t size)
+{
+  float scalar /* (error "must have an array type") */
+    __attribute__ ((heap_allocated));
+  float *ptr   /* (error "must have an array type") */
+    __attribute__ ((heap_allocated));
+  float incomp[]  /* (error "incomplete array type") */
+    __attribute__ ((heap_allocated));
+  float incomp2[size][3][]  /* (error "incomplete element type") */
+    __attribute__ ((heap_allocated));
+}

gcc-plugin/tests/heap-allocated.c

@@ -0,0 +1,88 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+
+static void
+foo (size_t size)
+{
+  /* See ISO C99, Section 6.7.5.2 ("Array Declarators") and Section 6.7.6
+     ("Type names").  */
+
+  float *test_array_parm (float m[size][23], int x, int y)
+  {
+    assert ((char *) &m[0][0] != (char *) &m);
+    return &m[x][y];
+  }
+
+  size_t minus_one = size - 1;
+  float *test_array_static_parm (float m[static minus_one][23], int x, int y)
+  {
+    assert ((char *) &m[0][0] != (char *) &m);
+    return &m[x][y];
+  }
+
+  float *test_pointer_parm (float *m, int x, int y)
+  {
+    return &m[23 * x + y];
+  }
+
+  expected_malloc_argument = size * 23 * sizeof (float);
+
+  /* The idea is that this code should be compilable both with and without
+     the attribute.  */
+  float m[size][23] __attribute__ ((heap_allocated));
+
+  assert (malloc_calls == 1);
+
+  /* `&m' points to the on-stack area that stores the underlying pointer,
+     whereas `m' and `m[0][0]' should point to the heap-allocated area.  */
+  assert ((char *) &m != (char *) m);
+  assert ((char *) &m[0][0] != (char *) &m);
+  assert ((char *) &m[0][0] == (char *) m);
+
+  /* Make sure "array arithmetic" works.  */
+  assert ((char *) &m[0][1] - (char *) &m[0][0] == sizeof m[0][0]);
+  assert ((char *) &m[1][0] - (char *) &m[0][22] == sizeof m[0][0]);
+
+  unsigned int x, y;
+  for (x = 0; x < size; x++)
+    for (y = 0; y < 23; y++)
+      {
+	assert (&m[x][y] == test_array_parm (m, x, y));
+	assert (&m[x][y] == test_array_static_parm (m, x, y));
+	assert (&m[x][y] == test_pointer_parm ((float *) m, x, y));
+      }
+
+  /* Freed when going out of scope.  */
+  expected_free_argument = m;
+}
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  /* Choose the size such that the process would most likely segfault if
+     `foo' tried to allocate this much data on the stack.  */
+  foo (100000);
+
+  assert (free_calls == 1);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/lib-user.c

@@ -0,0 +1,71 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Test whether tasks defined in another compilation unit can actually be
+   used.  */
+
+/* (instructions run (dependencies "my-lib.c")) */
+
+#include <mocks.h>
+
+#include <my-lib.h>
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  /* Align X so that the assumptions behind `dummy_pointer_to_handle'
+     hold.  */
+  static const signed char x[] __attribute__ ((aligned (8))) =
+    { 0, 1, 2, 3, 4, 5 };
+
+  float y[sizeof x];
+
+  static const char forty_two = 42;
+  static const int  sizeof_x = sizeof x;
+
+  struct insert_task_argument expected_pointer_task[] =
+    {
+      { STARPU_VALUE, &forty_two, sizeof forty_two },
+      { STARPU_R,  x },
+      { STARPU_RW, y },
+      { STARPU_VALUE, &sizeof_x, sizeof sizeof_x },
+      { 0, 0, 0 }
+    };
+
+  expected_insert_task_arguments = expected_pointer_task;
+
+  expected_register_arguments.pointer = (void *) x;
+  expected_register_arguments.elements = sizeof x / sizeof x[0];
+  expected_register_arguments.element_size = sizeof x[0];
+#pragma starpu register x
+
+  expected_register_arguments.pointer = y;
+  expected_register_arguments.elements = sizeof y / sizeof y[0];
+  expected_register_arguments.element_size = sizeof y[0];
+#pragma starpu register y
+
+  /* Invoke the task, which should make sure it gets called with
+     EXPECTED.  */
+  my_task (42, x, y, sizeof x);
+
+  assert (tasks_submitted == 1);
+
+#pragma starpu shutdown
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/mocks.h

@@ -0,0 +1,625 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Testing library, including stubs of StarPU functions.  */
+
+#ifndef STARPU_GCC_PLUGIN
+# error barf!
+#endif
+
+#ifndef STARPU_USE_CPU
+# error damn it!
+#endif
+
+#undef NDEBUG
+
+#include <stdlib.h>
+#include <stdarg.h>
+#include <stdint.h>
+#include <string.h>
+#include <assert.h>
+#include <common/uthash.h>
+#include <stdint.h>
+
+
+/* Typedefs as found in <CL/cl_platform.h>.  */
+
+typedef int8_t         cl_char;
+typedef uint8_t        cl_uchar;
+typedef int16_t        cl_short;
+typedef uint16_t       cl_ushort;
+typedef int32_t        cl_int;
+typedef uint32_t       cl_uint;
+#ifdef BREAK_CL_LONG
+/* Make `cl_long' different from `long' for test purposes.  */
+typedef int16_t        cl_long;
+typedef uint16_t       cl_ulong;
+#else
+typedef int64_t        cl_long;
+typedef uint64_t       cl_ulong;
+#endif
+
+typedef uint16_t       cl_half;
+typedef float          cl_float;
+typedef double         cl_double;
+
+
+/* Stub used for testing purposes.  */
+
+/* Number of tasks submitted.  */
+static unsigned int tasks_submitted;
+
+struct insert_task_argument
+{
+  /* `STARPU_VALUE', etc. */
+  int type;
+
+  /* Pointer to the expected value.  */
+  const void *pointer;
+
+  /* Size in bytes of the data pointed to.  */
+  size_t size;
+};
+
+/* Pointer to a zero-terminated array listing the expected
+   `starpu_insert_task' arguments.  */
+const struct insert_task_argument *expected_insert_task_arguments;
+
+/* Expected targets of the codelets submitted.  */
+static int expected_insert_task_targets = STARPU_CPU | STARPU_OPENCL;
+
+
+int
+starpu_insert_task (struct starpu_codelet *cl, ...)
+{
+  assert (cl->name != NULL && strlen (cl->name) > 0);
+  assert (cl->where == expected_insert_task_targets);
+
+  assert ((cl->where & STARPU_CPU) == 0
+	  ? cl->cpu_funcs[0] == NULL
+	  : cl->cpu_funcs[0] != NULL);
+  assert ((cl->where & STARPU_OPENCL) == 0
+	  ? cl->opencl_funcs[0] == NULL
+	  : cl->opencl_funcs[0] != NULL);
+  assert ((cl->where & STARPU_CUDA) == 0
+	  ? cl->cuda_funcs[0] == NULL
+	  : cl->cuda_funcs[0] != NULL);
+
+  va_list args;
+  size_t i, scalars, pointers, cl_args_offset;
+  void *pointer_args[123];
+  struct starpu_vector_interface pointer_args_ifaces[123];
+  unsigned char cl_args[234];
+
+  va_start (args, cl);
+
+  const struct insert_task_argument *expected;
+  for (expected = expected_insert_task_arguments,
+	 cl_args_offset = 1, scalars = 0, pointers = 0;
+       expected->type != 0;
+       expected++)
+    {
+      int type;
+
+      type = va_arg (args, int);
+      assert (type == expected->type);
+
+      switch (type)
+	{
+	case STARPU_VALUE:
+	  {
+	    void *arg;
+	    size_t size;
+
+	    arg = va_arg (args, void *);
+	    size = va_arg (args, size_t);
+
+	    assert (size == expected->size);
+	    assert (arg != NULL);
+	    assert (!memcmp (arg, expected->pointer, size));
+
+	    /* Pack ARG into CL_ARGS.  */
+	    assert (cl_args_offset + size + sizeof size < sizeof cl_args);
+	    memcpy (&cl_args[cl_args_offset], &size, sizeof size);
+	    cl_args_offset += sizeof size;
+	    memcpy (&cl_args[cl_args_offset], arg, size);
+	    cl_args_offset += size;
+
+	    scalars++;
+	    break;
+	  }
+
+	case STARPU_RW:
+	case STARPU_R:
+	case STARPU_W:
+	  {
+	    starpu_data_handle_t handle;
+	    handle = starpu_data_lookup (expected->pointer);
+
+	    assert (type == cl->modes[pointers]);
+	    assert (va_arg (args, void *) == handle);
+	    assert (pointers + 1
+		    < sizeof pointer_args_ifaces / sizeof pointer_args_ifaces[0]);
+
+	    pointer_args_ifaces[pointers].ptr = (uintptr_t) expected->pointer;
+	    pointer_args_ifaces[pointers].dev_handle =
+	      (uintptr_t) expected->pointer;	  /* for OpenCL */
+	    pointer_args_ifaces[pointers].elemsize = 1;
+	    pointer_args_ifaces[pointers].nx = 1;
+	    pointer_args_ifaces[pointers].offset = 0;
+
+	    pointers++;
+	    break;
+	  }
+
+	default:
+	  abort ();
+	}
+    }
+
+  va_end (args);
+
+  /* Make sure all the arguments were consumed.  */
+  assert (expected->type == 0);
+
+  tasks_submitted++;
+
+  /* Finish packing the scalar arguments in CL_ARGS.  */
+  cl_args[0] = (unsigned char) scalars;
+  for (i = 0; i < pointers; i++)
+    pointer_args[i] = &pointer_args_ifaces[i];
+
+  /* Call the codelets.  */
+  if (cl->where & STARPU_CPU)
+    cl->cpu_funcs[0] (pointer_args, cl_args);
+  if (cl->where & STARPU_OPENCL)
+    cl->opencl_funcs[0] (pointer_args, cl_args);
+  if (cl->where & STARPU_CUDA)
+    cl->cuda_funcs[0] (pointer_args, cl_args);
+
+  return 0;
+}
+
+/* Our own implementation of `starpu_codelet_unpack_args', for debugging
+   purposes.  */
+
+void
+starpu_codelet_unpack_args (void *cl_raw_arg, ...)
+{
+  va_list args;
+  size_t nargs, arg, offset, size;
+  unsigned char *cl_arg;
+
+  cl_arg = (unsigned char *) cl_raw_arg;
+
+  nargs = *cl_arg;
+
+  va_start (args, cl_raw_arg);
+
+  for (arg = 0, offset = 1;
+       arg < nargs;
+       arg++, offset += sizeof (size_t) + size)
+    {
+      void *argp;
+
+      argp = va_arg (args, void *);
+      size = *(size_t *) &cl_arg[offset];
+
+      memcpy (argp, &cl_arg[offset + sizeof size], size);
+    }
+
+  va_end (args);
+}
+
+
+/* Data handles.  A hash table mapping pointers to handles is maintained,
+   which allows us to mimic the actual behavior of libstarpu.  */
+
+/* Entry in the `registered_handles' hash table.  `starpu_data_handle_t' is
+   assumed to be a pointer to this structure.  */
+struct handle_entry
+{
+  UT_hash_handle hh;
+  void *pointer;
+  starpu_data_handle_t handle;
+};
+
+#define handle_to_entry(h) ((struct handle_entry *) (h))
+#define handle_to_pointer(h)				\
+  ({							\
+    assert ((h) != NULL);				\
+    assert (handle_to_entry (h)->handle == (h));	\
+    handle_to_entry (h)->pointer;			\
+   })
+
+static struct handle_entry *registered_handles;
+
+starpu_data_handle_t
+starpu_data_lookup (const void *ptr)
+{
+  starpu_data_handle_t result;
+
+  struct handle_entry *entry;
+
+  HASH_FIND_PTR (registered_handles, &ptr, entry);
+  if (STARPU_UNLIKELY (entry == NULL))
+    result = NULL;
+  else
+    result = entry->handle;
+
+  return result;
+}
+
+void *
+starpu_handle_get_local_ptr (starpu_data_handle_t handle)
+{
+  return handle_to_pointer (handle);
+}
+
+
+/* Data registration.  */
+
+struct data_register_arguments
+{
+  /* A pointer to the vector being registered.  */
+  void *pointer;
+
+  /* Number of elements in the vector.  */
+  size_t elements;
+
+  /* Size of individual elements.  */
+  size_t element_size;
+};
+
+/* Number of `starpu_vector_data_register' calls.  */
+static unsigned int data_register_calls;
+
+/* Variable describing the expected `starpu_vector_data_register'
+   arguments.  */
+struct data_register_arguments expected_register_arguments;
+
+void
+starpu_vector_data_register (starpu_data_handle_t *handle,
+			     uint32_t home_node, uintptr_t ptr,
+			     uint32_t count, size_t elemsize)
+{
+  /* Sometimes tests cannot tell what the pointer will be (for instance, for
+     the `registered' attribute), and thus pass NULL as the expected
+     pointer.  */
+  if (expected_register_arguments.pointer != NULL)
+    assert ((void *) ptr == expected_register_arguments.pointer);
+  else
+    /* Allow users to check the pointer afterward.  */
+    expected_register_arguments.pointer = (void *) ptr;
+
+  assert (count == expected_register_arguments.elements);
+  assert (elemsize == expected_register_arguments.element_size);
+
+  data_register_calls++;
+
+  /* Add PTR to the REGISTERED_HANDLES hash table.  */
+
+  struct handle_entry *entry = malloc (sizeof (*entry));
+  assert (entry != NULL);
+
+  entry->pointer = (void *) ptr;
+  entry->handle = (starpu_data_handle_t) entry;
+
+  HASH_ADD_PTR(registered_handles, pointer, entry);
+
+  *handle = (starpu_data_handle_t) entry;
+}
+
+
+/* Data acquisition.  */
+
+struct data_acquire_arguments
+{
+  /* Pointer to the data being acquired.  */
+  void *pointer;
+};
+
+struct data_release_arguments
+{
+  /* Pointer to the data being released.  */
+  void *pointer;
+};
+
+/* Number of `starpu_data_{acquire,release}' calls.  */
+static unsigned int data_acquire_calls, data_release_calls;
+
+/* Variable describing the expected `starpu_data_{acquire,release}'
+   arguments.  */
+struct data_acquire_arguments expected_acquire_arguments;
+struct data_release_arguments expected_release_arguments;
+
+int
+starpu_data_acquire (starpu_data_handle_t handle, enum starpu_access_mode mode)
+{
+  /* XXX: Currently only `STARPU_RW'.  */
+  assert (mode == STARPU_RW);
+
+  assert (handle_to_pointer (handle) == expected_acquire_arguments.pointer);
+  data_acquire_calls++;
+
+  return 0;
+}
+
+void
+starpu_data_release (starpu_data_handle_t handle)
+{
+  assert (handle_to_pointer (handle) == expected_release_arguments.pointer);
+  data_release_calls++;
+}
+
+
+/* Data acquisition.  */
+
+struct data_unregister_arguments
+{
+  /* Pointer to the data being unregistered.  */
+  void *pointer;
+};
+
+/* Number of `starpu_data_unregister' calls.  */
+static unsigned int data_unregister_calls;
+
+/* Variable describing the expected `starpu_data_unregister' arguments.  */
+struct data_unregister_arguments expected_unregister_arguments;
+
+void
+starpu_data_unregister (starpu_data_handle_t handle)
+{
+  assert (handle != NULL);
+
+  struct handle_entry *entry = handle_to_entry (handle);
+
+  assert (entry->pointer != NULL);
+  assert (entry->pointer == expected_unregister_arguments.pointer);
+
+  /* Remove the PTR -> HANDLE mapping.  If a mapping from PTR to another
+     handle existed before (e.g., when using filters), it becomes visible
+     again.  */
+  HASH_DEL (registered_handles, entry);
+  entry->pointer = NULL;
+  free (entry);
+
+  data_unregister_calls++;
+}
+
+
+/* Heap allocation.  */
+
+/* Number of `starpu_malloc' and `starpu_free' calls.  */
+static unsigned int malloc_calls, free_calls;
+
+static size_t expected_malloc_argument;
+static void *expected_free_argument;
+
+int
+starpu_malloc (void **ptr, size_t size)
+{
+  assert (size == expected_malloc_argument);
+
+  *ptr = malloc (size);
+  malloc_calls++;
+
+  return 0;
+}
+
+int
+starpu_free (void *ptr)
+{
+  assert (starpu_data_lookup (ptr) == NULL);
+  assert (ptr == expected_free_argument);
+  free_calls++;
+  return 0;
+}
+
+
+/* OpenCL support.  */
+
+#ifndef STARPU_USE_OPENCL
+
+# define STARPU_USE_OPENCL 1
+
+/* The `opencl' pragma needs this structure, so make sure it's defined.  */
+struct starpu_opencl_program
+{
+  /* Nothing.  */
+};
+
+typedef int cl_event;
+typedef int cl_kernel;
+typedef int cl_command_queue;
+
+extern cl_int clSetKernelArg (cl_kernel, cl_uint, size_t, const void *);
+
+extern cl_int
+clEnqueueNDRangeKernel(cl_command_queue /* command_queue */,
+                       cl_kernel        /* kernel */,
+                       cl_uint          /* work_dim */,
+                       const size_t *   /* global_work_offset */,
+                       const size_t *   /* global_work_size */,
+                       const size_t *   /* local_work_size */,
+                       cl_uint          /* num_events_in_wait_list */,
+                       const cl_event * /* event_wait_list */,
+                       cl_event *       /* event */);
+
+#endif
+
+
+/* Number of `load_opencl_from_string', `load_kernel', and `clSetKernelArg'
+   calls.  */
+static unsigned int load_opencl_calls, load_opencl_kernel_calls,
+  opencl_set_kernel_arg_calls, opencl_enqueue_calls, opencl_finish_calls,
+  opencl_collect_stats_calls, opencl_release_event_calls;
+
+struct load_opencl_arguments
+{
+  const char *source_file;
+  struct starpu_opencl_program *program;
+};
+
+/* Expected arguments.  */
+static struct load_opencl_arguments expected_load_opencl_arguments;
+
+struct cl_enqueue_kernel_arguments
+{
+  size_t * global_work_size;
+};
+
+/* Variable describing the expected `clEnqueueNDRangeKernel' arguments. */
+static struct cl_enqueue_kernel_arguments expected_cl_enqueue_kernel_arguments;
+
+
+int
+starpu_opencl_load_opencl_from_string (const char *source,
+				       struct starpu_opencl_program *program,
+				       const char *build_options)
+{
+  assert (source != NULL);		       /* FIXME: mmap file & check */
+  assert (program != expected_load_opencl_arguments.program);
+  load_opencl_calls++;
+  return 0;
+}
+
+int
+starpu_opencl_load_kernel (cl_kernel *kernel,
+			   cl_command_queue *queue,
+			   struct starpu_opencl_program *programs,
+			   const char *kernel_name, int devid)
+{
+  assert (kernel != NULL && queue != NULL && programs != NULL
+	  && kernel_name != NULL && devid == -42);
+  load_opencl_kernel_calls++;
+  return 0;
+}
+
+int
+starpu_worker_get_id (void)
+{
+  return 42;
+}
+
+int
+starpu_worker_get_devid (int id)
+{
+  return -id;
+}
+
+/* Set the INDEXth argument to KERNEL to the SIZE bytes pointed to by
+   VALUE.  */
+cl_int
+clSetKernelArg (cl_kernel kernel, cl_uint index, size_t size,
+		const void *value)
+{
+  size_t n;
+  const struct insert_task_argument *arg;
+
+  for (n = 0, arg = expected_insert_task_arguments;
+       n < index;
+       n++, arg++)
+    assert (arg->pointer != NULL);
+
+  switch (arg->type)
+    {
+    case STARPU_VALUE:
+      assert (size == arg->size);
+      assert (memcmp (arg->pointer, value, size) == 0);
+      break;
+
+    case STARPU_RW:
+    case STARPU_R:
+    case STARPU_W:
+      assert (size == sizeof (void *));
+      assert (* (void **) value == arg->pointer);
+      break;
+
+    default:
+      abort ();
+    }
+
+  opencl_set_kernel_arg_calls++;
+  return 0;
+}
+
+cl_int
+clEnqueueNDRangeKernel(cl_command_queue command_queue,
+                       cl_kernel        kernel,
+                       cl_uint          work_dim,
+                       const size_t *   global_work_offset,
+                       const size_t *   global_work_size,
+                       const size_t *   local_work_size,
+                       cl_uint          num_events_in_wait_list,
+                       const cl_event * event_wait_list,
+                       cl_event *       event)
+{
+  assert (*local_work_size == 1);
+  assert (*global_work_size == *expected_cl_enqueue_kernel_arguments.global_work_size);
+
+  opencl_enqueue_calls++;
+  return 0;
+}
+
+cl_int
+clFinish (cl_command_queue command_queue)
+{
+  opencl_finish_calls++;
+  return 0;
+}
+
+cl_int
+starpu_opencl_collect_stats (cl_event event)
+{
+  opencl_collect_stats_calls++;
+  return 0;
+}
+
+cl_int
+clReleaseEvent (cl_event event)
+{
+  opencl_release_event_calls++;
+  return 0;
+}
+
+
+const char *
+starpu_opencl_error_string (cl_int s)
+{
+  return "mock";
+}
+
+
+/* Initialization.  */
+
+static int initialized;
+
+int
+starpu_init (struct starpu_conf *config)
+{
+  initialized++;
+  return 0;
+}
+
+
+/* Shutdown.  */
+
+void
+starpu_shutdown (void)
+{
+  initialized--;
+}

gcc-plugin/tests/my-lib.c

@@ -0,0 +1,44 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Example library of tasks.  */
+
+/* (instructions compile) */
+
+#include <my-lib.h>
+
+
+/* Task implementations: one is `static', one is global.  The codelet
+   wrapper, codelet, and task body should be instantiated in this file.  */
+
+/* The implicit CPU implementation of `my_task'.  */
+void
+my_task (signed char a, const signed char *p, float *q, int b)
+{
+  int i;
+
+  for (i = 0; i < b; i++)
+    *q = *p + a;
+}
+
+void
+my_task_opencl (signed char a, const signed char *p, float *q, int b)
+{
+  int i;
+
+  for (i = 0; i < b; i++)
+    *q = *p + a;
+}

gcc-plugin/tests/my-lib.h

@@ -0,0 +1,30 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Example library of tasks.  */
+
+#ifndef MY_LIB_H
+#define MY_LIB_H
+
+extern void my_task (signed char, const signed char *, float *, int)
+  __attribute__ ((task));
+
+/* One of the implementations of MY_TASK.  Since it's `extern', this should
+   not trigger generation of the codelet, wrapper, etc.  */
+extern void my_task_opencl (signed char, const signed char *, float *, int)
+  __attribute__ ((task_implementation ("opencl", my_task)));
+
+#endif /* MY_LIB_H */

gcc-plugin/tests/no-initialize.c

@@ -0,0 +1,23 @@
+/* GCC-StarPU
+   Copyright (C) 2011 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions compile) */
+
+int
+main (int argc, char *argv[]) /* (warning "does not initialize") */
+{
+  return 0;
+}

gcc-plugin/tests/opencl-errors.c

@@ -0,0 +1,54 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#include <mocks.h>	    /* for `starpu_opencl_load_opencl_from_string' */
+
+/* Claim that OpenCL is supported.  */
+#pragma starpu add_target "opencl"
+
+
+void my_task (int x, float a[x])
+  __attribute__ ((task));
+
+static void my_task_cpu (int x, float a[x])
+  __attribute__ ((task_implementation ("cpu", my_task)));
+
+static void my_task_opencl (int x, float a[x])
+  __attribute__ ((task_implementation ("opencl", my_task)));
+
+static void
+my_task_cpu (int x, float a[x])
+{
+}
+
+
+#pragma starpu opencl my_task "test.cl" "kern" 1 /* (error "not a.* task impl") */
+#pragma starpu opencl my_task_cpu  /* (error "not a.* task impl") */	\
+                      "test.cl" "kern" 1
+#pragma starpu opencl my_task_opencl "/dev/null" "kern" 1 /* (error "empty") */
+#pragma starpu opencl my_task_opencl "/does-not-exist/" "kern" 1 /* (error "failed to access") */
+
+#pragma starpu opencl my_task_opencl	  /* (error "wrong number of arg") */
+#pragma starpu opencl my_task_opencl 123 "kern" 1 /* (error "string constant") */
+#pragma starpu opencl my_task_opencl "test.cl" 123 1 /* (error "string constant") */
+#pragma starpu opencl my_task_opencl "test.cl" "kern" "a" /* (error "integral type") */
+#pragma starpu opencl my_task_opencl "test.cl" "kern" 1 "foo" /* (error "junk after") */
+
+void
+foo (void)
+{
+#pragma starpu opencl my_task_opencl "test.cl" "kern" 1 /* (error "top-level") */
+}

gcc-plugin/tests/opencl-lacking.c

@@ -0,0 +1,26 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions compile) */
+
+void my_task (int x, float a[x])
+  __attribute__ ((task));
+
+static void my_task_opencl (int x, float a[x])
+  __attribute__ ((task_implementation ("opencl", my_task)));
+
+#pragma starpu opencl my_task_opencl  /* (note "not generated") */	\
+               "test.cl" "kern" 8

gcc-plugin/tests/opencl-types.c

@@ -0,0 +1,171 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Make sure use of `size_t' as a task argument type is flagged.  */
+
+/* (instructions compile) */
+
+#undef NDEBUG
+
+/* Please gimme a broken `cl_long'!  */
+#define BREAK_CL_LONG
+
+#include <mocks.h>
+#include <unistd.h>
+#include <sys/types.h>				  /* for `uint' & co. */
+
+
+/* Make sure `size_t' is flagged.  */
+
+static void my_task (size_t size, int x[size])
+  __attribute__ ((task));
+
+static void my_task_cpu (size_t size, int x[size])
+  __attribute__ ((task_implementation ("cpu", my_task)));
+static void my_task_opencl (size_t size, int x[size]) /* (warning "size_t.*not.*known OpenCL type") */
+  __attribute__ ((task_implementation ("opencl", my_task)));
+
+static void
+my_task_cpu (size_t size, int x[size])
+{
+}
+
+static void
+my_task_opencl (size_t size, int x[size])
+{
+}
+
+
+/* Make sure types that have the same name in C and OpenCL but are actually
+   different are flagged.  We assume `sizeof (long) == 4' here.  */
+
+static void my_long_task (unsigned long size, int x[size])
+  __attribute__ ((task));
+
+static void my_long_task_cpu (unsigned long size, int x[size])
+  __attribute__ ((task_implementation ("cpu", my_long_task)));
+static void my_long_task_opencl (unsigned long size,  /* (warning "differs from the same-named OpenCL type") */
+				 int x[size])
+  __attribute__ ((task_implementation ("opencl", my_long_task)));
+
+static void
+my_long_task_cpu (unsigned long size, int x[size])
+{
+}
+
+static void
+my_long_task_opencl (unsigned long size, int x[size])
+{
+}
+
+
+/* Same with a pointer-to-long.  */
+
+static void my_long_ptr_task (unsigned long *p)
+  __attribute__ ((task));
+
+static void my_long_ptr_task_cpu (unsigned long *p)
+  __attribute__ ((task_implementation ("cpu", my_long_ptr_task)));
+static void my_long_ptr_task_opencl (unsigned long *p) /* (warning "differs from the same-named OpenCL type") */
+  __attribute__ ((task_implementation ("opencl", my_long_ptr_task)));
+
+static void
+my_long_ptr_task_cpu (unsigned long *p)
+{
+}
+
+static void
+my_long_ptr_task_opencl (unsigned long *p)
+{
+}
+
+
+/* Same with an array of unsigned chars.  */
+
+static void my_uchar_task (char c[])
+  __attribute__ ((task));
+
+static void my_uchar_task_cpu (char c[])
+  __attribute__ ((task_implementation ("cpu", my_uchar_task)));
+static void my_uchar_task_opencl (char c[]) /* (warning "differs in signedness from the same-named OpenCL type") */
+  __attribute__ ((task_implementation ("opencl", my_uchar_task)));
+
+static void
+my_uchar_task_cpu (char c[])
+{
+}
+
+static void
+my_uchar_task_opencl (char c[])
+{
+}
+
+
+/* "unsigned int" is aka. "uint".  */
+
+static void my_uint_task (const uint *c)
+  __attribute__ ((task));
+static void my_uint_task_cpu (const uint *c)
+  __attribute__ ((task_implementation ("cpu", my_uint_task)));
+static void my_uint_task_opencl (const uint *c)	  /* no warning */
+  __attribute__ ((task_implementation ("opencl", my_uint_task)));
+
+static void
+my_uint_task_cpu (const uint *c)
+{
+}
+
+static void
+my_uint_task_opencl (const uint *c)
+{
+}
+
+
+/* "unsigned char" is aka. "uchar".  */
+
+typedef float uchar;				  /* not a real `uchar'! */
+
+static void my_fake_uchar_task (const uchar *c)
+  __attribute__ ((task));
+static void my_fake_uchar_task_cpu (const uchar *c)
+  __attribute__ ((task_implementation ("cpu", my_fake_uchar_task)));
+static void my_fake_uchar_task_opencl (const uchar *c) /* (warning "differs from the same-named OpenCL type") */
+  __attribute__ ((task_implementation ("opencl", my_fake_uchar_task)));
+
+static void
+my_fake_uchar_task_cpu (const uchar *c)
+{
+}
+
+static void
+my_fake_uchar_task_opencl (const uchar *c)
+{
+}
+
+
+
+/* No OpenCL, no problems.  */
+
+static void my_cool_task (size_t size, long long x[size])
+  __attribute__ ((task));
+
+static void my_cool_task_cpu (size_t size, long long x[size])
+  __attribute__ ((task_implementation ("cpu", my_cool_task)));
+
+static void
+my_cool_task_cpu (size_t size, long long x[size])
+{
+}

gcc-plugin/tests/opencl.c

@@ -0,0 +1,73 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+#include <stdlib.h>
+
+/* Claim that OpenCL is supported.  */
+#pragma starpu add_target "opencl"
+
+
+static void my_task (int x, float a[x])
+  __attribute__ ((task));
+
+static void my_task_opencl (int x, float a[x])
+  __attribute__ ((task_implementation ("opencl", my_task)));
+
+#pragma starpu opencl my_task_opencl "test.cl" "kern" 8
+
+int
+main ()
+{
+  static float a[123];
+
+#pragma starpu initialize
+
+  memset (a, 0, sizeof a);
+
+  expected_register_arguments.pointer = a;
+  expected_register_arguments.elements = sizeof a / sizeof a[0];
+  expected_register_arguments.element_size = sizeof a[0];
+#pragma starpu register a
+
+  static int x = 123;
+  struct insert_task_argument expected[] =
+    {
+      { STARPU_VALUE, &x, sizeof x },
+      { STARPU_RW, a },
+      { 0, 0, 0 }
+    };
+
+  expected_insert_task_arguments = expected;
+  expected_insert_task_targets = STARPU_OPENCL;
+  size_t y = 8; expected_cl_enqueue_kernel_arguments.global_work_size = &y;
+
+  my_task (123, a);
+  my_task (123, a);
+  my_task (123, a);
+
+  assert (tasks_submitted == 3);
+  assert (load_opencl_calls == 1);
+  assert (load_opencl_kernel_calls == 3);
+  assert (opencl_set_kernel_arg_calls == 3 * 2);
+  assert (opencl_enqueue_calls == 3);
+  assert (opencl_finish_calls == 3);
+  assert (opencl_release_event_calls == 3);
+  assert (opencl_collect_stats_calls == 3);
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/output-pointer-errors.c

@@ -0,0 +1,25 @@
+/* GCC-StarPU
+   Copyright (C) 2011 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#define __output __attribute__ ((output))
+
+/* XXX: Currently, since `output' is a type attribute, we have no way to
+   restrict its use to PARM_DECLs of tasks.  */
+
+void f (__output int x)				  /* (error "not allowed") */
+  __attribute__ ((task));
+
+__output void g (int x);			  /* (error "not allowed") */

gcc-plugin/tests/output-pointer.c

@@ -0,0 +1,107 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+
+#define __output __attribute__ ((output))
+
+/* The tasks under test.  */
+
+static void my_pointer_task (int size, __output int *x)
+  __attribute__ ((task));
+
+static void my_pointer_task_cpu (int size, __output int *x)
+  __attribute__ ((task_implementation ("cpu", my_pointer_task)));
+static void my_pointer_task_opencl (int size, __output int *x)
+  __attribute__ ((task_implementation ("opencl", my_pointer_task)));
+
+static void
+my_pointer_task_cpu (int size, __output int *x)
+{
+  printf ("%s: x = %p, size = %i\n", __func__, x, size);
+}
+
+static void
+my_pointer_task_opencl (int size, int *x)
+{
+  printf ("%s: x = %p, size = %i\n", __func__, x, size);
+}
+
+
+
+static void my_array_task (int size, __output int x[size])
+  __attribute__ ((task));
+
+static void my_array_task_cpu (int size, __output int x[size])
+  __attribute__ ((task_implementation ("cpu", my_array_task)));
+static void my_array_task_opencl (int size, __output int x[size])
+  __attribute__ ((task_implementation ("opencl", my_array_task)));
+
+static void
+my_array_task_cpu (int size, __output int x[size])
+{
+  printf ("%s: x = %p, size = %i\n", __func__, x, size);
+}
+
+static void
+my_array_task_opencl (int size, __output int x[size])
+{
+  printf ("%s: x = %p, size = %i\n", __func__, x, size);
+}
+
+
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  int size = 42;
+  int x[size];
+
+  /* Register X (don't use the pragma, to avoid mixing concerns in this
+     test.)  */
+
+  starpu_data_handle_t handle;
+
+  expected_register_arguments.pointer = x;
+  expected_register_arguments.elements = 42;
+  expected_register_arguments.element_size = sizeof x[0];
+  starpu_vector_data_register (&handle, 0, (uintptr_t) x, 42, sizeof x[0]);
+
+  struct insert_task_argument expected[] =
+    {
+      { STARPU_VALUE, &size, sizeof size },
+      { STARPU_W, x },
+      { 0, 0, 0 }
+    };
+
+  expected_insert_task_arguments = expected;
+
+  /* Invoke the task, which makes sure it gets called with EXPECTED.  */
+  my_pointer_task (size, x);
+
+  assert (tasks_submitted == 1);
+
+  /* Again.  */
+  my_array_task (size, x);
+
+  assert (tasks_submitted == 2);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/pointer-tasks.c

@@ -0,0 +1,82 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions run (ldflags "-lstarpu-1.0")) */
+
+#undef NDEBUG
+
+#include <stdlib.h>
+#include <assert.h>
+#include <string.h>
+
+
+/* The task under test.  */
+
+static void my_pointer_task (const int *x, char a, long long *y, int b)
+  __attribute__ ((task));
+
+static int implementations_called;
+
+/* The input arguments.  */
+static const int pointer_arg1[] = { 42, 1, 2, 3, 4, 5 };
+static long long *pointer_arg2;
+
+/* CPU implementation of `my_pointer_task'.  */
+static void
+my_pointer_task (const int *x, char a, long long *y, int b)
+{
+  implementations_called |= STARPU_CPU;
+  assert (x == pointer_arg1);
+  assert (y == pointer_arg2);
+  assert (a == 'S');
+  assert (b == 42);
+}
+
+
+
+int
+main (int argc, char *argv[])
+{
+#define COUNT 100
+  pointer_arg2 = malloc (COUNT * sizeof *pointer_arg2);
+  memset (pointer_arg2, 0x77, COUNT * sizeof *pointer_arg2);
+
+#pragma starpu initialize
+
+  /* Register POINTER_ARG1 and POINTER_ARG2.  */
+#pragma starpu register pointer_arg1
+#pragma starpu register pointer_arg2 COUNT
+
+  /* Invoke the task, which should make sure it gets called with
+     EXPECTED.  */
+  my_pointer_task (pointer_arg1, 'S', pointer_arg2, 42);
+#pragma starpu wait
+  assert (implementations_called == STARPU_CPU);
+
+  implementations_called = 0;
+
+  /* Same, but with implicit integer type conversion.  */
+  my_pointer_task (pointer_arg1, (long long) 'S', pointer_arg2, (char) 42);
+#pragma starpu wait
+  assert (implementations_called == STARPU_CPU);
+
+  starpu_shutdown ();
+
+  free (pointer_arg2);
+
+  return EXIT_SUCCESS;
+#undef COUNT
+}

gcc-plugin/tests/pointers.c

@@ -0,0 +1,130 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+
+
+/* The tasks under test.  */
+
+static void my_pointer_task (const int *x, short *y) __attribute__ ((task));
+
+static void my_pointer_task_cpu (const int *, short *)
+  __attribute__ ((task_implementation ("cpu", my_pointer_task)));
+static void my_pointer_task_opencl (const int *, short *)
+  __attribute__ ((task_implementation ("opencl", my_pointer_task)));
+
+static void
+my_pointer_task_cpu (const int *x, short *y)
+{
+  printf ("%s: x = %p, y = %p\n", __func__, x, y);
+  assert (*x == 42 && *y == 77);
+}
+
+static void
+my_pointer_task_opencl (const int *x, short *y)
+{
+  printf ("%s: x = %p, y = %p\n", __func__, x, y);
+  assert (*x == 42 && *y == 77);
+}
+
+
+
+static void my_mixed_task (int *x, unsigned char z, const short *y)
+  __attribute__ ((task));
+static void my_mixed_task_cpu (int *, unsigned char, const short *)
+  __attribute__ ((task_implementation ("cpu", my_mixed_task)));
+static void my_mixed_task_opencl (int *, unsigned char, const short *)
+  __attribute__ ((task_implementation ("opencl", my_mixed_task)));
+
+static void
+my_mixed_task_cpu (int *x, unsigned char z, const short *y)
+{
+  printf ("%s: x = %p, y = %p, z = %i\n", __func__, x, y, (int) z);
+}
+
+static void
+my_mixed_task_opencl (int *x, unsigned char z, const short *y)
+{
+  printf ("%s: x = %p, y = %p, z = %i\n", __func__, x, y, (int) z);
+}
+
+
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  static const unsigned char z = 0x77;
+  static int x[] = { 42 };
+  short *y;
+
+  y = malloc (sizeof *y);
+  *y = 77;
+
+  /* Register X and Y (don't use the pragma, to avoid mixing concerns in this
+     test.)  */
+
+  starpu_data_handle_t handle;
+
+  expected_register_arguments.pointer = x;
+  expected_register_arguments.elements = 1;
+  expected_register_arguments.element_size = sizeof x[0];
+  starpu_vector_data_register (&handle, 0, (uintptr_t) x, 1, sizeof x[0]);
+
+  expected_register_arguments.pointer = y;
+  expected_register_arguments.elements = 1;
+  expected_register_arguments.element_size = sizeof *y;
+  starpu_vector_data_register (&handle, 0, (uintptr_t) y, 1, sizeof *y);
+
+  struct insert_task_argument expected_pointer_task[] =
+    {
+      { STARPU_R,  x },
+      { STARPU_RW, y },
+      { 0, 0, 0 }
+    };
+
+  expected_insert_task_arguments = expected_pointer_task;
+
+  /* Invoke the task, which should make sure it gets called with
+     EXPECTED.  */
+  my_pointer_task (x, y);
+
+  assert (tasks_submitted == 1);
+
+
+  /* Likewise with `my_mixed_task'.  */
+
+  struct insert_task_argument expected_mixed_task[] =
+    {
+      { STARPU_RW, x },
+      { STARPU_VALUE, &z, sizeof z },
+      { STARPU_R,  y },
+      { 0, 0, 0 }
+    };
+
+  expected_insert_task_arguments = expected_mixed_task;
+
+  my_mixed_task (x, 0x77, y);
+
+  assert (tasks_submitted == 2);
+
+  free (y);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/register-errors.c

@@ -0,0 +1,60 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Test error handling for `#pragma starpu register ...'.  */
+
+#undef NDEBUG
+
+extern void *void_pointer;
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+#pragma starpu register /* (error "parse error") */
+
+#pragma starpu register argv 234 junk here /* (error "junk after") *//* (error "unbound") *//* (error "unbound") */
+
+  static int x[123] __attribute__ ((unused));
+#pragma starpu register x 234 /* (note "can be omitted") *//* (error "differs from actual size") */
+
+  size_t x_size __attribute__ ((unused)) = sizeof x / sizeof x[0];
+#pragma starpu register x x_size /* (note "can be omitted") *//* (error "known at compile-time") */
+
+#pragma starpu register does_not_exit 123  /* (error "unbound variable") */
+#pragma starpu register argv does_not_exit /* (error "unbound variable") */
+
+#pragma starpu register argv /* (error "cannot determine size") */
+#pragma starpu register &argv[2] /* (error "cannot determine size") */
+#pragma starpu register &x[2] /* (error "cannot determine size") */
+
+#pragma starpu register argc /* (error "neither a pointer nor an array") */
+
+#pragma starpu register argv[2][3] 3 /* (error "neither a pointer nor an array") */
+
+#pragma starpu register argv[does_not_exist] 3 /* (error "unbound variable") */
+
+  char **p = argv;
+  size_t ps = argc;
+#pragma starpu register p ps  /* No unused variable warning, please! */
+
+#pragma starpu register void_pointer 123 /* (error "not allowed") */
+
+#pragma starpu register "hello"		   /* (error "invalid .*argument") */
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/register.c

@@ -0,0 +1,183 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Test whether `#pragma starpu register ...' generates the right code.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+
+static void
+foo (void)
+{
+  int x[] = { 1, 2, 3 };
+
+  expected_register_arguments.pointer = x;
+  expected_register_arguments.elements = sizeof x / sizeof x[0];
+  expected_register_arguments.element_size = sizeof x[0];
+#pragma starpu register x /* (warning "considered unsafe") */
+}
+
+static void
+bar (float *p, int s)
+{
+  expected_register_arguments.pointer = p;
+  expected_register_arguments.elements = s;
+  expected_register_arguments.element_size = sizeof *p;
+#pragma starpu register p s
+}
+
+/* Same as above, but with arguments reversed, to make sure using S doesn't
+   mutate the parameter list.  */
+static void
+baz (int s, float *p)
+{
+  expected_register_arguments.pointer = p;
+  expected_register_arguments.elements = s;
+  expected_register_arguments.element_size = sizeof *p;
+#pragma starpu register p s
+}
+
+/* Check the interaction between `register' and `heap_allocated'.  This test
+   assumes `heap_allocated' works as expected.  */
+
+static void
+heap_alloc (int x, int y)
+{
+  data_register_calls = data_unregister_calls = 0;
+
+  expected_malloc_argument = x * y * sizeof (float);
+
+  float m[x][y] __attribute__ ((heap_allocated));
+
+  expected_register_arguments.pointer = m;
+  expected_register_arguments.elements = x;
+  expected_register_arguments.element_size = y * sizeof m[0][0];
+#pragma starpu register m
+
+  expected_unregister_arguments.pointer = m;
+#pragma starpu unregister m
+
+  assert (data_register_calls == 1);
+  assert (data_unregister_calls == 1);
+
+  expected_free_argument = m;
+}
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  int x[123];
+  double *y;
+  static char z[345];
+  static float m[7][42];
+  static float m3d[14][11][80];
+  short w[] = { 1, 2, 3 };
+  size_t y_size = 234;
+
+  y = malloc (234 * sizeof *y);
+
+  expected_register_arguments.pointer = x;
+  expected_register_arguments.elements = 123;
+  expected_register_arguments.element_size = sizeof x[0];
+#pragma starpu register x 123 /* (note "can be omitted") */
+
+  expected_register_arguments.pointer = y;
+  expected_register_arguments.elements = 234;
+  expected_register_arguments.element_size = sizeof *y;
+#pragma starpu register y 234
+
+  expected_register_arguments.pointer = y;
+  expected_register_arguments.elements = y_size;
+  expected_register_arguments.element_size = sizeof *y;
+#pragma starpu register y y_size
+
+  expected_register_arguments.pointer = z;
+  expected_register_arguments.elements = 345;
+  expected_register_arguments.element_size = sizeof z[0];
+#pragma starpu register z
+
+  expected_register_arguments.pointer = w;
+  expected_register_arguments.elements = 3;
+  expected_register_arguments.element_size = sizeof w[0];
+#pragma starpu register w
+
+  expected_register_arguments.pointer = argv;
+  expected_register_arguments.elements = 456;
+  expected_register_arguments.element_size = sizeof argv[0];
+#pragma starpu register argv 456
+
+#define ARGV argv
+#define N 456
+  expected_register_arguments.pointer = argv;
+  expected_register_arguments.elements = N;
+  expected_register_arguments.element_size = sizeof argv[0];
+#pragma starpu register   ARGV /* hello, world! */  N
+#undef ARGV
+#undef N
+
+  foo ();
+  bar ((float *) argv, argc);
+  baz (argc, (float *) argv);
+
+  expected_register_arguments.pointer = argv;
+  expected_register_arguments.elements = argc;
+  expected_register_arguments.element_size = sizeof argv[0];
+
+  int chbouib = argc;
+#pragma starpu register argv chbouib
+
+  expected_register_arguments.pointer = &argv[2];
+  expected_register_arguments.elements = 3;
+  expected_register_arguments.element_size = sizeof argv[0];
+#pragma starpu register &argv[2] 3
+
+  expected_register_arguments.pointer = &argv[argc + 3 / 2];
+  expected_register_arguments.elements = argc * 4;
+  expected_register_arguments.element_size = sizeof argv[0];
+#pragma starpu register &argv[argc + 3 / 2] (argc * 4)
+
+  expected_register_arguments.pointer = &y[y_size / 2];
+  expected_register_arguments.elements = (y_size / 2 - 7);
+  expected_register_arguments.element_size = sizeof y[0];
+#pragma starpu register &y[y_size / 2] (y_size / 2 - 7)
+
+  expected_register_arguments.pointer = m[6];
+  expected_register_arguments.elements = 42;
+  expected_register_arguments.element_size = sizeof m[0][0];
+#pragma starpu register m[6]
+
+  expected_register_arguments.pointer = m;
+  expected_register_arguments.elements = 7;
+  expected_register_arguments.element_size = sizeof m[0];
+#pragma starpu register m
+
+  expected_register_arguments.pointer = m3d;
+  expected_register_arguments.elements = 14;
+  expected_register_arguments.element_size = sizeof m3d[0];
+#pragma starpu register m3d
+
+  assert (data_register_calls == 17);
+
+  free (y);
+
+  heap_alloc (42, 77);
+  assert (free_calls == 1);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/registered-errors.c

@@ -0,0 +1,36 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions compile (cflags "-Wno-unused-variable")) */
+
+static int global[123]              /* (error "cannot be used") */
+  __attribute__ ((registered, used));
+
+extern int external[123]            /* (error "cannot be used") */
+  __attribute__ ((registered));
+
+void
+foo (size_t size)
+{
+  float scalar /* (error "must have an array type") */
+    __attribute__ ((registered));
+  float *ptr   /* (error "must have an array type") */
+    __attribute__ ((registered));
+  float incomp[]  /* (error "incomplete array type") */
+    __attribute__ ((registered));
+  float incomp2[size][3][]  /* (error "incomplete element type") */
+    __attribute__ ((registered));
+}

gcc-plugin/tests/registered.c

@@ -0,0 +1,128 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+#include <stdlib.h>
+
+static void
+test_vec (void)
+{
+  data_register_calls = data_unregister_calls = 0;
+  expected_register_arguments.pointer = NULL;
+  expected_register_arguments.elements = 123;
+  expected_register_arguments.element_size = sizeof (float);
+
+  float vec[123]		    /* FIXME: warning: "considered unsafe" */
+    __attribute__ ((registered));
+
+  assert (data_register_calls == 1);
+  assert (expected_register_arguments.pointer == vec);
+
+  expected_unregister_arguments.pointer = vec;
+}
+
+static void
+test_matrix (void)
+{
+  data_register_calls = data_unregister_calls = 0;
+  expected_register_arguments.pointer = NULL;
+  expected_register_arguments.elements = 123;
+  expected_register_arguments.element_size = 234 * sizeof (double);
+
+  double matrix[123][234]	     /* FIXME: warning "considered unsafe" */
+    __attribute__ ((registered));
+
+  assert (data_register_calls == 1);
+  assert (expected_register_arguments.pointer == matrix);
+
+  expected_unregister_arguments.pointer = matrix;
+}
+
+static void
+test_with_heap_alloc (void)
+{
+  data_register_calls = data_unregister_calls = 0;
+  malloc_calls = free_calls = 0;
+
+  expected_register_arguments.pointer = NULL;
+  expected_register_arguments.elements = 123;
+  expected_register_arguments.element_size =
+     234 * 77 * sizeof (int);
+  expected_malloc_argument =
+    expected_register_arguments.elements
+    * expected_register_arguments.element_size;
+
+  int matrix[123][234][77]
+    __attribute__ ((registered, heap_allocated));
+
+  assert (data_register_calls == 1);
+  assert (expected_register_arguments.pointer == matrix);
+  assert (malloc_calls == 1);
+
+  expected_unregister_arguments.pointer = matrix;
+  expected_free_argument = matrix;
+}
+
+/* Same as above, but with the attributes in reverse order.  */
+static void
+test_with_heap_alloc_reversed (void)
+{
+  data_register_calls = data_unregister_calls = 0;
+  malloc_calls = free_calls = 0;
+
+  expected_register_arguments.pointer = NULL;
+  expected_register_arguments.elements = 123;
+  expected_register_arguments.element_size =
+     234 * 77 * sizeof (int);
+  expected_malloc_argument =
+    expected_register_arguments.elements
+    * expected_register_arguments.element_size;
+
+  int matrix[123][234][77]
+    __attribute__ ((heap_allocated, registered));
+
+  assert (data_register_calls == 1);
+  assert (expected_register_arguments.pointer == matrix);
+  assert (malloc_calls == 1);
+
+  expected_unregister_arguments.pointer = matrix;
+  expected_free_argument = matrix;
+}
+
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  test_vec ();
+  assert (data_unregister_calls == 1);
+
+  test_matrix ();
+  assert (data_unregister_calls == 1);
+
+  test_with_heap_alloc ();
+  assert (data_unregister_calls == 1);
+  assert (free_calls == 1);
+
+  test_with_heap_alloc_reversed ();
+  assert (data_unregister_calls == 1);
+  assert (free_calls == 1);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/release-errors.c

@@ -0,0 +1,42 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+extern float *y;
+
+static const double a[123];
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  int x[123] __attribute__ ((unused));
+  static char z[345] __attribute__ ((unused));
+
+#pragma starpu register x
+
+#pragma starpu release /* (error "parse error") */
+#pragma starpu release 123 /* (error "neither a pointer nor an array") */
+#pragma starpu release does_not_exit /* (error "unbound variable") */
+
+#pragma starpu release argc /* (error "neither a pointer nor an array") */
+#pragma starpu release y
+#pragma starpu release x
+
+#pragma starpu release x z			  /* (error "junk after") */
+
+  return 1;
+}

gcc-plugin/tests/release.c

@@ -0,0 +1,61 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Test whether `#pragma starpu release ...' generates the right code.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+
+static void
+foo (char *x, int foo)
+{
+  expected_release_arguments.pointer = x;
+#pragma starpu release x
+}
+
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  int x[123];
+  static char z[345];
+
+  expected_register_arguments.pointer = x;
+  expected_register_arguments.elements = 123;
+  expected_register_arguments.element_size = sizeof x[0];
+#pragma starpu register x
+
+  expected_release_arguments.pointer = x;
+#pragma starpu release x
+
+  expected_register_arguments.pointer = z;
+  expected_register_arguments.elements = sizeof z;
+  expected_register_arguments.element_size = sizeof z[0];
+#pragma starpu register z
+
+  expected_release_arguments.pointer = z;
+#pragma starpu release z
+
+  foo (z, 345);
+
+  assert (data_register_calls == 2);
+  assert (data_release_calls == 3);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/run-test.in

@@ -0,0 +1,462 @@
+#!/bin/sh
+# -*- mode: scheme; coding: utf-8; -*-
+GUILE_AUTO_COMPILE=0
+export GUILE_AUTO_COMPILE
+main='(@ (run-test) build/run)'
+exec "${GUILE-@GUILE@}" -l "$0"    \
+         -c "(apply $main (cdr (command-line)))" "$@"
+!#
+;;; GCC-StarPU
+;;; Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+;;;
+;;; GCC-StarPU is free software: you can redistribute it and/or modify
+;;; it under the terms of the GNU General Public License as published by
+;;; the Free Software Foundation, either version 3 of the License, or
+;;; (at your option) any later version.
+;;;
+;;; GCC-StarPU is distributed in the hope that it will be useful,
+;;; but WITHOUT ANY WARRANTY; without even the implied warranty of
+;;; MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+;;; GNU General Public License for more details.
+;;;
+;;; You should have received a copy of the GNU General Public License
+;;; along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.
+
+;;;
+;;; Written by Ludovic Courtès <ludovic.courtes@inria.fr>.
+;;;
+
+(define-module (run-test)
+  #:use-module (ice-9 regex)
+  #:use-module (ice-9 popen)
+  #:use-module (ice-9 rdelim)
+  #:use-module (ice-9 format)
+  #:use-module (ice-9 match)
+  #:use-module (srfi srfi-1)
+  #:use-module (srfi srfi-9)
+  #:use-module (srfi srfi-11)
+  #:use-module (srfi srfi-13)
+  #:use-module (srfi srfi-14)
+  #:use-module (srfi srfi-26)
+  #:export (build/run))
+
+;;; Commentary:
+;;;
+;;; Test machinery similar to the DejaGNU-based test framework used in GCC.
+;;; In a nutshell, this program compiles code with GCC and makes sure
+;;; warnings and errors are as appear in source code comments.
+;;;
+;;; This module should work with both Guile 1.8 and Guile 2.0.
+;;;
+;;; Code:
+
+;; Make sure the reader gets position information.
+(read-enable 'positions)
+
+(define %log-port
+  ;; Output port for debugging messages.
+  (current-output-port))
+
+(define (log fmt . args)
+  "Write an informational message."
+  (apply format %log-port (string-append fmt "\n") args))
+
+
+;;;
+;;; Compiling code.
+;;;
+
+(define %srcdir "@srcdir@")
+(define %builddir "@builddir@")
+(define %gcc "@CC@")
+
+(define %cuda-cppflags
+  (string-tokenize "@STARPU_CUDA_CPPFLAGS@"))
+
+(define %opencl-cppflags
+  (string-tokenize "@STARPU_OPENCL_CPPFLAGS@"))
+
+(define %default-cflags
+  `("-I" ,%srcdir
+    "-I" ,(string-append %srcdir "/../../src")    ; for <common/uthash.h>
+    "-I" ,(string-append %srcdir "/../../include")
+    "-I" ,(string-append %builddir "/../../include")
+    "-I" ,(string-append %builddir "/../..")
+
+    ,@%cuda-cppflags
+    ,@%opencl-cppflags
+
+    ;; Unfortunately `libtool --mode=execute' doesn't help here, so hard-code
+    ;; the real file name.
+    ,(string-append "-fplugin=" %builddir "/../src/.libs/starpu.so")
+
+    ;; Use the non-installed headers.
+    "-fplugin-arg-starpu-include-dir=@top_srcdir@/include"
+
+    ;; Find OpenCL source files under $srcdir.
+    ,(string-append "-fplugin-arg-starpu-opencl-include-dir=" %srcdir)
+
+    "-g"
+    "-fdump-tree-gimple" "-Wall"))
+
+(define %default-ldflags
+  `(,(string-append "-L" %builddir "/../../src")))
+
+(define %libtool
+  (string-append %builddir "/../../libtool"))
+
+
+(define (compile-starpu-code file cc cflags ldflags)
+  "Compile and link FILE with CC, using CFLAGS and LDFLAGS.  Return the
+compiler status and the list of lines printed on stdout/stderr."
+  (let* ((compile? (member "-c" cflags))
+         (ldflags  (if compile?
+                       (remove (cut string-prefix? "-L" <>) ldflags)
+                       ldflags))
+         (mode     (if compile?
+                       "compile"
+                       "link"))
+         (command  (format #f "LC_ALL=C ~a --mode=~a ~a ~{~a ~} \"~a\" ~{~a ~} 2>&1"
+                           %libtool mode cc cflags file ldflags))
+         (pipe     (begin
+                     (log "running `~a'" command)
+                     (open-input-pipe command))))
+    (let loop ((line    (read-line pipe))
+               (result '()))
+      (if (eof-object? line)
+          (values (close-pipe pipe) (reverse result))
+          (loop (read-line pipe)
+                (cons line result))))))
+
+(define (run-starpu-code executable)
+  "Run EXECUTABLE using Libtool; return its exit status."
+  (let* ((exe     (if (string-index executable #\/)
+                      executable
+                      (string-append (getcwd) "/" executable)))
+         (command (string-append %libtool " --mode=execute "
+                                 exe)))
+    (log "running `~a'" command)
+    (system command)))
+
+
+;;;
+;;; GCC diagnostics.
+;;;
+
+(define-record-type <location>
+  (make-location file line column)
+  location?
+  (file     location-file)
+  (line     location-line)
+  (column   location-column))
+
+(define (location=? loc1 loc2)
+  "Return #t if LOC1 and LOC2 refer roughly to the same file and line
+number."
+  (and (location-file loc1) (location-file loc2)
+       (string=? (basename (location-file loc1))
+                 (basename (location-file loc2)))
+       (= (location-line loc1) (location-line loc2))))
+
+(define-record-type <diagnostic>
+  (make-diagnostic location kind message)
+  diagnostic?
+  (location diagnostic-location)
+  (kind     diagnostic-kind)
+  (message  diagnostic-message))
+
+(define %diagnostic-with-location-rx
+  ;; "FILE:LINE:COL: KIND: MESSAGE..."
+  (make-regexp "^(.+):([[:digit:]]+):([[:digit:]]+): ([^:]+): (.*)$"))
+
+(define (string->diagnostic str)
+  "Parse STR and return the corresponding `diagnostic' object."
+  (cond ((regexp-exec %diagnostic-with-location-rx str)
+         =>
+         (lambda (m)
+           (let ((loc  (make-location (match:substring m 1)
+                                      (string->number (match:substring m 2))
+                                      (string->number (match:substring m 3))))
+                 (kind (string->symbol (match:substring m 4))))
+            (make-diagnostic loc kind (match:substring m 5)))))
+        (else
+         (make-diagnostic #f #f str))))
+
+
+;;;
+;;; Reading test directives.
+;;;
+
+(define (read-test-directives port)
+  "Read test directives from PORT.  Return a list of location/directive
+pairs."
+  (define (consume-whitespace p)
+    (let loop ((chr (peek-char p)))
+      (cond ((char-set-contains? char-set:whitespace chr)
+             (read-char p) ;; consume CHR
+             (loop (peek-char p)))
+            (else chr))))
+
+  (define (read-until-*/ p)
+    (let loop ((chr (read-char p)))
+      (cond ((eof-object? chr)
+             (error "unterminated C comment"))
+            ((eq? chr #\*)
+             (let ((next (peek-char p)))
+               (if (eq? next #\/)
+                   (read-char p) ;; consume CHR
+                   (loop (read-char p)))))
+            (else
+             (loop (read-char p))))))
+
+  (let loop ((chr        (read-char port))
+             (directives '()))
+    (cond ((eof-object? chr)
+           (reverse directives))
+          ((eq? chr #\/)
+           (let ((chr (read-char port)))
+             (if (eq? chr #\*)
+                 (let ((chr (consume-whitespace port)))
+                   (if (eq? chr #\()
+                       (let ((loc  (make-location (port-filename port)
+                                                  (1+ (port-line port))
+                                                  (port-column port)))
+                             (sexp (read port)))
+                         (read-until-*/ port)
+                         (loop (peek-char port)
+                               (cons (cons loc sexp)
+                                     directives)))
+                       (begin
+                         (read-until-*/ port)
+                         (loop (peek-char port) directives))))
+                 (loop chr directives))))
+          (else
+           (loop (read-char port) directives)))))
+
+(define (diagnostic-matches-directive? diagnostic directive location
+                                       cflags ldflags)
+  "Return #t if DIAGNOSTIC matches DIRECTIVE, which is at LOCATION."
+  (define optimizing?
+    (let ((opt (find (cut string-prefix? "-O" <>) cflags)))
+      (match opt
+        ((or #f "-O0") #f)
+        (_ #t))))
+
+  (let loop ((directive directive))
+    (match directive
+      (('if 'optimizing? directive)
+       (or (not optimizing?)
+           (loop directive)))
+      (('unless 'optimizing? directive)
+       (or optimizing?
+           (loop directive)))
+      ((kind message)
+       (and (eq? kind (diagnostic-kind diagnostic))
+            (location? (diagnostic-location diagnostic))
+            (location=? (diagnostic-location diagnostic) location)
+            (string-match message (diagnostic-message diagnostic)))))))
+
+
+;;;
+;;; Compiling and matching diagnostics against directives.
+;;;
+
+(define (compile/match* file directives cc cflags ldflags)
+  "Compile FILE and check whether GCC's diagnostics match DIRECTIVES.  Return
+3 values: the compiler's status code, the unmatched diagnostics, and the
+unsatisfied directives."
+  (let-values (((status diagnostics)
+                (compile-starpu-code file cc cflags ldflags)))
+    (let loop ((diagnostics (map string->diagnostic diagnostics))
+               (directives  directives)
+               (unsatisfied '()))
+      (if (null? directives)
+          (values status diagnostics unsatisfied)
+          (let* ((dir  (car directives))
+                 (diag (find (cute diagnostic-matches-directive?
+                                   <> (cdr dir) (car dir)
+                                   cflags ldflags)
+                             diagnostics)))
+            (if diag
+                (loop (delq diag diagnostics)
+                      (cdr directives)
+                      unsatisfied)
+                (loop diagnostics
+                      (cdr directives)
+                      (cons dir unsatisfied))))))))
+
+(define (executable-file source)
+  "Return the name of the executable file corresponding to SOURCE."
+  (let* ((dot (string-rindex source #\.))
+         (exe (if dot
+                  (substring source 0 dot)
+                  (string-append source ".exe")))
+         )
+  (if (string-prefix? %srcdir exe)
+      (string-append %builddir (substring exe (string-length %srcdir)))
+      exe
+      )))
+
+(define (compile/match file cc cflags ldflags)
+  "Read directives from FILE, and compiler/link/run it.  Make sure directives
+are matched, and report any errors otherwise.  Return #t on success and #f
+otherwise."
+  (define directives
+    (call-with-input-file file read-test-directives))
+
+  (define exe
+    (executable-file file))
+
+  (define (c->o c-file)
+    (string-append (substring c-file 0 (- (string-length c-file) 2))
+                   ".lo"))
+
+  (log "~a directives found in `~a'" (length directives) file)
+
+  (let*-values (((error-expected?)
+                 (find (lambda (l+d)
+                         (match l+d
+                           (((? location?) 'error _)
+                            #t)
+                           (_ #f)))
+                       directives))
+                ((instructions)
+                 (or (any (lambda (l+d)
+                            (match l+d
+                              (((? location?) 'instructions x ...)
+                               x)
+                              (_ #f)))
+                          directives)
+                     '(run)))
+                ((options)
+                 (match instructions
+                   ((_ options ...)
+                    options)
+                   (_ '())))
+                ((dependencies)
+                 (or (assq-ref options 'dependencies)
+                     '())))
+
+    (or (null? dependencies)
+        (format (current-output-port) "~s has ~a dependencies: ~{~s ~}~%"
+                file (length dependencies) dependencies))
+
+    (and (every (cut compile/match <> cc cflags ldflags)
+                (map (cut string-append %srcdir "/" <>) dependencies))
+         (let*-values (((goal)
+                        (if error-expected?
+                            'compile
+                            (car instructions)))
+                       ((cflags)
+                        `(,@cflags
+                          ,@(or (assq-ref options 'cflags) '())
+                          ,@(if (memq goal '(link run))
+                                `("-o" ,exe)
+                                '("-c"))))
+                       ((ldflags)
+                        `(,@(map c->o dependencies)
+                          ,@ldflags
+                          ,@(or (assq-ref options 'ldflags)
+                                '())))
+                       ((directives)
+                        (remove (lambda (l+d)
+                                  (match l+d
+                                    (((? location?) 'instructions _ ...)
+                                     #t)
+                                    (_ #f)))
+                                directives))
+                       ((status diagnostics unsatisfied)
+                        (compile/match* file directives cc cflags ldflags))
+                       ((unmatched)
+                        ;; Consider unmatched only diagnostics that have a
+                        ;; kind, to avoid taking into account messages like
+                        ;; "In file included from", "In function 'main'",
+                        ;; etc.
+                        (filter diagnostic-kind diagnostics)))
+
+           (or (null? unmatched)
+               (begin
+                 (format (current-error-port)
+                         "error: ~a unmatched GCC diagnostics:~%"
+                         (length unmatched))
+                 (for-each (lambda (d)
+                             (format (current-error-port)
+                                     "  ~a:~a:~a: ~a: ~a~%"
+                                     (and=> (diagnostic-location d)
+                                            location-file)
+                                     (and=> (diagnostic-location d)
+                                            location-line)
+                                     (and=> (diagnostic-location d)
+                                            location-column)
+                                     (diagnostic-kind d)
+                                     (diagnostic-message d)))
+                           unmatched)
+                 #f))
+
+           (if (null? unsatisfied)
+               (or (null? directives)
+                   (log "~a directives satisfied" (length directives)))
+               (begin
+                 (format (current-error-port)
+                         "error: ~a unsatisfied directives:~%"
+                         (length unsatisfied))
+                 (for-each (lambda (l+d)
+                             (let ((loc (car l+d))
+                                   (dir (cdr l+d)))
+                               (format (current-error-port)
+                                       "  ~a:~a:~a: ~a: ~s~%"
+                                       (location-file loc)
+                                       (location-line loc)
+                                       (location-column loc)
+                                       (car dir)
+                                       (cadr dir))))
+                           unsatisfied)
+                 #f))
+
+           (if error-expected?
+               (if (= 0 status)
+                   (format (current-error-port)
+                           "error: compilation succeeded~%"))
+               (if (= 0 status)
+                   (or (eq? goal 'compile)
+                       (file-exists? exe)
+                       (begin
+                         (format (current-error-port)
+                                 "error: executable file `~a' not found~%" exe)
+                         #f))
+                   (format (current-error-port)
+                           "error: compilation failed (compiler exit code ~a)~%~{  ~a~%~}"
+                           status
+                           (map diagnostic-message diagnostics))))
+
+           (and (null? unmatched)
+                (null? unsatisfied)
+                (if error-expected?
+                    (not (= 0 status))
+                    (and (= 0 status)
+                         (or (eq? goal 'compile) (file-exists? exe))
+                         (or (not (eq? goal 'run))
+                             (let ((status (run-starpu-code exe)))
+                               (or (= 0 status)
+                                   (begin
+                                     (format (current-error-port)
+                                             "error: program `~a' failed \
+(exit code ~a)~%"
+                                             exe status)
+                                     #f)))))))))))
+
+
+;;;
+;;; Entry point.
+;;;
+
+(define (build/run . file)
+  (exit (every (lambda (file)
+                 ;; For each file, check that everything works both with and
+                 ;; without optimizations.
+                 (every (cut compile/match file %gcc <> %default-ldflags)
+                        `((,"-O0" ,@%default-cflags)
+                          (,"-O2" ,@%default-cflags))))
+               file)))
+
+;;; run-test.in ends here

gcc-plugin/tests/scalar-tasks.c

@@ -0,0 +1,71 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions run (ldflags "-lstarpu-1.0")) */
+
+#undef NDEBUG
+
+#include <stdlib.h>
+#include <assert.h>
+
+
+/* The task under test.  */
+
+static void my_scalar_task (int x, int y) __attribute__ ((task));
+
+static void my_scalar_task_opencl (int, int)
+  __attribute__ ((task_implementation ("opencl", my_scalar_task), noinline));
+
+static int implementations_called;
+
+/* CPU implementation of `my_scalar_task'.  */
+static void
+my_scalar_task (int x, int y)
+{
+  implementations_called |= STARPU_CPU;
+  assert (x == 42);
+  assert (y == 77);
+}
+
+static void
+my_scalar_task_opencl (int x, int y)
+{
+  implementations_called |= STARPU_OPENCL;
+  assert (x == 42);
+  assert (y == 77);
+}
+
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  /* Invoke the task, which should make sure it gets called with
+     EXPECTED.  */
+  my_scalar_task (42, 77);
+
+#pragma starpu wait
+
+  assert ((implementations_called & STARPU_CPU)
+	  || (implementations_called & STARPU_OPENCL));
+
+  assert ((implementations_called & ~(STARPU_CPU | STARPU_OPENCL)) == 0);
+
+  starpu_shutdown ();
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/shutdown-errors.c

@@ -0,0 +1,23 @@
+/* GCC-StarPU
+   Copyright (C) 2011 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+#pragma starpu shutdown foo bar /* (error "junk after") */
+  return 0;
+}

gcc-plugin/tests/task-errors.c

@@ -0,0 +1,142 @@
+/* GCC-StarPU
+   Copyright (C) 2011, 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Test error handling for the `task' and `task_implementation' attributes.  */
+
+extern void my_external_task (int foo, char *bar) __attribute__ ((task));
+
+void my_task (int foo, char *bar) /* (error "none of the implementations") */
+  __attribute__ ((task));
+static void my_task_cpu (int foo, float *bar)    /* (error "type differs") */
+  __attribute__ ((task_implementation ("cpu", my_task)));
+
+static void my_task_opencl (long foo, char *bar) /* (error "type differs") */
+  __attribute__ ((task_implementation ("opencl", my_task)));
+
+static void my_task_nowhere (int foo, char *bar) /* (warning "unsupported target") */
+  __attribute__ ((task_implementation ("does-not-exist", my_task)));
+
+static void my_task_not_quite (int foo, char *bar) /* (error "lacks the 'task' attribute") */
+  __attribute__ ((task_implementation ("cpu", my_task_nowhere)));
+
+static int foo /* (error "only applies to function") */
+  __attribute__ ((task_implementation ("cpu", my_task)));
+
+static int bar /* (error "only applies to function") */
+  __attribute__ ((task, unused));
+
+static int not_a_task __attribute__ ((unused));
+
+static void my_task_almost (int foo, char *bar)    /* (error "not a function") */
+  __attribute__ ((task_implementation ("cpu", not_a_task)));
+
+static void my_task_wrong_task_arg (int foo, char *bar)   /* (error "not a function") */
+  __attribute__ ((task_implementation ("cpu", 123)));
+
+static void my_task_wrong_target_arg (int foo, char *bar) /* (error "string constant expected") */
+  __attribute__ ((task_implementation (123, my_task)));
+
+extern int my_task_not_void (int foo) /* (error "return type") */
+  __attribute__ ((task));
+
+void my_task_that_invokes_task (int x, char *y)
+  __attribute__ ((task));
+
+void my_task_that_invokes_task_cpu (int x, char *y)
+  __attribute__ ((task_implementation ("cpu", my_task_that_invokes_task)));
+
+/* XXX: The assumption behind this test is that STARPU_USE_GORDON is not
+   defined.  */
+void my_task_with_no_usable_implementation (int x) /* (error "none of the implementations") */
+  __attribute__ ((task));
+
+static void my_task_with_no_usable_implementation_gordon (int x)
+  __attribute__ ((task_implementation ("gordon",
+				       my_task_with_no_usable_implementation)));
+
+/* XXX: In practice this test fails for large values of `STARPU_NMAXBUFS'.  */
+void my_task_with_too_many_pointer_params (/* (error "maximum .* exceeded") */
+					   char *x1, char *x2, char *x3,
+					   char *x4, char *x5, char *x6,
+					   char *x7, char *x8, char *x9,
+					   char *xa, char *xb, char *xc,
+					   char *xd, char *xe, char *xf,
+					   char *xg, char *xh, char *xi)
+  __attribute__ ((task));
+
+
+static void my_task_without_any_parameters (void)
+  __attribute__ ((task));
+
+static void my_task_without_any_parameters_gordon (void)
+  __attribute__ ((task_implementation ("gordon", my_task_without_any_parameters)));
+
+void
+my_task_without_any_parameters (void)
+{
+}
+
+void
+my_task_without_any_parameters_gordon (void)
+{
+}
+
+
+static void
+my_task_cpu (int foo, float *bar)
+{
+}
+
+static void
+my_task_opencl (long foo, char *bar)
+{
+}
+
+static void
+my_task_nowhere (int foo, char *bar)
+{
+}
+
+static void
+my_task_not_quite (int foo, char *bar)
+{
+}
+
+static void
+my_task_almost (int foo, char *bar)
+{
+}
+
+static void
+my_task_wrong_task_arg (int foo, char *bar)
+{
+}
+
+static void
+my_task_wrong_target_arg (int foo, char *bar)
+{
+}
+
+void
+my_task_that_invokes_task_cpu (int x, char *y)
+{
+  my_external_task (x, y); /* (error "cannot be invoked from task implementation") */
+}
+
+static void
+my_task_with_no_usable_implementation_gordon (int x)
+{
+}

gcc-plugin/tests/test.cl

@@ -0,0 +1 @@
+/* This is an almost empty OpenCL file.  */

gcc-plugin/tests/unregister-errors.c

@@ -0,0 +1,47 @@
+/* GCC-StarPU
+   Copyright (C) 2011 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+extern float *y;
+
+static const double a[123];
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  int x[123] __attribute__ ((unused));
+  static char z[345] __attribute__ ((unused));
+
+#pragma starpu register x
+
+#pragma starpu unregister /* (error "parse error") */
+#pragma starpu unregister 123 /* (error "neither a pointer nor an array") */
+#pragma starpu unregister does_not_exit /* (error "unbound variable") */
+
+#pragma starpu unregister argc /* (error "neither a pointer nor an array") */
+#pragma starpu unregister y
+#pragma starpu unregister x
+
+  /* XXX: Uncomment below when this is supported.  */
+#if 0
+#pragma starpu unregister z /* error "not registered" */
+#pragma starpu unregister a /* error "not registered" */
+#pragma starpu unregister argv /* error "not registered" */
+#endif
+
+  return 1;
+}

gcc-plugin/tests/unregister.c

@@ -0,0 +1,51 @@
+/* GCC-StarPU
+   Copyright (C) 2011 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Test whether `#pragma starpu unregister ...' generates the right code.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+
+int
+main (int argc, char *argv[])
+{
+#pragma starpu initialize
+
+  int x[123];
+  static char z[345];
+
+  expected_register_arguments.pointer = x;
+  expected_register_arguments.elements = 123;
+  expected_register_arguments.element_size = sizeof x[0];
+#pragma starpu register x
+
+  expected_unregister_arguments.pointer = x;
+#pragma starpu unregister x
+
+  expected_register_arguments.pointer = z;
+  expected_register_arguments.elements = sizeof z;
+  expected_register_arguments.element_size = sizeof z[0];
+#pragma starpu register z
+
+  expected_unregister_arguments.pointer = z;
+#pragma starpu unregister z
+
+  assert (data_register_calls == 2);
+  assert (data_unregister_calls == 2);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/verbose.c

@@ -0,0 +1,46 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions compile (cflags "-fplugin-arg-starpu-verbose")) */
+
+#ifndef STARPU_GCC_PLUGIN
+# error barf!
+#endif
+
+static void my_task (int x)	      /* (note "implementations for task") */
+  __attribute__ ((task));
+
+static void my_task_cpu (int x)
+  __attribute__ ((task_implementation ("cpu", my_task)));
+extern void my_task_cpu_sse (int x)	       /* (note "my_task_cpu_sse") */
+  __attribute__ ((task_implementation ("cpu", my_task)));
+extern void my_task_opencl (int x)		/* (note "my_task_opencl") */
+  __attribute__ ((task_implementation ("opencl", my_task)));
+extern void my_task_cuda (int x)		  /* (note "my_task_cuda") */
+  __attribute__ ((task_implementation ("cuda", my_task)));
+
+static void
+my_task_cpu (int x)				  /* (note "my_task_cpu") */
+{
+  /* Nothing.  */
+}
+
+int
+bar (int x)
+{
+  my_task (x);					  /* (note "calls task") */
+  return 42;
+}

gcc-plugin/tests/wait-errors.c

@@ -0,0 +1,27 @@
+/* GCC-StarPU
+   Copyright (C) 2011 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* The task under test.  */
+
+void task (int x, char y, int z) __attribute__ ((task));
+static void task_cpu (int x, char y, int z)
+  __attribute__ ((task_implementation ("cpu", task)));
+
+static void
+task_cpu (int x, char y, int z)
+{
+#pragma starpu wait /* (error "not allowed") */
+}

gcc-plugin/tests/warn-unregistered.c

@@ -0,0 +1,211 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-StarPU is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* (instructions compile) */
+
+/* Make sure warnings get raised when pointer variables are likely never
+   registered.  */
+
+extern void my_task (double *x, double *y)
+  __attribute__ ((task));
+
+extern double *global1;
+double global2[123];
+
+void
+parm_decl_no_warn (double *parm1, double *parm2)
+{
+  my_task (parm1, parm2);	  /* no warning, because these are parameters
+				     so we cannot tell anything */
+}
+
+void
+global_decl_no_warn (void)
+{
+  my_task (global1, global2);
+}
+
+void
+two_unregistered_pointers (void)
+{
+  double *p, *q;
+
+  p = malloc (12 * sizeof *p);
+  q = malloc (23 * sizeof *q);
+
+  my_task (p, q); /* (warning "p.* used unregistered") *//* (warning "q.* used unregistered") */
+}
+
+void
+one_unregistered_pointer (void)
+{
+  double *p, *q;
+
+  p = malloc (12 * sizeof *p);
+  q = malloc (23 * sizeof *q);
+
+#pragma starpu register p 12
+  my_task (p, q);		      /* (if optimizing?
+					     (warning "q.* used unregistered")) */
+}
+
+void
+another_unregistered_pointer (void)
+{
+  double X[] = { 1, 2, 3, 4 };
+  double *Y;
+
+  Y = malloc (123 * sizeof *Y);
+  if (Y == NULL)
+    return;
+  else
+    {
+      extern void frob (double *);
+      frob (Y);
+    }
+  X[0] = 42;
+
+#pragma starpu register Y 123
+  my_task (X, Y);		      /* (warning "X.* used unregistered") */
+}
+
+void
+zero_unregistered_pointers (void)
+{
+  double *p, *q;
+
+  p = malloc (12 * sizeof *p);
+  q = malloc (23 * sizeof *q);
+
+#pragma starpu register p 12
+#pragma starpu register q 23
+  my_task (p, q);				  /* no warning */
+}
+
+void
+two_pointers_unregistered_before_call (void)
+{
+  double *p, *q;
+
+  p = malloc (12 * sizeof *p);
+  q = malloc (23 * sizeof *q);
+
+  my_task (p, q); /* (warning "p.* used unregistered") *//* (warning "q.* used unregistered") */
+
+#pragma starpu register p 12
+#pragma starpu register q 23
+}
+
+void
+one_unregistered_array (void)
+{
+  double PPP[12], QQQ[23];
+
+#pragma starpu register PPP	      /* (warning "on-stack .* unsafe") */
+  my_task (PPP, QQQ);		      /* (warning "QQQ.* used unregistered") */
+}
+
+void
+not_the_ones_registered (void)
+{
+  double a[12], b[23], p[12], q[23];
+
+#pragma starpu register a		 /* (warning "on-stack .* unsafe") */
+#pragma starpu register b		 /* (warning "on-stack .* unsafe") */
+
+  my_task (p, q); /* (warning "p.* used unregistered") */ /* (warning "q.* used unregistered") */
+}
+
+void
+registered_pointers_with_aliases (void)
+{
+  double *a, *b, *p, *q;
+
+  a = malloc (123 * sizeof *a);
+  b = malloc (234 * sizeof *b);
+
+#pragma starpu register a 123
+#pragma starpu register b 234
+  p = a;
+  q = b;
+  my_task (p, q);				  /* no warning */
+}
+
+void
+one_unregistered_array_attrs (void)
+{
+  double p[12];
+  double q[23] __attribute__ ((heap_allocated, registered));
+
+  my_task (p, q);		      /* (warning "p.* used unregistered") */
+}
+
+void
+unregistered_on_one_path (int x)
+{
+  double p[12], q[34];
+
+  if (x > 42)
+    {
+#pragma starpu register p	      /* (warning "on-stack .* unsafe") */
+    }
+
+#pragma starpu register q	      /* (warning "on-stack .* unsafe") */
+
+  my_task (p, q);		      /* (warning "p.* used unregistered") */
+}
+
+void
+registered_via_two_paths (int x)
+{
+  double p[12], q[34];
+
+  if (x > 42)
+    {
+#pragma starpu register p	      /* (warning "on-stack .* unsafe") */
+    }
+  else
+    {
+#pragma starpu register p	      /* (warning "on-stack .* unsafe") */
+    }
+
+#pragma starpu register q	      /* (warning "on-stack .* unsafe") */
+
+  my_task (p, q);				  /* no warning */
+}
+
+#if 0
+
+/* FIXME: This case currently triggers a false positives.  */
+
+void
+registered_and_used_in_loop (void)
+{
+  int i;
+  double *p[123];
+  static double q[234];
+
+  for (i = 0; i < 123; i++)
+    {
+      p[i] = malloc (123 * sizeof *p[i]);
+#pragma starpu register p[i] 123
+    }
+
+  for (i = 0; i < 123; i++)
+    my_task (p[i], q);
+}
+
+#endif