Merge SOCL into StarPU repo

Makefile.am

@@ -26,6 +26,10 @@ if COND_OPT
 SUBDIRS += tests/opt examples/opt
 endif
 
+if BUILD_SOCL
+SUBDIRS += socl
+endif
+
 if BUILD_GCC_PLUGIN
 SUBDIRS += gcc-plugin
 endif

configure.ac

@@ -1109,6 +1109,31 @@ AM_CONDITIONAL([HAVE_GUILE], [test "x$GUILE" != "x"])
 
 ###############################################################################
 #                                                                             #
+#                               OpenCL interface                              #
+#                                                                             #
+###############################################################################
+
+AC_ARG_ENABLE([socl],
+  [AS_HELP_STRING([--enable-socl],
+    [build the OpenCL interfacce (SOCL)])],
+  [enable_socl="$enableval"],
+  [enable_socl="no"])
+
+#TODO: check that OpenCL is enabled
+
+if test "x$enable_socl" = "xyes"; then
+   STARPU_SOCL_SUPPORT
+
+   build_socl="yes"
+else
+   build_socl="no"
+   run_socl_test_suite="no"
+fi
+
+AM_CONDITIONAL([BUILD_SOCL], [test "x$build_socl" = "xyes"])
+
+###############################################################################
+#                                                                             #
 #                                  Examples                                   #
 #                                                                             #
 ###############################################################################
@@ -1345,6 +1370,8 @@ AC_OUTPUT([
 	gcc-plugin/src/Makefile
 	gcc-plugin/tests/Makefile
 	gcc-plugin/tests/run-test
+	socl/Makefile
+	socl/src/Makefile
 	libstarpu.pc
 	examples/Makefile
         examples/opt/Makefile
@@ -1381,6 +1408,7 @@ AC_MSG_NOTICE([
 	Allocation cache:  $enable_allocation_cache
 
 	MPI enabled:   $use_mpi
+	SOCL enabled:  $build_socl
 	Magma enabled: $have_magma
 	BLAS library:  $blas_lib
 	hwloc:         $have_valid_hwloc

socl/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>.

socl/Makefile.am

@@ -0,0 +1,30 @@
+# StarPU --- Runtime system for heterogeneous multicore architectures.
+#
+# Copyright (C) 2011 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
+
+EXTRA_DIST = COPYING README
+
+libsocl_la_includedir=$(includedir)/starpu/CL
+
+libsocl_la_include_HEADERS = \
+  include/CL/cl.h \
+  include/CL/cl_ext.h \
+  include/CL/cl_gl.h \
+  include/CL/cl_gl_ext.h \
+  include/CL/cl_platform.h \
+  include/CL/cl_starpu.h \
+  include/CL/opencl.h \
+  include/CL/cl.hpp

socl/README

@@ -0,0 +1,7 @@
+StarPU's OpenCL interface
+=========================
+
+This directory contains an OpenCL implementation that can
+be used as a replacement of the classic StarPU's API.
+
+OpenCL applications need to be compiled using provided headers.

socl/include/CL/cl.h

@@ -0,0 +1,879 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2009 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+/* $Revision: 10424 $ on $Date: 2010-02-17 14:34:49 -0800 (Wed, 17 Feb 2010) $ */
+
+#ifndef __OPENCL_CL_H
+#define __OPENCL_CL_H
+
+#ifdef __APPLE__
+#include <OpenCL/cl_platform.h>
+#else
+#include <CL/cl_platform.h>
+#endif	
+
+#include "CL/cl_starpu.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/******************************************************************************/
+
+typedef struct _cl_platform_id *    cl_platform_id;
+typedef struct _cl_device_id *      cl_device_id;
+typedef struct _cl_context *        cl_context;
+typedef struct _cl_command_queue *  cl_command_queue;
+typedef struct _cl_mem *            cl_mem;
+typedef struct _cl_program *        cl_program;
+typedef struct _cl_kernel *         cl_kernel;
+typedef struct _cl_event *          cl_event;
+typedef struct _cl_sampler *        cl_sampler;
+
+typedef cl_uint             cl_bool;                     /* WARNING!  Unlike cl_ types in cl_platform.h, cl_bool is not guaranteed to be the same size as the bool in kernels. */ 
+typedef cl_ulong            cl_bitfield;
+typedef cl_bitfield         cl_device_type;
+typedef cl_uint             cl_platform_info;
+typedef cl_uint             cl_device_info;
+typedef cl_bitfield         cl_device_address_info;
+typedef cl_bitfield         cl_device_fp_config;
+typedef cl_uint             cl_device_mem_cache_type;
+typedef cl_uint             cl_device_local_mem_type;
+typedef cl_bitfield         cl_device_exec_capabilities;
+typedef cl_bitfield         cl_command_queue_properties;
+
+typedef intptr_t			cl_context_properties;
+typedef cl_uint             cl_context_info;
+typedef cl_uint             cl_command_queue_info;
+typedef cl_uint             cl_channel_order;
+typedef cl_uint             cl_channel_type;
+typedef cl_bitfield         cl_mem_flags;
+typedef cl_uint             cl_mem_object_type;
+typedef cl_uint             cl_mem_info;
+typedef cl_uint             cl_image_info;
+typedef cl_uint             cl_addressing_mode;
+typedef cl_uint             cl_filter_mode;
+typedef cl_uint             cl_sampler_info;
+typedef cl_bitfield         cl_map_flags;
+typedef cl_uint             cl_program_info;
+typedef cl_uint             cl_program_build_info;
+typedef cl_int              cl_build_status;
+typedef cl_uint             cl_kernel_info;
+typedef cl_uint             cl_kernel_work_group_info;
+typedef cl_uint             cl_event_info;
+typedef cl_uint             cl_command_type;
+typedef cl_uint             cl_profiling_info;
+
+typedef struct _cl_image_format {
+    cl_channel_order        image_channel_order;
+    cl_channel_type         image_channel_data_type;
+} cl_image_format;
+
+
+
+/******************************************************************************/
+
+/* Error Codes */
+#define CL_SUCCESS                                  0
+#define CL_DEVICE_NOT_FOUND                         -1
+#define CL_DEVICE_NOT_AVAILABLE                     -2
+#define CL_COMPILER_NOT_AVAILABLE                   -3
+#define CL_MEM_OBJECT_ALLOCATION_FAILURE            -4
+#define CL_OUT_OF_RESOURCES                         -5
+#define CL_OUT_OF_HOST_MEMORY                       -6
+#define CL_PROFILING_INFO_NOT_AVAILABLE             -7
+#define CL_MEM_COPY_OVERLAP                         -8
+#define CL_IMAGE_FORMAT_MISMATCH                    -9
+#define CL_IMAGE_FORMAT_NOT_SUPPORTED               -10
+#define CL_BUILD_PROGRAM_FAILURE                    -11
+#define CL_MAP_FAILURE                              -12
+
+#define CL_INVALID_VALUE                            -30
+#define CL_INVALID_DEVICE_TYPE                      -31
+#define CL_INVALID_PLATFORM                         -32
+#define CL_INVALID_DEVICE                           -33
+#define CL_INVALID_CONTEXT                          -34
+#define CL_INVALID_QUEUE_PROPERTIES                 -35
+#define CL_INVALID_COMMAND_QUEUE                    -36
+#define CL_INVALID_HOST_PTR                         -37
+#define CL_INVALID_MEM_OBJECT                       -38
+#define CL_INVALID_IMAGE_FORMAT_DESCRIPTOR          -39
+#define CL_INVALID_IMAGE_SIZE                       -40
+#define CL_INVALID_SAMPLER                          -41
+#define CL_INVALID_BINARY                           -42
+#define CL_INVALID_BUILD_OPTIONS                    -43
+#define CL_INVALID_PROGRAM                          -44
+#define CL_INVALID_PROGRAM_EXECUTABLE               -45
+#define CL_INVALID_KERNEL_NAME                      -46
+#define CL_INVALID_KERNEL_DEFINITION                -47
+#define CL_INVALID_KERNEL                           -48
+#define CL_INVALID_ARG_INDEX                        -49
+#define CL_INVALID_ARG_VALUE                        -50
+#define CL_INVALID_ARG_SIZE                         -51
+#define CL_INVALID_KERNEL_ARGS                      -52
+#define CL_INVALID_WORK_DIMENSION                   -53
+#define CL_INVALID_WORK_GROUP_SIZE                  -54
+#define CL_INVALID_WORK_ITEM_SIZE                   -55
+#define CL_INVALID_GLOBAL_OFFSET                    -56
+#define CL_INVALID_EVENT_WAIT_LIST                  -57
+#define CL_INVALID_EVENT                            -58
+#define CL_INVALID_OPERATION                        -59
+#define CL_INVALID_GL_OBJECT                        -60
+#define CL_INVALID_BUFFER_SIZE                      -61
+#define CL_INVALID_MIP_LEVEL                        -62
+#define CL_INVALID_GLOBAL_WORK_SIZE                 -63
+
+/* OpenCL Version */
+#define CL_VERSION_1_0                              1
+
+/* cl_bool */
+#define CL_FALSE                                    0
+#define CL_TRUE                                     1
+
+/* cl_platform_info */
+#define CL_PLATFORM_PROFILE                         0x0900
+#define CL_PLATFORM_VERSION                         0x0901
+#define CL_PLATFORM_NAME                            0x0902
+#define CL_PLATFORM_VENDOR                          0x0903
+#define CL_PLATFORM_EXTENSIONS                      0x0904
+
+/* cl_device_type - bitfield */
+#define CL_DEVICE_TYPE_DEFAULT                      (1 << 0)
+#define CL_DEVICE_TYPE_CPU                          (1 << 1)
+#define CL_DEVICE_TYPE_GPU                          (1 << 2)
+#define CL_DEVICE_TYPE_ACCELERATOR                  (1 << 3)
+#define CL_DEVICE_TYPE_ALL                          0xFFFFFFFF
+
+/* cl_device_info */
+#define CL_DEVICE_TYPE                              0x1000
+#define CL_DEVICE_VENDOR_ID                         0x1001
+#define CL_DEVICE_MAX_COMPUTE_UNITS                 0x1002
+#define CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS          0x1003
+#define CL_DEVICE_MAX_WORK_GROUP_SIZE               0x1004
+#define CL_DEVICE_MAX_WORK_ITEM_SIZES               0x1005
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR       0x1006
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT      0x1007
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT        0x1008
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG       0x1009
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT      0x100A
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE     0x100B
+#define CL_DEVICE_MAX_CLOCK_FREQUENCY               0x100C
+#define CL_DEVICE_ADDRESS_BITS                      0x100D
+#define CL_DEVICE_MAX_READ_IMAGE_ARGS               0x100E
+#define CL_DEVICE_MAX_WRITE_IMAGE_ARGS              0x100F
+#define CL_DEVICE_MAX_MEM_ALLOC_SIZE                0x1010
+#define CL_DEVICE_IMAGE2D_MAX_WIDTH                 0x1011
+#define CL_DEVICE_IMAGE2D_MAX_HEIGHT                0x1012
+#define CL_DEVICE_IMAGE3D_MAX_WIDTH                 0x1013
+#define CL_DEVICE_IMAGE3D_MAX_HEIGHT                0x1014
+#define CL_DEVICE_IMAGE3D_MAX_DEPTH                 0x1015
+#define CL_DEVICE_IMAGE_SUPPORT                     0x1016
+#define CL_DEVICE_MAX_PARAMETER_SIZE                0x1017
+#define CL_DEVICE_MAX_SAMPLERS                      0x1018
+#define CL_DEVICE_MEM_BASE_ADDR_ALIGN               0x1019
+#define CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE          0x101A
+#define CL_DEVICE_SINGLE_FP_CONFIG                  0x101B
+#define CL_DEVICE_GLOBAL_MEM_CACHE_TYPE             0x101C
+#define CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE         0x101D
+#define CL_DEVICE_GLOBAL_MEM_CACHE_SIZE             0x101E
+#define CL_DEVICE_GLOBAL_MEM_SIZE                   0x101F
+#define CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE          0x1020
+#define CL_DEVICE_MAX_CONSTANT_ARGS                 0x1021
+#define CL_DEVICE_LOCAL_MEM_TYPE                    0x1022
+#define CL_DEVICE_LOCAL_MEM_SIZE                    0x1023
+#define CL_DEVICE_ERROR_CORRECTION_SUPPORT          0x1024
+#define CL_DEVICE_PROFILING_TIMER_RESOLUTION        0x1025
+#define CL_DEVICE_ENDIAN_LITTLE                     0x1026
+#define CL_DEVICE_AVAILABLE                         0x1027
+#define CL_DEVICE_COMPILER_AVAILABLE                0x1028
+#define CL_DEVICE_EXECUTION_CAPABILITIES            0x1029
+#define CL_DEVICE_QUEUE_PROPERTIES                  0x102A
+#define CL_DEVICE_NAME                              0x102B
+#define CL_DEVICE_VENDOR                            0x102C
+#define CL_DRIVER_VERSION                           0x102D
+#define CL_DEVICE_PROFILE                           0x102E
+#define CL_DEVICE_VERSION                           0x102F
+#define CL_DEVICE_EXTENSIONS                        0x1030
+#define CL_DEVICE_PLATFORM                          0x1031
+/* 0x1032 reserved for CL_DEVICE_DOUBLE_FP_CONFIG */
+/* 0x1033 reserved for CL_DEVICE_HALF_FP_CONFIG */
+
+/* cl_device_fp_config - bitfield */
+#define CL_FP_DENORM                                (1 << 0)
+#define CL_FP_INF_NAN                               (1 << 1)
+#define CL_FP_ROUND_TO_NEAREST                      (1 << 2)
+#define CL_FP_ROUND_TO_ZERO                         (1 << 3)
+#define CL_FP_ROUND_TO_INF                          (1 << 4)
+#define CL_FP_FMA                                   (1 << 5)
+
+/* cl_device_mem_cache_type */
+#define CL_NONE                                     0x0
+#define CL_READ_ONLY_CACHE                          0x1
+#define CL_READ_WRITE_CACHE                         0x2
+
+/* cl_device_local_mem_type */
+#define CL_LOCAL                                    0x1
+#define CL_GLOBAL                                   0x2
+
+/* cl_device_exec_capabilities - bitfield */
+#define CL_EXEC_KERNEL                              (1 << 0)
+#define CL_EXEC_NATIVE_KERNEL                       (1 << 1)
+
+/* cl_command_queue_properties - bitfield */
+#define CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE      (1 << 0)
+#define CL_QUEUE_PROFILING_ENABLE                   (1 << 1)
+
+/* cl_context_info  */
+#define CL_CONTEXT_REFERENCE_COUNT                  0x1080
+#define CL_CONTEXT_DEVICES                          0x1081
+#define CL_CONTEXT_PROPERTIES                       0x1082
+
+/* cl_context_info + cl_context_properties */
+#define CL_CONTEXT_PLATFORM                         0x1084
+
+/* cl_command_queue_info */
+#define CL_QUEUE_CONTEXT                            0x1090
+#define CL_QUEUE_DEVICE                             0x1091
+#define CL_QUEUE_REFERENCE_COUNT                    0x1092
+#define CL_QUEUE_PROPERTIES                         0x1093
+
+/* cl_mem_flags - bitfield */
+#define CL_MEM_READ_WRITE                           (1 << 0)
+#define CL_MEM_WRITE_ONLY                           (1 << 1)
+#define CL_MEM_READ_ONLY                            (1 << 2)
+#define CL_MEM_USE_HOST_PTR                         (1 << 3)
+#define CL_MEM_ALLOC_HOST_PTR                       (1 << 4)
+#define CL_MEM_COPY_HOST_PTR                        (1 << 5)
+
+/* cl_channel_order */
+#define CL_R                                        0x10B0
+#define CL_A                                        0x10B1
+#define CL_RG                                       0x10B2
+#define CL_RA                                       0x10B3
+#define CL_RGB                                      0x10B4
+#define CL_RGBA                                     0x10B5
+#define CL_BGRA                                     0x10B6
+#define CL_ARGB                                     0x10B7
+#define CL_INTENSITY                                0x10B8
+#define CL_LUMINANCE                                0x10B9
+
+/* cl_channel_type */
+#define CL_SNORM_INT8                               0x10D0
+#define CL_SNORM_INT16                              0x10D1
+#define CL_UNORM_INT8                               0x10D2
+#define CL_UNORM_INT16                              0x10D3
+#define CL_UNORM_SHORT_565                          0x10D4
+#define CL_UNORM_SHORT_555                          0x10D5
+#define CL_UNORM_INT_101010                         0x10D6
+#define CL_SIGNED_INT8                              0x10D7
+#define CL_SIGNED_INT16                             0x10D8
+#define CL_SIGNED_INT32                             0x10D9
+#define CL_UNSIGNED_INT8                            0x10DA
+#define CL_UNSIGNED_INT16                           0x10DB
+#define CL_UNSIGNED_INT32                           0x10DC
+#define CL_HALF_FLOAT                               0x10DD
+#define CL_FLOAT                                    0x10DE
+
+/* cl_mem_object_type */
+#define CL_MEM_OBJECT_BUFFER                        0x10F0
+#define CL_MEM_OBJECT_IMAGE2D                       0x10F1
+#define CL_MEM_OBJECT_IMAGE3D                       0x10F2
+
+/* cl_mem_info */
+#define CL_MEM_TYPE                                 0x1100
+#define CL_MEM_FLAGS                                0x1101
+#define CL_MEM_SIZE                                 0x1102
+#define CL_MEM_HOST_PTR                             0x1103
+#define CL_MEM_MAP_COUNT                            0x1104
+#define CL_MEM_REFERENCE_COUNT                      0x1105
+#define CL_MEM_CONTEXT                              0x1106
+
+/* cl_image_info */
+#define CL_IMAGE_FORMAT                             0x1110
+#define CL_IMAGE_ELEMENT_SIZE                       0x1111
+#define CL_IMAGE_ROW_PITCH                          0x1112
+#define CL_IMAGE_SLICE_PITCH                        0x1113
+#define CL_IMAGE_WIDTH                              0x1114
+#define CL_IMAGE_HEIGHT                             0x1115
+#define CL_IMAGE_DEPTH                              0x1116
+
+/* cl_addressing_mode */
+#define CL_ADDRESS_NONE                             0x1130
+#define CL_ADDRESS_CLAMP_TO_EDGE                    0x1131
+#define CL_ADDRESS_CLAMP                            0x1132
+#define CL_ADDRESS_REPEAT                           0x1133
+
+/* cl_filter_mode */
+#define CL_FILTER_NEAREST                           0x1140
+#define CL_FILTER_LINEAR                            0x1141
+
+/* cl_sampler_info */
+#define CL_SAMPLER_REFERENCE_COUNT                  0x1150
+#define CL_SAMPLER_CONTEXT                          0x1151
+#define CL_SAMPLER_NORMALIZED_COORDS                0x1152
+#define CL_SAMPLER_ADDRESSING_MODE                  0x1153
+#define CL_SAMPLER_FILTER_MODE                      0x1154
+
+/* cl_map_flags - bitfield */
+#define CL_MAP_READ                                 (1 << 0)
+#define CL_MAP_WRITE                                (1 << 1)
+
+/* cl_program_info */
+#define CL_PROGRAM_REFERENCE_COUNT                  0x1160
+#define CL_PROGRAM_CONTEXT                          0x1161
+#define CL_PROGRAM_NUM_DEVICES                      0x1162
+#define CL_PROGRAM_DEVICES                          0x1163
+#define CL_PROGRAM_SOURCE                           0x1164
+#define CL_PROGRAM_BINARY_SIZES                     0x1165
+#define CL_PROGRAM_BINARIES                         0x1166
+
+/* cl_program_build_info */
+#define CL_PROGRAM_BUILD_STATUS                     0x1181
+#define CL_PROGRAM_BUILD_OPTIONS                    0x1182
+#define CL_PROGRAM_BUILD_LOG                        0x1183
+
+/* cl_build_status */
+#define CL_BUILD_SUCCESS                            0
+#define CL_BUILD_NONE                               -1
+#define CL_BUILD_ERROR                              -2
+#define CL_BUILD_IN_PROGRESS                        -3
+
+/* cl_kernel_info */
+#define CL_KERNEL_FUNCTION_NAME                     0x1190
+#define CL_KERNEL_NUM_ARGS                          0x1191
+#define CL_KERNEL_REFERENCE_COUNT                   0x1192
+#define CL_KERNEL_CONTEXT                           0x1193
+#define CL_KERNEL_PROGRAM                           0x1194
+
+/* cl_kernel_work_group_info */
+#define CL_KERNEL_WORK_GROUP_SIZE                   0x11B0
+#define CL_KERNEL_COMPILE_WORK_GROUP_SIZE           0x11B1
+#define CL_KERNEL_LOCAL_MEM_SIZE                    0x11B2
+
+/* cl_event_info  */
+#define CL_EVENT_COMMAND_QUEUE                      0x11D0
+#define CL_EVENT_COMMAND_TYPE                       0x11D1
+#define CL_EVENT_REFERENCE_COUNT                    0x11D2
+#define CL_EVENT_COMMAND_EXECUTION_STATUS           0x11D3
+
+/* cl_command_type */
+#define CL_COMMAND_NDRANGE_KERNEL                   0x11F0
+#define CL_COMMAND_TASK                             0x11F1
+#define CL_COMMAND_NATIVE_KERNEL                    0x11F2
+#define CL_COMMAND_READ_BUFFER                      0x11F3
+#define CL_COMMAND_WRITE_BUFFER                     0x11F4
+#define CL_COMMAND_COPY_BUFFER                      0x11F5
+#define CL_COMMAND_READ_IMAGE                       0x11F6
+#define CL_COMMAND_WRITE_IMAGE                      0x11F7
+#define CL_COMMAND_COPY_IMAGE                       0x11F8
+#define CL_COMMAND_COPY_IMAGE_TO_BUFFER             0x11F9
+#define CL_COMMAND_COPY_BUFFER_TO_IMAGE             0x11FA
+#define CL_COMMAND_MAP_BUFFER                       0x11FB
+#define CL_COMMAND_MAP_IMAGE                        0x11FC
+#define CL_COMMAND_UNMAP_MEM_OBJECT                 0x11FD
+#define CL_COMMAND_MARKER                           0x11FE
+#define CL_COMMAND_ACQUIRE_GL_OBJECTS               0x11FF
+#define CL_COMMAND_RELEASE_GL_OBJECTS               0x1200
+
+/* command execution status */
+#define CL_COMPLETE                                 0x0
+#define CL_RUNNING                                  0x1
+#define CL_SUBMITTED                                0x2
+#define CL_QUEUED                                   0x3
+  
+/* cl_profiling_info  */
+#define CL_PROFILING_COMMAND_QUEUED                 0x1280
+#define CL_PROFILING_COMMAND_SUBMIT                 0x1281
+#define CL_PROFILING_COMMAND_START                  0x1282
+#define CL_PROFILING_COMMAND_END                    0x1283
+
+/********************************************************************************************************/
+
+/* Platform API */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetPlatformIDs(cl_uint          /* num_entries */,
+                 cl_platform_id * /* platforms */,
+                 cl_uint *        /* num_platforms */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL 
+soclGetPlatformInfo(cl_platform_id   /* platform */, 
+                  cl_platform_info /* param_name */,
+                  size_t           /* param_value_size */, 
+                  void *           /* param_value */,
+                  size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Device APIs */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetDeviceIDs(cl_platform_id   /* platform */,
+               cl_device_type   /* device_type */, 
+               cl_uint          /* num_entries */, 
+               cl_device_id *   /* devices */, 
+               cl_uint *        /* num_devices */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetDeviceInfo(cl_device_id    /* device */,
+                cl_device_info  /* param_name */, 
+                size_t          /* param_value_size */, 
+                void *          /* param_value */,
+                size_t *        /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Context APIs  */
+extern CL_API_ENTRY cl_context CL_API_CALL
+soclCreateContext(const cl_context_properties * /* properties */,
+                cl_uint                       /* num_devices */,
+                const cl_device_id *          /* devices */,
+                void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */,
+                void *                        /* user_data */,
+                cl_int *                      /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_context CL_API_CALL
+soclCreateContextFromType(const cl_context_properties * /* properties */,
+                        cl_device_type                /* device_type */,
+                        void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */,
+                        void *                        /* user_data */,
+                        cl_int *                      /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainContext(cl_context /* context */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseContext(cl_context /* context */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetContextInfo(cl_context         /* context */, 
+                 cl_context_info    /* param_name */, 
+                 size_t             /* param_value_size */, 
+                 void *             /* param_value */, 
+                 size_t *           /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Command Queue APIs */
+extern CL_API_ENTRY cl_command_queue CL_API_CALL
+soclCreateCommandQueue(cl_context                     /* context */, 
+                     cl_device_id                   /* device */, 
+                     cl_command_queue_properties    /* properties */,
+                     cl_int *                       /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainCommandQueue(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseCommandQueue(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetCommandQueueInfo(cl_command_queue      /* command_queue */,
+                      cl_command_queue_info /* param_name */,
+                      size_t                /* param_value_size */,
+                      void *                /* param_value */,
+                      size_t *              /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclSetCommandQueueProperty(cl_command_queue              /* command_queue */,
+                          cl_command_queue_properties   /* properties */, 
+                          cl_bool                        /* enable */,
+                          cl_command_queue_properties * /* old_properties */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Memory Object APIs  */
+extern CL_API_ENTRY cl_mem CL_API_CALL
+soclCreateBuffer(cl_context   /* context */,
+               cl_mem_flags /* flags */,
+               size_t       /* size */,
+               void *       /* host_ptr */,
+               cl_int *     /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+soclCreateImage2D(cl_context              /* context */,
+                cl_mem_flags            /* flags */,
+                const cl_image_format * /* image_format */,
+                size_t                  /* image_width */,
+                size_t                  /* image_height */,
+                size_t                  /* image_row_pitch */, 
+                void *                  /* host_ptr */,
+                cl_int *                /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+                        
+extern CL_API_ENTRY cl_mem CL_API_CALL
+soclCreateImage3D(cl_context              /* context */,
+                cl_mem_flags            /* flags */,
+                const cl_image_format * /* image_format */,
+                size_t                  /* image_width */, 
+                size_t                  /* image_height */,
+                size_t                  /* image_depth */, 
+                size_t                  /* image_row_pitch */, 
+                size_t                  /* image_slice_pitch */, 
+                void *                  /* host_ptr */,
+                cl_int *                /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+                        
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainMemObject(cl_mem /* memobj */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseMemObject(cl_mem /* memobj */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetSupportedImageFormats(cl_context           /* context */,
+                           cl_mem_flags         /* flags */,
+                           cl_mem_object_type   /* image_type */,
+                           cl_uint              /* num_entries */,
+                           cl_image_format *    /* image_formats */,
+                           cl_uint *            /* num_image_formats */) CL_API_SUFFIX__VERSION_1_0;
+                                    
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetMemObjectInfo(cl_mem           /* memobj */,
+                   cl_mem_info      /* param_name */, 
+                   size_t           /* param_value_size */,
+                   void *           /* param_value */,
+                   size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetImageInfo(cl_mem           /* image */,
+               cl_image_info    /* param_name */, 
+               size_t           /* param_value_size */,
+               void *           /* param_value */,
+               size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Sampler APIs  */
+extern CL_API_ENTRY cl_sampler CL_API_CALL
+soclCreateSampler(cl_context          /* context */,
+                cl_bool             /* normalized_coords */, 
+                cl_addressing_mode  /* addressing_mode */, 
+                cl_filter_mode      /* filter_mode */,
+                cl_int *            /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainSampler(cl_sampler /* sampler */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseSampler(cl_sampler /* sampler */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetSamplerInfo(cl_sampler         /* sampler */,
+                 cl_sampler_info    /* param_name */,
+                 size_t             /* param_value_size */,
+                 void *             /* param_value */,
+                 size_t *           /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+                            
+/* Program Object APIs  */
+extern CL_API_ENTRY cl_program CL_API_CALL
+soclCreateProgramWithSource(cl_context        /* context */,
+                          cl_uint           /* count */,
+                          const char **     /* strings */,
+                          const size_t *    /* lengths */,
+                          cl_int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_program CL_API_CALL
+soclCreateProgramWithBinary(cl_context                     /* context */,
+                          cl_uint                        /* num_devices */,
+                          const cl_device_id *           /* device_list */,
+                          const size_t *                 /* lengths */,
+                          const unsigned char **         /* binaries */,
+                          cl_int *                       /* binary_status */,
+                          cl_int *                       /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainProgram(cl_program /* program */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseProgram(cl_program /* program */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclBuildProgram(cl_program           /* program */,
+               cl_uint              /* num_devices */,
+               const cl_device_id * /* device_list */,
+               const char *         /* options */, 
+               void (*pfn_notify)(cl_program /* program */, void * /* user_data */),
+               void *               /* user_data */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclUnloadCompiler(void) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetProgramInfo(cl_program         /* program */,
+                 cl_program_info    /* param_name */,
+                 size_t             /* param_value_size */,
+                 void *             /* param_value */,
+                 size_t *           /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetProgramBuildInfo(cl_program            /* program */,
+                      cl_device_id          /* device */,
+                      cl_program_build_info /* param_name */,
+                      size_t                /* param_value_size */,
+                      void *                /* param_value */,
+                      size_t *              /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+                            
+/* Kernel Object APIs */
+extern CL_API_ENTRY cl_kernel CL_API_CALL
+soclCreateKernel(cl_program      /* program */,
+               const char *    /* kernel_name */,
+               cl_int *        /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclCreateKernelsInProgram(cl_program     /* program */,
+                         cl_uint        /* num_kernels */,
+                         cl_kernel *    /* kernels */,
+                         cl_uint *      /* num_kernels_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainKernel(cl_kernel    /* kernel */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseKernel(cl_kernel   /* kernel */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclSetKernelArg(cl_kernel    /* kernel */,
+               cl_uint      /* arg_index */,
+               size_t       /* arg_size */,
+               const void * /* arg_value */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetKernelInfo(cl_kernel       /* kernel */,
+                cl_kernel_info  /* param_name */,
+                size_t          /* param_value_size */,
+                void *          /* param_value */,
+                size_t *        /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetKernelWorkGroupInfo(cl_kernel                  /* kernel */,
+                         cl_device_id               /* device */,
+                         cl_kernel_work_group_info  /* param_name */,
+                         size_t                     /* param_value_size */,
+                         void *                     /* param_value */,
+                         size_t *                   /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Event Object APIs  */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclWaitForEvents(cl_uint             /* num_events */,
+                const cl_event *    /* event_list */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetEventInfo(cl_event         /* event */,
+               cl_event_info    /* param_name */,
+               size_t           /* param_value_size */,
+               void *           /* param_value */,
+               size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+                            
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclRetainEvent(cl_event /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseEvent(cl_event /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Profiling APIs  */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclGetEventProfilingInfo(cl_event            /* event */,
+                        cl_profiling_info   /* param_name */,
+                        size_t              /* param_value_size */,
+                        void *              /* param_value */,
+                        size_t *            /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+                                
+/* Flush and Finish APIs */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclFlush(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclFinish(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Enqueued Commands APIs */
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueReadBuffer(cl_command_queue    /* command_queue */,
+                    cl_mem              /* buffer */,
+                    cl_bool             /* blocking_read */,
+                    size_t              /* offset */,
+                    size_t              /* cb */, 
+                    void *              /* ptr */,
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+                            
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueWriteBuffer(cl_command_queue   /* command_queue */, 
+                     cl_mem             /* buffer */, 
+                     cl_bool            /* blocking_write */, 
+                     size_t             /* offset */, 
+                     size_t             /* cb */, 
+                     const void *       /* ptr */, 
+                     cl_uint            /* num_events_in_wait_list */, 
+                     const cl_event *   /* event_wait_list */, 
+                     cl_event *         /* event */) CL_API_SUFFIX__VERSION_1_0;
+                            
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyBuffer(cl_command_queue    /* command_queue */, 
+                    cl_mem              /* src_buffer */,
+                    cl_mem              /* dst_buffer */, 
+                    size_t              /* src_offset */,
+                    size_t              /* dst_offset */,
+                    size_t              /* cb */, 
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+                            
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueReadImage(cl_command_queue     /* command_queue */,
+                   cl_mem               /* image */,
+                   cl_bool              /* blocking_read */, 
+                   const size_t *       /* origin[3] */,
+                   const size_t *       /* region[3] */,
+                   size_t               /* row_pitch */,
+                   size_t               /* slice_pitch */, 
+                   void *               /* ptr */,
+                   cl_uint              /* num_events_in_wait_list */,
+                   const cl_event *     /* event_wait_list */,
+                   cl_event *           /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueWriteImage(cl_command_queue    /* command_queue */,
+                    cl_mem              /* image */,
+                    cl_bool             /* blocking_write */, 
+                    const size_t *      /* origin[3] */,
+                    const size_t *      /* region[3] */,
+                    size_t              /* input_row_pitch */,
+                    size_t              /* input_slice_pitch */, 
+                    const void *        /* ptr */,
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyImage(cl_command_queue     /* command_queue */,
+                   cl_mem               /* src_image */,
+                   cl_mem               /* dst_image */, 
+                   const size_t *       /* src_origin[3] */,
+                   const size_t *       /* dst_origin[3] */,
+                   const size_t *       /* region[3] */, 
+                   cl_uint              /* num_events_in_wait_list */,
+                   const cl_event *     /* event_wait_list */,
+                   cl_event *           /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyImageToBuffer(cl_command_queue /* command_queue */,
+                           cl_mem           /* src_image */,
+                           cl_mem           /* dst_buffer */, 
+                           const size_t *   /* src_origin[3] */,
+                           const size_t *   /* region[3] */, 
+                           size_t           /* dst_offset */,
+                           cl_uint          /* num_events_in_wait_list */,
+                           const cl_event * /* event_wait_list */,
+                           cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyBufferToImage(cl_command_queue /* command_queue */,
+                           cl_mem           /* src_buffer */,
+                           cl_mem           /* dst_image */, 
+                           size_t           /* src_offset */,
+                           const size_t *   /* dst_origin[3] */,
+                           const size_t *   /* region[3] */, 
+                           cl_uint          /* num_events_in_wait_list */,
+                           const cl_event * /* event_wait_list */,
+                           cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY void * CL_API_CALL
+soclEnqueueMapBuffer(cl_command_queue /* command_queue */,
+                   cl_mem           /* buffer */,
+                   cl_bool          /* blocking_map */, 
+                   cl_map_flags     /* map_flags */,
+                   size_t           /* offset */,
+                   size_t           /* cb */,
+                   cl_uint          /* num_events_in_wait_list */,
+                   const cl_event * /* event_wait_list */,
+                   cl_event *       /* event */,
+                   cl_int *         /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY void * CL_API_CALL
+soclEnqueueMapImage(cl_command_queue  /* command_queue */,
+                  cl_mem            /* image */, 
+                  cl_bool           /* blocking_map */, 
+                  cl_map_flags      /* map_flags */, 
+                  const size_t *    /* origin[3] */,
+                  const size_t *    /* region[3] */,
+                  size_t *          /* image_row_pitch */,
+                  size_t *          /* image_slice_pitch */,
+                  cl_uint           /* num_events_in_wait_list */,
+                  const cl_event *  /* event_wait_list */,
+                  cl_event *        /* event */,
+                  cl_int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueUnmapMemObject(cl_command_queue /* command_queue */,
+                        cl_mem           /* memobj */,
+                        void *           /* mapped_ptr */,
+                        cl_uint          /* num_events_in_wait_list */,
+                        const cl_event *  /* event_wait_list */,
+                        cl_event *        /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueNDRangeKernel(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 */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueTask(cl_command_queue  /* command_queue */,
+              cl_kernel         /* kernel */,
+              cl_uint           /* num_events_in_wait_list */,
+              const cl_event *  /* event_wait_list */,
+              cl_event *        /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueNativeKernel(cl_command_queue  /* command_queue */,
+					  void (*user_func)(void *), 
+                      void *            /* args */,
+                      size_t            /* cb_args */, 
+                      cl_uint           /* num_mem_objects */,
+                      const cl_mem *    /* mem_list */,
+                      const void **     /* args_mem_loc */,
+                      cl_uint           /* num_events_in_wait_list */,
+                      const cl_event *  /* event_wait_list */,
+                      cl_event *        /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueMarker(cl_command_queue    /* command_queue */,
+                cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueWaitForEvents(cl_command_queue /* command_queue */,
+                       cl_uint          /* num_events */,
+                       const cl_event * /* event_list */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueBarrier(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+/* Extension function access
+ *
+ * Returns the extension function address for the given function name,
+ * or NULL if a valid function can not be found.  The client must
+ * check to make sure the address is not NULL, before using or 
+ * calling the returned function address.
+ */
+extern CL_API_ENTRY void * CL_API_CALL
+soclGetExtensionFunctionAddress(const char * /* func_name */) CL_API_SUFFIX__VERSION_1_0;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_H */
+

socl/include/CL/cl_ext.h

@@ -0,0 +1,60 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2009 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+/* $Revision: 10424 $ on $Date: 2010-02-17 14:34:49 -0800 (Wed, 17 Feb 2010) $ */
+
+#ifndef __CL_EXT_H
+#define __CL_EXT_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/* cl_khr_fp64 extension - no extension #define since it has no functions  */
+#define CL_DEVICE_DOUBLE_FP_CONFIG                  0x1032
+
+
+/* cl_khr_fp16 extension - no extension #define since it has no functions  */
+#define CL_DEVICE_HALF_FP_CONFIG                    0x1033
+
+
+/* cl_khr_icd extension                                                    */
+#define cl_khr_icd 1
+
+/* cl_platform_info                                                        */
+#define CL_PLATFORM_ICD_SUFFIX_KHR                  0x0920
+
+/* Additional Error Codes                                                  */
+#define CL_PLATFORM_NOT_FOUND_KHR                   -1001
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clIcdGetPlatformIDsKHR(cl_uint          /* num_entries */,
+                       cl_platform_id * /* platforms */,
+                       cl_uint *        /* num_platforms */);
+
+#ifdef __cplusplus
+}
+#endif
+
+
+#endif /* __CL_EXT_H */

socl/include/CL/cl_gl.h

@@ -0,0 +1,146 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2009 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 10424 $ on $Date: 2010-02-17 14:34:49 -0800 (Wed, 17 Feb 2010) $ */
+
+/*
+ * cl_gl.h contains Khronos-approved (KHR) OpenCL extensions which have
+ * OpenGL dependencies. The application is responsible for #including
+ * OpenGL or OpenGL ES headers before #including cl_gl.h.
+ */
+
+#ifndef __OPENCL_CL_GL_H
+#define __OPENCL_CL_GL_H
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif	
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef cl_uint     cl_gl_object_type;
+typedef cl_uint     cl_gl_texture_info;
+typedef cl_uint     cl_gl_platform_info;
+
+/* cl_gl_object_type */
+#define CL_GL_OBJECT_BUFFER             0x2000
+#define CL_GL_OBJECT_TEXTURE2D          0x2001
+#define CL_GL_OBJECT_TEXTURE3D          0x2002
+#define CL_GL_OBJECT_RENDERBUFFER       0x2003
+
+/* cl_gl_texture_info */
+#define CL_GL_TEXTURE_TARGET            0x2004
+#define CL_GL_MIPMAP_LEVEL              0x2005
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLBuffer(cl_context     /* context */,
+                     cl_mem_flags   /* flags */,
+                     cl_GLuint      /* bufobj */,
+                     int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLTexture2D(cl_context      /* context */,
+                        cl_mem_flags    /* flags */,
+                        cl_GLenum       /* target */,
+                        cl_GLint        /* miplevel */,
+                        cl_GLuint       /* texture */,
+                        cl_int *        /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLTexture3D(cl_context      /* context */,
+                        cl_mem_flags    /* flags */,
+                        cl_GLenum       /* target */,
+                        cl_GLint        /* miplevel */,
+                        cl_GLuint       /* texture */,
+                        cl_int *        /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLRenderbuffer(cl_context   /* context */,
+                           cl_mem_flags /* flags */,
+                           cl_GLuint    /* renderbuffer */,
+                           cl_int *     /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLObjectInfo(cl_mem                /* memobj */,
+                  cl_gl_object_type *   /* gl_object_type */,
+                  cl_GLuint *              /* gl_object_name */) CL_API_SUFFIX__VERSION_1_0;
+                  
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLTextureInfo(cl_mem               /* memobj */,
+                   cl_gl_texture_info   /* param_name */,
+                   size_t               /* param_value_size */,
+                   void *               /* param_value */,
+                   size_t *             /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueAcquireGLObjects(cl_command_queue      /* command_queue */,
+                          cl_uint               /* num_objects */,
+                          const cl_mem *        /* mem_objects */,
+                          cl_uint               /* num_events_in_wait_list */,
+                          const cl_event *      /* event_wait_list */,
+                          cl_event *            /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueReleaseGLObjects(cl_command_queue      /* command_queue */,
+                          cl_uint               /* num_objects */,
+                          const cl_mem *        /* mem_objects */,
+                          cl_uint               /* num_events_in_wait_list */,
+                          const cl_event *      /* event_wait_list */,
+                          cl_event *            /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+/* cl_khr_gl_sharing extension  */
+
+#define cl_khr_gl_sharing 1
+
+typedef cl_uint     cl_gl_context_info;
+
+/* Additional Error Codes  */
+#define CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR  -1000
+
+/* cl_gl_context_info  */
+#define CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR    0x2006
+#define CL_DEVICES_FOR_GL_CONTEXT_KHR           0x2007
+
+/* Additional cl_context_properties  */
+#define CL_GL_CONTEXT_KHR                       0x2008
+#define CL_EGL_DISPLAY_KHR                      0x2009
+#define CL_GLX_DISPLAY_KHR                      0x200A
+#define CL_WGL_HDC_KHR                          0x200B
+#define CL_CGL_SHAREGROUP_KHR                   0x200C
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLContextInfoKHR(const cl_context_properties * /* properties */,
+                      cl_gl_context_info            /* param_name */,
+                      size_t                        /* param_value_size */,
+                      void *                        /* param_value */,
+                      size_t *                      /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_GL_H  */

socl/include/CL/cl_gl_ext.h

@@ -0,0 +1,52 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2009 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 10424 $ on $Date: 2010-02-17 14:34:49 -0800 (Wed, 17 Feb 2010) $ */
+
+/* cl_gl_ext.h contains vendor (non-KHR) OpenCL extensions which have           */
+/* OpenGL dependencies.                                                         */
+
+#ifndef __OPENCL_CL_GL_EXT_H
+#define __OPENCL_CL_GL_EXT_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/*
+ * For each extension, follow this template
+ * // cl_VEN_extname extension 
+ * #define cl_VEN_extname 1
+ * ... define new types, if any
+ * ... define new tokens, if any
+ * ... define new APIs, if any
+ *
+ *  If you need GLtypes here, mirror them with a cl_GLtype, rather than including a GL header
+ *  This allows us to avoid having to decide whether to include GL headers or GLES here.
+ */
+ 
+#ifdef __cplusplus
+}
+#endif
+
+#endif	/* __OPENCL_CL_GL_EXT_H  */

socl/include/CL/cl_starpu.h

@@ -0,0 +1,80 @@
+#define SOCL
+
+#define clGetPlatformIDs soclGetPlatformIDs
+#define clGetPlatformInfo soclGetPlatformInfo
+
+#define clGetDeviceIDs soclGetDeviceIDs
+#define clGetDeviceInfo soclGetDeviceInfo
+
+#define clCreateContext soclCreateContext
+#define clCreateContextFromType soclCreateContextFromType
+#define clRetainContext soclRetainContext
+#define clReleaseContext soclReleaseContext
+#define clGetContextInfo soclGetContextInfo
+
+#define clCreateCommandQueue soclCreateCommandQueue
+#define clRetainCommandQueue soclRetainCommandQueue
+#define clReleaseCommandQueue soclReleaseCommandQueue
+#define clGetCommandQueueInfo soclGetCommandQueueInfo
+#define clSetCommandQueueProperty soclSetCommandQueueProperty
+
+#define clCreateBuffer soclCreateBuffer
+#define clCreateImage2D soclCreateImage2D
+#define clCreateImage3D soclCreateImage3D
+#define clRetainMemObject soclRetainMemObject
+#define clReleaseMemObject soclReleaseMemObject
+#define clGetSupportedImageFormats soclGetSupportedImageFormats
+#define clGetMemObjectInfo soclGetMemObjectInfo
+#define clGetImageInfo soclGetImageInfo
+
+#define clCreateSampler soclCreateSampler
+#define clRetainSampler soclRetainSampler
+#define clReleaseSampler soclReleaseSampler
+#define clGetSamplerInfo soclGetSamplerInfo
+
+#define clCreateProgramWithSource soclCreateProgramWithSource
+#define clCreateProgramWithBinary soclCreateProgramWithBinary
+#define clRetainProgram soclRetainProgram
+#define clReleaseProgram soclReleaseProgram
+#define clBuildProgram soclBuildProgram
+#define clUnloadCompiler soclUnloadCompiler
+#define clGetProgramInfo soclGetProgramInfo
+#define clGetProgramBuildInfo soclGetProgramBuildInfo
+
+#define clCreateKernel soclCreateKernel
+#define clCreateKernelsInProgram soclCreateKernelsInProgram
+#define clRetainKernel soclRetainKernel
+#define clReleaseKernel soclReleaseKernel
+#define clSetKernelArg soclSetKernelArg
+#define clGetKernelInfo soclGetKernelInfo
+#define clGetKernelWorkGroupInfo soclGetKernelWorkGroupInfo
+
+#define clWaitForEvents soclWaitForEvents
+#define clGetEventInfo soclGetEventInfo
+#define clRetainEvent soclRetainEvent
+#define clReleaseEvent soclReleaseEvent
+
+#define clGetEventProfilingInfo soclGetEventProfilingInfo
+
+#define clFlush soclFlush
+#define clFinish soclFinish
+
+#define clEnqueueReadBuffer soclEnqueueReadBuffer
+#define clEnqueueWriteBuffer soclEnqueueWriteBuffer
+#define clEnqueueCopyBuffer soclEnqueueCopyBuffer
+#define clEnqueueReadImage soclEnqueueReadImage
+#define clEnqueueWriteImage soclEnqueueWriteImage
+#define clEnqueueCopyImage soclEnqueueCopyImage
+#define clEnqueueCopyImageToBuffer soclEnqueueCopyImageToBuffer
+#define clEnqueueCopyBufferToImage soclEnqueueCopyBufferToImage
+#define clEnqueueMapBuffer soclEnqueueMapBuffer
+#define clEnqueueMapImage soclEnqueueMapImage
+#define clEnqueueUnmapMemObject soclEnqueueUnmapMemObject
+#define clEnqueueNDRangeKernel soclEnqueueNDRangeKernel
+#define clEnqueueTask soclEnqueueTask
+#define clEnqueueNativeKernel soclEnqueueNativeKernel
+#define clEnqueueMarker soclEnqueueMarker
+#define clEnqueueWaitForEvents soclEnqueueWaitForEvents
+#define clEnqueueBarrier soclEnqueueBarrier
+
+#define clGetExtensionFunctionAddress soclGetExtensionFunctionAddress

socl/include/CL/opencl.h

@@ -0,0 +1,54 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2009 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+/* $Revision: 10424 $ on $Date: 2010-02-17 14:34:49 -0800 (Wed, 17 Feb 2010) $ */
+
+#ifndef __OPENCL_H
+#define __OPENCL_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifdef __APPLE__
+
+#include <OpenCL/cl.h>
+#include <OpenCL/cl_gl.h>
+#include <OpenCL/cl_gl_ext.h>
+#include <OpenCL/cl_ext.h>
+
+#else
+
+#include <CL/cl.h>
+#include <CL/cl_gl.h>
+#include <CL/cl_gl_ext.h>
+#include <CL/cl_ext.h>
+
+#endif
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_H   */
+

socl/src/Makefile.am

@@ -0,0 +1,28 @@
+# StarPU --- Runtime system for heterogeneous multicore architectures.
+#
+# Copyright (C) 2009, 2010, 2011  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.
+
+AM_CFLAGS = -Wall -Wextra
+LIBS = $(top_builddir)/src/libstarpu.la
+AM_CPPFLAGS = -I$(top_srcdir)/include/ -I$(top_builddir)/include
+AM_LDFLAGS = $(STARPU_CUDA_LDFLAGS) $(STARPU_OPENCL_LDFLAGS)
+
+SUBDIRS =
+
+lib_LTLIBRARIES = libsocl.la
+
+libsocl_la_SOURCES = 						\
+  opencl.c
+

socl/src/cl_buildprogram.c.inc

@@ -0,0 +1,76 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+struct bp_data {
+   cl_program program;
+   char * options;
+};
+
+static void soclBuildProgram_task(void *data) {
+   struct bp_data *d = (struct bp_data*)data;
+   cl_device_id device;
+   cl_int err;
+
+   int wid = starpu_worker_get_id();
+   int range = starpu_worker_get_range();
+   starpu_opencl_get_device(wid, &device);
+
+   DEBUG_MSG("[Worker %d] Building program...\n", wid);
+
+   err = clBuildProgram(d->program->cl_programs[range], 1, &device, d->options, NULL, NULL);
+   if (err != CL_SUCCESS) {
+      size_t len;
+      static char buffer[4096];
+      clGetProgramBuildInfo(d->program->cl_programs[range], device, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
+      DEBUG_CL("clBuildProgram", err);
+      ERROR_MSG("clBuildProgram: %s\n Aborting.\n", buffer);
+   }
+
+   DEBUG_MSG("[Worker %d] Done building.\n", wid);
+}
+
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclBuildProgram(cl_program         program,
+               cl_uint              UNUSED(num_devices),
+               const cl_device_id * UNUSED(device_list),
+               const char *         options, 
+               void (*pfn_notify)(cl_program program, void * user_data),
+               void *               user_data) CL_API_SUFFIX__VERSION_1_0
+{
+   struct bp_data *data;
+
+   program->options = options != NULL ? strdup(options) : NULL;
+   program->options_size = options != NULL ? strlen(options)+1 : 0;
+
+   data = (struct bp_data*)malloc(sizeof(struct bp_data));
+   gc_entity_store(&data->program, program);
+   data->options = (char*)options;
+
+   /*FIXME: starpu_execute_on_each_worker is synchronous.
+    * However pfn_notify may be useful only because build is supposed to be asynchronous
+    */
+   starpu_execute_on_each_worker(soclBuildProgram_task, data, STARPU_OPENCL);
+
+
+   if (pfn_notify != NULL)
+      pfn_notify(program, user_data);
+
+   gc_entity_unstore(&data->program);
+   free(data);
+
+   return CL_SUCCESS;
+}

socl/src/cl_createbuffer.c.inc

@@ -0,0 +1,130 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+/**
+ * \brief Create a buffer
+ *
+ * A buffer has always an allocated region in host memory. If CL_MEM_USE_HOST_PTR
+ * is set, we use memory pointed by host_ptr, otherwise some host memory is
+ * allocated.
+ *
+ * If CL_MEM_USE_HOST_PTR or CL_MEM_ALLOC_HOST_PTR are set, memory pointed by host_ptr
+ * is not coherent. To enforce coherency, you have to map the buffer (clEnqueueMapBuffer).
+ * 
+ * If CL_MEM_COPY_HOST_PTR is set, the buffer will be duplicated in host memory. You
+ * should avoid it.
+ *
+ */
+CL_API_ENTRY cl_mem CL_API_CALL
+soclCreateBuffer(cl_context   context,
+               cl_mem_flags flags,
+               size_t       size,
+               void *       host_ptr,
+               cl_int *     errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (errcode_ret != NULL)
+      *errcode_ret = CL_SUCCESS;
+
+   //Check flags
+   if (((flags & CL_MEM_READ_ONLY) && (flags & CL_MEM_WRITE_ONLY))
+      || ((flags & CL_MEM_READ_WRITE) && (flags & CL_MEM_READ_ONLY))
+      || ((flags & CL_MEM_READ_WRITE) && (flags & CL_MEM_WRITE_ONLY))
+      || ((flags & CL_MEM_USE_HOST_PTR) && (flags & CL_MEM_ALLOC_HOST_PTR))
+      || ((flags & CL_MEM_USE_HOST_PTR) && (flags & CL_MEM_COPY_HOST_PTR))) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_INVALID_VALUE;
+      return NULL;
+   }
+
+   if (size == 0) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_INVALID_BUFFER_SIZE;
+      return NULL;
+   }
+
+   if ((host_ptr == NULL && (flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR)))
+      || (host_ptr != NULL && !(flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR)))) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_INVALID_HOST_PTR;
+      return NULL;
+   }
+
+   {
+      cl_mem mem;
+
+      //Alloc cl_mem structure
+      mem = (cl_mem)gc_entity_alloc(sizeof(struct _cl_mem), release_callback_memobject);
+      if (mem == NULL) {
+         if (errcode_ret != NULL)
+            *errcode_ret = CL_OUT_OF_HOST_MEMORY;
+         return NULL;
+      }
+      
+      mem->ptr = NULL;
+      mem->map_count = 0;
+      gc_entity_store(&mem->context, context);
+      mem->flags = flags;
+      mem->size = size;
+      mem->host_ptr = host_ptr;
+
+      #ifdef DEBUG
+      static int id = 0;
+      mem->id = id++;
+      #endif
+
+      mem_object_store(mem);
+
+      //TODO: we shouldn't allocate the buffer ourselves. StarPU allocates it if a NULL pointer is given
+
+      // If not MEM_USE_HOST_PTR, we need to alloc the buffer ourselves
+      if (!(flags & CL_MEM_USE_HOST_PTR)) {
+         mem->ptr = valloc(size);
+         if (mem->ptr == NULL) {
+            if (errcode_ret != NULL)
+               *errcode_ret = CL_MEM_OBJECT_ALLOCATION_FAILURE;
+            free(mem);
+            return NULL;
+         }
+         //The buffer doesn't contain meaningful data
+         mem->scratch = 1;
+      }
+      else {
+         //The buffer may contain meaningful data
+         mem->scratch = 0;
+         mem->ptr = host_ptr;
+      }
+
+      // Access mode
+      if (flags & CL_MEM_READ_ONLY)
+         mem->mode = CL_MEM_READ_ONLY;
+      else if (flags & CL_MEM_WRITE_ONLY)
+         mem->mode = CL_MEM_WRITE_ONLY;
+      else
+         mem->mode = CL_MEM_READ_WRITE;
+
+      // Perform data copy if necessary
+      if (flags & CL_MEM_COPY_HOST_PTR)
+         memcpy(mem->ptr, host_ptr, size);
+      
+      // Create StarPU buffer (on home node? what's this?)
+      starpu_variable_data_register(&mem->handle, 0, (uintptr_t)mem->ptr, size); 
+
+      DEBUG_MSG("[Buffer %d] Initialized (cl_mem %p handle %p)\n", mem->id, mem, mem->handle);
+      
+      return mem;
+   }
+}

socl/src/cl_createcommandqueue.c.inc

@@ -0,0 +1,58 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+CL_API_ENTRY cl_command_queue CL_API_CALL
+soclCreateCommandQueue(cl_context                   context, 
+                     cl_device_id                   device, 
+                     cl_command_queue_properties    properties,
+                     cl_int *                       errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   cl_command_queue cq;
+
+   cq = (cl_command_queue)gc_entity_alloc(sizeof(struct _cl_command_queue),
+                                          release_callback_command_queue);
+   if (cq == NULL) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_OUT_OF_HOST_MEMORY;
+      return NULL;
+   }
+
+   cq->properties = properties;
+   gc_entity_store(&cq->context, context);
+   cq->device = device;
+
+   #ifdef DEBUG
+   static int id = 0;
+   cq->id = id++;
+   #endif
+
+   //Enable StarPU profiling if necessary
+   if (properties & CL_QUEUE_PROFILING_ENABLE) {
+      if (profiling_queue_count == 0)
+         starpu_profiling_status_set(STARPU_PROFILING_ENABLE);
+      profiling_queue_count += 1;
+   }
+
+   cq->events = NULL;
+   cq->barrier = NULL;
+   pthread_spin_init(&cq->spin, 0);
+
+   if (errcode_ret != NULL)
+      *errcode_ret = CL_SUCCESS;
+
+   return cq;
+}

socl/src/cl_createcontext.c.inc

@@ -0,0 +1,94 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+CL_API_ENTRY cl_context CL_API_CALL
+soclCreateContext(const cl_context_properties * properties,
+                cl_uint                       num_devices,
+                const cl_device_id *          devices,
+                void (*pfn_notify)(const char *, const void *, size_t, void *),
+                void *                        user_data,
+                cl_int *                      errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+
+   if (pfn_notify == NULL && user_data != NULL) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_INVALID_VALUE;
+      return NULL;
+   }
+
+   //Check properties
+   if (properties != NULL) {
+      const cl_context_properties *p = properties;
+      int i = 0;
+      while (p[i] != 0) {
+         switch (p[i]) {
+            case CL_CONTEXT_PLATFORM:
+               i++;
+               if (p[i] != (cl_context_properties)&socl_platform) {
+                  if (errcode_ret != NULL)
+                     *errcode_ret = CL_INVALID_PLATFORM;
+                  return NULL;
+               }
+               break;
+         }
+         i++;
+      }
+   }
+
+
+   cl_context ctx;
+   ctx = (cl_context)gc_entity_alloc(sizeof(struct _cl_context), release_callback_context);
+   if (ctx == NULL) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_OUT_OF_HOST_MEMORY;
+      return NULL;
+   }
+
+   ctx->num_properties = 0;
+   ctx->properties = NULL;
+
+   // Cache properties
+   if (properties != NULL) {
+      //Count properties
+      const cl_context_properties * p = properties;
+      do {
+         ctx->num_properties++;
+         p++;
+      } while (*p != 0);
+
+      //Copy properties
+      ctx->properties = malloc(sizeof(cl_context_properties) * ctx->num_properties);
+      memcpy(ctx->properties, properties, sizeof(cl_context_properties) * ctx->num_properties);
+   }
+
+   ctx->pfn_notify = pfn_notify;
+   ctx->user_data = user_data;
+   ctx->num_devices = num_devices;
+
+   #ifdef DEBUG
+   static int id = 0;
+   ctx->id = id++;
+   #endif
+
+   ctx->devices = malloc(sizeof(cl_device_id) * num_devices);
+   memcpy(ctx->devices, devices, sizeof(cl_device_id)*num_devices);
+
+   if (errcode_ret != NULL)
+      *errcode_ret = CL_SUCCESS;
+
+   return ctx;
+}

socl/src/cl_createcontextfromtype.c.inc

@@ -0,0 +1,27 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_context CL_API_CALL
+soclCreateContextFromType(const cl_context_properties * properties,
+                        cl_device_type                UNUSED(device_type),
+                        void (*pfn_notify)(const char *, const void *, size_t, void *),
+                        void *                        user_data,
+                        cl_int *                      errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   //We assume clCreateContext doesn't support devices
+   //TODO:use devices
+   return soclCreateContext(properties, 0, NULL, pfn_notify, user_data, errcode_ret);
+}

socl/src/cl_createimage2d.c.inc

@@ -0,0 +1,30 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_mem CL_API_CALL
+soclCreateImage2D(cl_context              UNUSED(context),
+                cl_mem_flags            UNUSED(flags),
+                const cl_image_format * UNUSED(image_format),
+                size_t                  UNUSED(image_width),
+                size_t                  UNUSED(image_height),
+                size_t                  UNUSED(image_row_pitch), 
+                void *                  UNUSED(host_ptr),
+                cl_int *                errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (errcode_ret != NULL)
+      *errcode_ret = CL_INVALID_OPERATION;
+   return NULL;
+}

socl/src/cl_createimage3d.c.inc

@@ -0,0 +1,32 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_mem CL_API_CALL
+soclCreateImage3D(cl_context              UNUSED(context),
+                cl_mem_flags            UNUSED(flags),
+                const cl_image_format * UNUSED(image_format),
+                size_t                  UNUSED(image_width), 
+                size_t                  UNUSED(image_height),
+                size_t                  UNUSED(image_depth), 
+                size_t                  UNUSED(image_row_pitch), 
+                size_t                  UNUSED(image_slice_pitch), 
+                void *                  UNUSED(host_ptr),
+                cl_int *                errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (errcode_ret != NULL)
+      *errcode_ret = CL_INVALID_OPERATION;
+   return NULL;
+}

socl/src/cl_createkernel.c.inc

@@ -0,0 +1,125 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+static void soclCreateKernel_task(void *data) {
+   struct _cl_kernel *k = (struct _cl_kernel *)data;
+
+   int range = starpu_worker_get_range();
+   cl_int err;
+
+   DEBUG_MSG("[Device %d] Creating kernel...\n", starpu_worker_get_id());
+   k->cl_kernels[range] = clCreateKernel(k->program->cl_programs[range], k->kernel_name, &err);
+   if (err != CL_SUCCESS) {
+      k->errcodes[range] = err;
+      ERROR_STOP("[Device %d] Unable to create kernel. Aborting.\n", starpu_worker_get_id());
+      return;
+   }
+
+   /* One worker creates argument structures */
+   if (__sync_bool_compare_and_swap(&k->arg_count, 0, 666)) {
+      unsigned int i;
+      cl_uint arg_count;
+
+      err = clGetKernelInfo(k->cl_kernels[range], CL_KERNEL_NUM_ARGS, sizeof(arg_count), &arg_count, NULL);
+      if (err != CL_SUCCESS) {
+         DEBUG_CL("clGetKernelInfo", err);
+         ERROR_STOP("Unable to get kernel argument count. Aborting.\n");
+      }
+      k->arg_count = arg_count;
+      DEBUG_MSG("Kernel has %d arguments\n", arg_count);
+
+      k->arg_size = (size_t*)malloc(sizeof(size_t) * arg_count);
+      k->arg_value = (void**)malloc(sizeof(void*) * arg_count);
+      k->arg_type = (enum kernel_arg_type*)malloc(sizeof(enum kernel_arg_type) * arg_count);
+      /* Settings default type to NULL */
+      for (i=0; i<arg_count; i++) {
+         k->arg_value[i] = NULL;
+         k->arg_type[i] = Null;
+      }
+   }
+}
+
+CL_API_ENTRY cl_kernel CL_API_CALL
+soclCreateKernel(cl_program    program,
+               const char *    kernel_name,
+               cl_int *        errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   cl_kernel k;
+   int device_count;
+
+   if (program == NULL) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_INVALID_PROGRAM;
+      return NULL;
+   }
+
+   //TODO: check programs (see opencl specs)
+
+   /* Create Kernel structure */
+   k = (cl_kernel)gc_entity_alloc(sizeof(struct _cl_kernel), release_callback_kernel);
+   if (k == NULL) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_OUT_OF_HOST_MEMORY;
+      return NULL;
+   }
+   
+   gc_entity_store(&k->program, program);
+   k->kernel_name = strdup(kernel_name);
+   k->arg_count = 0;
+   k->arg_value = NULL;
+   k->arg_size = NULL;
+
+   #ifdef DEBUG
+   static int id = 0;
+   k->id = id++;
+   #endif
+   
+   device_count = starpu_opencl_worker_get_count();
+   k->cl_kernels = (cl_kernel*)malloc(device_count * sizeof(cl_kernel));
+   k->errcodes = (cl_int*)malloc(device_count * sizeof(cl_int));
+   {
+      int i;
+      for (i=0; i<device_count; i++) {
+         k->cl_kernels[i] = NULL;
+         k->errcodes[i] = -9999;
+      }
+   }
+
+   /* Create kernel on each device */
+   DEBUG_MSG("[Kernel %d] Create %d kernels (name \"%s\")\n", k->id, starpu_opencl_worker_get_count(), kernel_name);
+   starpu_execute_on_each_worker(soclCreateKernel_task, k, STARPU_OPENCL);
+
+   if (errcode_ret != NULL) {
+      int i;
+      *errcode_ret = CL_SUCCESS;
+      for (i=0; i<device_count; i++) {
+         switch (k->errcodes[i]) {
+            #define CASE_RET(e) case e: *errcode_ret = e; return k;
+            CASE_RET(CL_INVALID_PROGRAM)
+            CASE_RET(CL_INVALID_PROGRAM_EXECUTABLE)
+            CASE_RET(CL_INVALID_KERNEL_NAME)
+            CASE_RET(CL_INVALID_KERNEL_DEFINITION)
+            CASE_RET(CL_INVALID_VALUE)
+            CASE_RET(CL_OUT_OF_RESOURCES)
+            CASE_RET(CL_OUT_OF_HOST_MEMORY)
+            #undef CASE_RET
+         }
+      }
+   }
+
+   return k;
+}

socl/src/cl_createkernelsinprogram.c.inc

@@ -0,0 +1,25 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclCreateKernelsInProgram(cl_program   UNUSED(program),
+                         cl_uint        UNUSED(num_kernels),
+                         cl_kernel *    UNUSED(kernels),
+                         cl_uint *      UNUSED(num_kernels_ret)) CL_API_SUFFIX__VERSION_1_0
+{
+   //TODO
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_createprogramwithbinary.c.inc

@@ -0,0 +1,31 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_program CL_API_CALL
+soclCreateProgramWithBinary(cl_context                     UNUSED(context),
+                          cl_uint                        UNUSED(num_devices),
+                          const cl_device_id *           UNUSED(device_list),
+                          const size_t *                 UNUSED(lengths),
+                          const unsigned char **         UNUSED(binaries),
+                          cl_int *                       UNUSED(binary_status),
+                          cl_int *                       errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   //TODO
+   if (errcode_ret != NULL)
+      *errcode_ret = CL_INVALID_OPERATION;
+
+   return NULL;
+}

socl/src/cl_createprogramwithsource.c.inc

@@ -0,0 +1,129 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+struct cpws_data {
+   struct _cl_program *program;
+   cl_int *errcodes;
+   cl_uint count;
+   char **strings;
+   size_t *lengths;
+};
+
+static void soclCreateProgramWithSource_task(void *data) {
+   struct cpws_data *d = (struct cpws_data*)data;
+
+   cl_context context;
+   int wid = starpu_worker_get_id();
+   DEBUG_MSG("Worker id: %d\n", wid);
+
+   int range = starpu_worker_get_range();
+
+   starpu_opencl_get_context(wid, &context);
+
+   d->program->cl_programs[range] = clCreateProgramWithSource(context, d->count, (const char**)d->strings, d->lengths, &d->errcodes[range]);
+
+}
+
+
+CL_API_ENTRY cl_program CL_API_CALL
+soclCreateProgramWithSource(cl_context      context,
+                          cl_uint           count,
+                          const char **     strings,
+                          const size_t *    lengths,
+                          cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   cl_program p;
+   struct cpws_data *data;
+   int device_count;
+
+   if (errcode_ret != NULL)
+      *errcode_ret = CL_SUCCESS;
+
+   device_count = starpu_opencl_worker_get_count();
+   DEBUG_MSG("Worker count: %d\n", device_count);
+
+   /* Check arguments */
+   if (count == 0 || strings == NULL) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_INVALID_VALUE;
+      return NULL;
+   }
+
+   /* Alloc cl_program structure */
+   p = (cl_program)gc_entity_alloc(sizeof(struct _cl_program), release_callback_program);
+   if (p == NULL) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_OUT_OF_HOST_MEMORY;
+      return NULL;
+   }
+
+   gc_entity_store(&p->context, context);
+   p->options = NULL;
+
+   #ifdef DEBUG
+   static int id = 0;
+   p->id = id++;
+   #endif
+
+
+   p->cl_programs = (cl_program*)malloc(sizeof(cl_program) * device_count);
+   if (p->cl_programs == NULL) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_OUT_OF_HOST_MEMORY;
+      return NULL;
+   }
+
+   {
+      int i;
+      for (i=0; i<device_count; i++)
+         p->cl_programs[i] = NULL;
+   }
+
+   /* Construct structure to pass arguments to workers */
+   data = (struct cpws_data*)malloc(sizeof(struct cpws_data));
+   if (data == NULL) {
+      if (errcode_ret != NULL)
+         *errcode_ret = CL_OUT_OF_HOST_MEMORY;
+      return NULL;
+   }
+   data->count = count;
+   data->program = p;
+   data->strings = (char**)strings;
+   data->lengths = (size_t*)lengths;
+
+   data->errcodes = (cl_int*)malloc(sizeof(cl_int) * device_count);
+
+   /* Init real cl_program for each OpenCL device */
+   starpu_execute_on_each_worker(soclCreateProgramWithSource_task, data, STARPU_OPENCL);
+
+   if (errcode_ret != NULL) {
+      int i;
+      *errcode_ret = CL_SUCCESS;
+      for (i=0; i<device_count; i++) {
+         if (data->errcodes[i] != CL_SUCCESS) {
+            DEBUG_MSG("Worker [%d] failed\n", i);
+            DEBUG_CL("clCreateProgramWithSource", data->errcodes[i]);
+            *errcode_ret = data->errcodes[i];
+            break;
+         }
+      }
+   }
+
+   free(data->errcodes);
+   free(data);
+   
+   return p;
+}

socl/src/cl_createsampler.c.inc

@@ -0,0 +1,27 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_sampler CL_API_CALL
+soclCreateSampler(cl_context          UNUSED(context),
+                cl_bool             UNUSED(normalized_coords), 
+                cl_addressing_mode  UNUSED(addressing_mode), 
+                cl_filter_mode      UNUSED(filter_mode),
+                cl_int *            errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (errcode_ret != NULL)
+      *errcode_ret = CL_INVALID_OPERATION;
+   return NULL;
+}

socl/src/cl_enqueuebarrier.c.inc

@@ -0,0 +1,24 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueBarrier(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
+{
+   cl_event ev = enqueueBarrier(cq);   
+   gc_entity_release(ev);
+
+   return CL_SUCCESS;
+}

socl/src/cl_enqueuecopybuffer.c.inc

@@ -0,0 +1,109 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+struct arg_copybuffer {
+   size_t src_offset, dst_offset;
+   cl_mem src_buffer, dst_buffer;
+   size_t cb;
+};
+
+static void soclEnqueueCopyBuffer_opencl_task(void *descr[], void *args) {
+   struct arg_copybuffer *arg;
+   int wid;
+   cl_command_queue cq;
+   cl_event ev;
+
+   arg = (struct arg_copybuffer*)args;
+   wid = starpu_worker_get_id();
+   starpu_opencl_get_queue(wid, &cq);
+
+   cl_mem src = (cl_mem)STARPU_VARIABLE_GET_PTR(descr[0]);
+   cl_mem dst = (cl_mem)STARPU_VARIABLE_GET_PTR(descr[1]);
+
+   clEnqueueCopyBuffer(cq, src,dst, arg->src_offset, arg->dst_offset, arg->cb, 0, NULL, &ev);
+   clWaitForEvents(1, &ev);
+   clReleaseEvent(ev);
+
+   gc_entity_unstore(&arg->src_buffer);
+   gc_entity_unstore(&arg->dst_buffer);
+
+   free(arg);
+}
+
+static void soclEnqueueCopyBuffer_cpu_task(void *descr[], void *args) {
+   struct arg_copybuffer *arg;
+   arg = (struct arg_copybuffer*)args;
+   void * src = (void*)STARPU_VARIABLE_GET_PTR(descr[0]);
+   void * dst = (void*)STARPU_VARIABLE_GET_PTR(descr[1]);
+   memcpy(dst+arg->dst_offset, src+arg->src_offset, arg->cb);
+
+   gc_entity_unstore(&arg->src_buffer);
+   gc_entity_unstore(&arg->dst_buffer);
+
+   free(arg);
+}
+
+static starpu_codelet codelet_copybuffer = {
+   .where = STARPU_CPU | STARPU_OPENCL,
+   .model = NULL,
+   .cpu_func = &soclEnqueueCopyBuffer_cpu_task,
+   .opencl_func = &soclEnqueueCopyBuffer_opencl_task,
+   .nbuffers = 2
+};
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyBuffer(cl_command_queue  cq, 
+                    cl_mem              src_buffer,
+                    cl_mem              dst_buffer, 
+                    size_t              src_offset,
+                    size_t              dst_offset,
+                    size_t              cb, 
+                    cl_uint             num_events,
+                    const cl_event *    events,
+                    cl_event *          event) CL_API_SUFFIX__VERSION_1_0
+{
+   struct starpu_task *task;
+   struct arg_copybuffer *arg;
+   cl_event ev;
+
+   task = task_create(CL_COMMAND_COPY_BUFFER);
+   ev = task_event(task);
+
+   task->buffers[0].handle = src_buffer->handle;
+   task->buffers[0].mode = STARPU_R;
+   task->buffers[1].handle = dst_buffer->handle;
+   task->buffers[1].mode = STARPU_RW;
+   task->cl = &codelet_copybuffer;
+
+   arg = (struct arg_copybuffer*)malloc(sizeof(struct arg_copybuffer));
+   arg->src_offset = src_offset;
+   arg->dst_offset = dst_offset;
+   arg->cb = cb;
+   gc_entity_store(&arg->src_buffer, src_buffer);
+   gc_entity_store(&arg->dst_buffer, dst_buffer);
+   task->cl_arg = arg;
+   task->cl_arg_size = sizeof(struct arg_copybuffer);
+
+   dst_buffer->scratch = 0;
+
+   DEBUG_MSG("Submitting CopyBuffer task (event %d)\n", ev->id);
+
+   cl_int ret = command_queue_enqueue(cq, task, 0, num_events, events);
+
+   RETURN_EVENT(ev, event);
+
+   return ret;
+}

socl/src/cl_enqueuecopybuffertoimage.c.inc

@@ -0,0 +1,29 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyBufferToImage(cl_command_queue UNUSED(command_queue),
+                           cl_mem           UNUSED(src_buffer),
+                           cl_mem           UNUSED(dst_image), 
+                           size_t           UNUSED(src_offset),
+                           const size_t *   UNUSED(dst_origin),
+                           const size_t *   UNUSED(region), 
+                           cl_uint          UNUSED(num_events_in_wait_list),
+                           const cl_event * UNUSED(event_wait_list),
+                           cl_event *       UNUSED(event)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_enqueuecopyimage.c.inc

@@ -0,0 +1,29 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyImage(cl_command_queue   UNUSED(command_queue),
+                   cl_mem               UNUSED(src_image),
+                   cl_mem               UNUSED(dst_image), 
+                   const size_t *       UNUSED(src_origin),
+                   const size_t *       UNUSED(dst_origin),
+                   const size_t *       UNUSED(region), 
+                   cl_uint              UNUSED(num_events_in_wait_list),
+                   const cl_event *     UNUSED(event_wait_list),
+                   cl_event *           UNUSED(event)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_enqueuecopyimagetobuffer.c.inc

@@ -0,0 +1,29 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueCopyImageToBuffer(cl_command_queue UNUSED(command_queue),
+                           cl_mem           UNUSED(src_image),
+                           cl_mem           UNUSED(dst_buffer), 
+                           const size_t *   UNUSED(src_origin),
+                           const size_t *   UNUSED(region), 
+                           size_t           UNUSED(dst_offset),
+                           cl_uint          UNUSED(num_events_in_wait_list),
+                           const cl_event * UNUSED(event_wait_list),
+                           cl_event *       UNUSED(event)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_enqueuemapbuffer.c.inc

@@ -0,0 +1,89 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+struct mb_data {
+  cl_event ev;
+  cl_mem buffer;
+  cl_map_flags map_flags;
+};
+
+static void mapbuffer_callback(void *args) {
+  struct mb_data *arg = (struct mb_data*)args;
+
+  starpu_tag_notify_from_apps(arg->ev->id);
+  arg->ev->status = CL_COMPLETE;
+
+  gc_entity_unstore(&arg->ev);
+  gc_entity_unstore(&arg->buffer);
+  free(args);
+}
+
+static void mapbuffer_task(void *args) {
+  struct mb_data *arg = (struct mb_data*)args;
+
+  starpu_access_mode mode = (arg->map_flags == CL_MAP_READ ? STARPU_R : STARPU_RW);
+
+  starpu_data_acquire_cb(arg->buffer->handle, mode, mapbuffer_callback, arg);
+}
+
+CL_API_ENTRY void * CL_API_CALL
+soclEnqueueMapBuffer(cl_command_queue cq,
+                   cl_mem           buffer,
+                   cl_bool          blocking_map, 
+                   cl_map_flags     map_flags,
+                   size_t           offset, 
+                   size_t           UNUSED(cb),
+                   cl_uint          num_events,
+                   const cl_event * events,
+                   cl_event *       event,
+                   cl_int *         errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   struct starpu_task *task;
+   struct mb_data *arg;
+   cl_event ev;
+   cl_int err;
+
+   /* Create custom event that will be triggered when map is complete */
+   ev = event_create();
+
+   /* Store arguments */
+   arg = (struct mb_data*)malloc(sizeof(struct mb_data));
+   arg->map_flags = map_flags;
+   gc_entity_store(&arg->ev, ev);
+   gc_entity_store(&arg->buffer, buffer);
+
+   /* Create StarPU task */
+   task = task_create_cpu(CL_COMMAND_MAP_BUFFER, mapbuffer_task, arg, 0);
+   cl_event map_event = task_event(task);
+
+   /* Enqueue task */
+   DEBUG_MSG("Submitting MapBuffer task (event %d)\n", ev->id);
+   err = command_queue_enqueue_fakeevent(cq, task, 0, num_events, events, ev);
+   gc_entity_release(map_event);
+
+   if (errcode_ret != NULL)
+      *errcode_ret = err;
+
+   if (err != CL_SUCCESS)
+      return NULL;
+
+   if (blocking_map == CL_TRUE)
+      soclWaitForEvents(1, &ev);
+
+   RETURN_EVENT(ev, event);
+
+   return (void*)(starpu_variable_get_local_ptr(buffer->handle) + offset);
+}

socl/src/cl_enqueuemapimage.c.inc

@@ -0,0 +1,35 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY void * CL_API_CALL
+soclEnqueueMapImage(cl_command_queue  UNUSED(command_queue),
+                  cl_mem            UNUSED(image), 
+                  cl_bool           UNUSED(blocking_map), 
+                  cl_map_flags      UNUSED(map_flags), 
+                  const size_t *    UNUSED(origin),
+                  const size_t *    UNUSED(region),
+                  size_t *          UNUSED(image_row_pitch),
+                  size_t *          UNUSED(image_slice_pitch),
+                  cl_uint           UNUSED(num_events_in_wait_list),
+                  const cl_event *  UNUSED(event_wait_list),
+                  cl_event *        UNUSED(event),
+                  cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (errcode_ret != NULL)
+      *errcode_ret = CL_INVALID_OPERATION;
+
+   return NULL;
+}

socl/src/cl_enqueuemarker.c.inc

@@ -0,0 +1,28 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueMarker(cl_command_queue  cq,
+                cl_event *          event) CL_API_SUFFIX__VERSION_1_0
+{
+   if (event == NULL)
+      return CL_INVALID_VALUE;
+
+   starpu_task * task = task_create(CL_COMMAND_MARKER);
+   *event = task_event(task);
+
+   return command_queue_enqueue(cq, task, 0, 0, NULL);
+}

socl/src/cl_enqueuenativekernel.c.inc

@@ -0,0 +1,30 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueNativeKernel(cl_command_queue  UNUSED(command_queue),
+					       __attribute__((unused)) void (*user_func)(void *), 
+                      void *            UNUSED(args),
+                      size_t            UNUSED(cb_args), 
+                      cl_uint           UNUSED(num_mem_objects),
+                      const cl_mem *    UNUSED(mem_list),
+                      const void **     UNUSED(args_mem_loc),
+                      cl_uint           UNUSED(num_events_in_wait_list),
+                      const cl_event *  UNUSED(event_wait_list),
+                      cl_event *        UNUSED(event)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_enqueuendrangekernel.c.inc

@@ -0,0 +1,283 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+typedef struct running_kernel * running_kernel;
+
+struct running_kernel {
+  cl_kernel kernel;
+  cl_mem *buffers;
+  unsigned int buffer_count;
+  starpu_codelet *codelet;
+  cl_uint work_dim;
+  size_t * global_work_offset;
+  size_t * global_work_size;
+  size_t * local_work_size;
+  /* Arguments */
+  unsigned int arg_count;
+  size_t *arg_size;
+  enum kernel_arg_type  *arg_type;
+  void  **arg_value;
+};
+
+static void soclEnqueueNDRangeKernel_task(void *descr[], void *args) {
+   running_kernel d;
+   cl_command_queue cq;
+   int wid;
+   cl_int err;
+
+   d = (running_kernel)args;
+   wid = starpu_worker_get_id();
+   starpu_opencl_get_queue(wid, &cq);
+
+   DEBUG_MSG("[worker %d] [kernel %d] Executing kernel...\n", wid, d->kernel->id);
+
+   int range = starpu_worker_get_range();
+
+   /* Set arguments */
+   {
+      unsigned int i;
+      int buf = 0;
+      for (i=0; i<d->arg_count; i++) {
+         switch (d->arg_type[i]) {
+            case Null:
+               err = clSetKernelArg(d->kernel->cl_kernels[range], i, d->arg_size[i], NULL);
+               break;
+            case Buffer: {
+                  cl_mem mem;  
+                  mem = (cl_mem)STARPU_VARIABLE_GET_PTR(descr[buf]);
+                  err = clSetKernelArg(d->kernel->cl_kernels[range], i, d->arg_size[i], &mem);
+                  buf++;
+               }
+               break;
+            case Immediate:
+               err = clSetKernelArg(d->kernel->cl_kernels[range], i, d->arg_size[i], d->arg_value[i]);
+               break;
+         }
+         if (err != CL_SUCCESS) {
+            DEBUG_CL("clSetKernelArg", err);
+            DEBUG_ERROR("Aborting\n");
+         }
+      }
+   }
+
+   /* Calling Kernel */
+   cl_event event;
+   err = clEnqueueNDRangeKernel(cq, d->kernel->cl_kernels[range], d->work_dim, d->global_work_offset, d->global_work_size, d->local_work_size, 0, NULL, &event);
+
+   if (err != CL_SUCCESS) {
+      ERROR_MSG("Worker[%d] Unable to Enqueue kernel (error %d)\n", wid, err);
+      DEBUG_CL("clEnqueueNDRangeKernel", err);
+      DEBUG_MSG("Workdim %d, global_work_offset %p, global_work_size %p, local_work_size %p\n",
+                d->work_dim, d->global_work_offset, d->global_work_size, d->local_work_size);
+      DEBUG_MSG("Global work size: %ld %ld %ld\n", d->global_work_size[0],
+            (d->work_dim > 1 ? d->global_work_size[1] : 1), (d->work_dim > 2 ? d->global_work_size[2] : 1)); 
+      if (d->local_work_size != NULL)
+         DEBUG_MSG("Local work size: %ld %ld %ld\n", d->local_work_size[0],
+               (d->work_dim > 1 ? d->local_work_size[1] : 1), (d->work_dim > 2 ? d->local_work_size[2] : 1)); 
+      ERROR_MSG("Aborting.\n");
+      exit(1);
+   }
+
+   /* Waiting for kernel to terminate */
+   clWaitForEvents(1, &event);
+   clReleaseEvent(event);
+}
+
+static void cleaning_task_callback(void *args) {
+   running_kernel arg = (running_kernel)args;
+
+   free(arg->arg_size);
+   free(arg->arg_type);
+
+   unsigned int i;
+   for (i=0; i<arg->arg_count; i++) {
+      free(arg->arg_value[i]);
+   }
+   free(arg->arg_value);
+
+   for (i=0; i<arg->buffer_count; i++)
+      gc_entity_unstore(&arg->buffers[i]);
+
+   gc_entity_unstore(&arg->kernel);
+
+   free(arg->buffers);
+   free(arg->global_work_offset);
+   free(arg->global_work_size);
+   free(arg->local_work_size);
+   void * co = arg->codelet;
+   arg->codelet = NULL;
+   free(co);
+}
+
+static struct starpu_perfmodel_t perf_model = {
+  .type = STARPU_HISTORY_BASED,
+  .symbol = "perf_model"
+};
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueNDRangeKernel(cl_command_queue cq,
+                       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,
+                       const cl_event * events,
+                       cl_event *       event) CL_API_SUFFIX__VERSION_1_0
+{
+   struct starpu_task *task;
+   running_kernel arg;
+   starpu_codelet *codelet;
+   cl_event ev;
+
+   /***********************
+    * Allocate structures *
+    ***********************/
+
+   /* Codelet */
+   codelet = (starpu_codelet*)malloc(sizeof(starpu_codelet));
+   if (codelet == NULL)
+      return CL_OUT_OF_HOST_MEMORY;
+
+   /* Codelet arguments */
+   arg = (running_kernel)malloc(sizeof(struct running_kernel));
+   if (arg == NULL) {
+      free(codelet);
+      return CL_OUT_OF_HOST_MEMORY;
+   }
+
+   /* StarPU task */
+   task = task_create(&work_dim != &et_work_dim ? CL_COMMAND_NDRANGE_KERNEL: CL_COMMAND_TASK);
+   ev = task_event(task);
+
+   /*******************
+    * Initializations *
+    *******************/
+
+   /* ------- *
+    * Codelet *
+    * ------- */
+   codelet->where = STARPU_OPENCL;
+   codelet->power_model = NULL;
+   codelet->opencl_func = &soclEnqueueNDRangeKernel_task;
+   //codelet->model = NULL;
+   codelet->model = &perf_model;
+
+   /* ---------------- *
+    * Codelet argument *
+    * ---------------- */
+   gc_entity_store(&arg->kernel, kernel);
+   arg->work_dim = work_dim;
+   arg->codelet = codelet;
+
+   /* Global work offset */
+   if (global_work_offset != NULL) {
+      arg->global_work_offset = (size_t*)malloc(sizeof(size_t)*work_dim);
+      memcpy(arg->global_work_offset, global_work_offset, work_dim*sizeof(size_t));
+   }
+   else arg->global_work_offset = NULL;
+
+   /* Global work size */
+   arg->global_work_size = (size_t*)malloc(sizeof(size_t)*work_dim);
+   memcpy(arg->global_work_size, global_work_size, work_dim*sizeof(size_t));
+
+   /* Local work size */
+   if (local_work_size != NULL) {
+      arg->local_work_size = (size_t*)malloc(sizeof(size_t)*work_dim);
+      memcpy(arg->local_work_size, local_work_size, work_dim*sizeof(size_t));
+   }
+   else arg->local_work_size = NULL;
+
+   /* ----------- *
+    * StarPU task *
+    * ----------- */
+   task->cl = codelet;
+   task->cl_arg = arg;
+   task->cl_arg_size = sizeof(struct running_kernel);
+
+   /* Convert OpenCL's memory objects to StarPU buffers */
+   codelet->nbuffers = 0;
+   {
+      arg->buffers = malloc(sizeof(cl_mem) * kernel->arg_count);
+      arg->buffer_count = 0;
+
+      unsigned int i;
+      for (i=0; i<kernel->arg_count; i++) {
+         if (kernel->arg_type[i] == Buffer) {
+
+            cl_mem buf = (cl_mem)kernel->arg_value[i];
+
+            /* We save cl_mem references in order to properly release them after kernel termination */
+            gc_entity_store(&arg->buffers[arg->buffer_count], buf);
+            arg->buffer_count += 1;
+
+            codelet->nbuffers++;
+            task->buffers[codelet->nbuffers-1].handle = buf->handle;
+
+            /* Determine best StarPU buffer access mode */
+            int mode;
+            if (buf->mode == CL_MEM_READ_ONLY)
+               mode = STARPU_R;
+            else if (buf->mode == CL_MEM_WRITE_ONLY) {
+               mode = STARPU_W;
+               buf->scratch = 0;
+            }
+            else if (buf->scratch) { //RW but never accessed in RW or W mode
+               mode = STARPU_W;
+               buf->scratch = 0;
+            }
+            else {
+               mode = STARPU_RW;
+               buf->scratch = 0;
+            }
+            task->buffers[codelet->nbuffers-1].mode = mode; 
+         }
+      }
+   }
+
+   /* Copy arguments as kernel args can be modified by the time we launch the kernel */
+   {
+      arg->arg_count = kernel->arg_count;
+      arg->arg_size = malloc(sizeof(size_t) * kernel->arg_count);
+      memcpy(arg->arg_size, kernel->arg_size, sizeof(size_t) * kernel->arg_count);
+      arg->arg_type = malloc(sizeof(enum kernel_arg_type) * kernel->arg_count);
+      memcpy(arg->arg_type, kernel->arg_type, sizeof(enum kernel_arg_type) * kernel->arg_count);
+      arg->arg_value = malloc(sizeof(void*) * kernel->arg_count);
+      unsigned int i;
+      for (i=0; i<kernel->arg_count; i++) {
+         if (kernel->arg_value[i] != NULL) {
+           arg->arg_value[i] = malloc(arg->arg_size[i]);
+           memcpy(arg->arg_value[i], kernel->arg_value[i], arg->arg_size[i]);
+         }
+         else arg->arg_value[i] = NULL;
+      }
+   }
+
+   DEBUG_MSG("Submitting NDRange task (event %d)\n", ev->id);
+
+   cl_int ret = command_queue_enqueue(cq, task, 0, num_events, events);
+
+   /* Enqueue a cleaning task */
+   starpu_task * cleaning_task = task_create_cpu(0, cleaning_task_callback, arg,1);
+   cl_event cleaning_event = task_event(cleaning_task);
+   command_queue_enqueue(cq, cleaning_task, 0, 1, &ev);
+
+   gc_entity_release(cleaning_event);
+  
+   RETURN_EVENT(ev, event);
+
+   return ret;
+}

socl/src/cl_enqueuereadbuffer.c.inc

@@ -0,0 +1,107 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+struct arg_readbuffer {
+   size_t offset;
+   size_t cb;
+   void * ptr;
+   cl_mem buffer;
+};
+
+static void soclEnqueueReadBuffer_cpu_task(void *descr[], void *args) {
+   struct arg_readbuffer *arg;
+   arg = (struct arg_readbuffer*)args;
+   void * ptr = (void*)STARPU_VARIABLE_GET_PTR(descr[0]);
+   DEBUG_MSG("[Buffer %d] Reading %ld bytes from %p to %p\n", arg->buffer->id, arg->cb, ptr+arg->offset, arg->ptr);
+
+   //This fix is for people who use USE_HOST_PTR and still use ReadBuffer to sync the buffer in host mem at host_ptr.
+   //They should use buffer mapping facilities instead.
+   if (ptr+arg->offset != arg->ptr)
+      memcpy(arg->ptr, ptr+arg->offset, arg->cb);
+
+   gc_entity_unstore(&arg->buffer);
+   free(args);
+}
+
+static void soclEnqueueReadBuffer_opencl_task(void *descr[], void *args) {
+   struct arg_readbuffer *arg;
+   arg = (struct arg_readbuffer*)args;
+
+   cl_mem mem = (cl_mem)STARPU_VARIABLE_GET_PTR(descr[0]);
+
+   DEBUG_MSG("[Buffer %d] Reading %ld bytes from offset %ld into %p\n", arg->buffer->id, arg->cb, arg->offset, arg->ptr);
+
+   int wid = starpu_worker_get_id();
+   cl_command_queue cq;
+   starpu_opencl_get_queue(wid, &cq);
+
+   cl_int ret = clEnqueueReadBuffer(cq, mem, CL_TRUE, arg->offset, arg->cb, arg->ptr, 0, NULL, NULL);
+   if (ret != CL_SUCCESS)
+      DEBUG_CL("clEnqueueReadBuffer", ret);
+
+   gc_entity_unstore(&arg->buffer);
+   free(args);
+}
+
+static starpu_codelet codelet_readbuffer = {
+   .where = STARPU_OPENCL,
+   .model = NULL,
+   .cpu_func = &soclEnqueueReadBuffer_cpu_task,
+   .opencl_func = &soclEnqueueReadBuffer_opencl_task,
+   .nbuffers = 1
+};
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueReadBuffer(cl_command_queue  cq,
+                    cl_mem              buffer,
+                    cl_bool             blocking,
+                    size_t              offset,
+                    size_t              cb, 
+                    void *              ptr,
+                    cl_uint             num_events,
+                    const cl_event *    events,
+                    cl_event *          event) CL_API_SUFFIX__VERSION_1_0
+{ 
+   struct starpu_task *task;
+   struct arg_readbuffer *arg;
+   cl_event ev;
+
+   task = task_create(CL_COMMAND_READ_BUFFER);
+   ev = task_event(task);
+
+   task->buffers[0].handle = buffer->handle;
+   task->buffers[0].mode = STARPU_R;
+   task->cl = &codelet_readbuffer;
+
+   arg = (struct arg_readbuffer*)malloc(sizeof(struct arg_readbuffer));
+   arg->offset = offset;
+   arg->cb = cb;
+   arg->ptr = ptr;
+   task->cl_arg = arg;
+   task->cl_arg_size = sizeof(struct arg_readbuffer);
+
+   gc_entity_store(&arg->buffer, buffer);
+
+   task->synchronous = (blocking == CL_TRUE);
+
+   DEBUG_MSG("Submitting EnqueueRWBuffer task (event %d)\n", ev->id);
+
+   cl_int ret = command_queue_enqueue(cq, task, 0, num_events, events);
+
+   RETURN_EVENT(ev, event);
+
+   return ret;
+}

socl/src/cl_enqueuereadimage.c.inc

@@ -0,0 +1,31 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueReadImage(cl_command_queue   UNUSED(command_queue),
+                   cl_mem               UNUSED(image),
+                   cl_bool              UNUSED(blocking_read), 
+                   const size_t *       UNUSED(origin),
+                   const size_t *       UNUSED(region),
+                   size_t               UNUSED(row_pitch),
+                   size_t               UNUSED(slice_pitch), 
+                   void *               UNUSED(ptr),
+                   cl_uint              UNUSED(num_events_in_wait_list),
+                   const cl_event *     UNUSED(event_wait_list),
+                   cl_event *           UNUSED(event)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_enqueuetask.c.inc

@@ -0,0 +1,33 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+static cl_uint et_work_dim = 3;
+static const size_t et_global_work_offset[3] = {0,0,0};
+static const size_t et_global_work_size[3] = {1,1,1};
+static const size_t * et_local_work_size = NULL;
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueNDRangeKernel(cl_command_queue, cl_kernel, cl_uint, const size_t *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *) CL_API_SUFFIX__VERSION_1_0;
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueTask(cl_command_queue cq,
+              cl_kernel         kernel,
+              cl_uint           num_events,
+              const cl_event *  events,
+              cl_event *        event) CL_API_SUFFIX__VERSION_1_0
+{
+   return soclEnqueueNDRangeKernel(cq, kernel, et_work_dim, et_global_work_offset, et_global_work_size, et_local_work_size, num_events, events, event); 
+}

socl/src/cl_enqueueunmapmemobject.c.inc

@@ -0,0 +1,40 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueUnmapMemObject(cl_command_queue cq,
+                        cl_mem            memobj,
+                        void *            UNUSED(mapped_ptr),
+                        cl_uint           num_events,
+                        const cl_event *  events,
+                        cl_event *        event) CL_API_SUFFIX__VERSION_1_0
+{
+   struct starpu_task *task;
+   cl_int err;
+   cl_event ev;
+
+   /* Create StarPU task */
+   task = task_create_cpu(CL_COMMAND_UNMAP_MEM_OBJECT, (void(*)(void*))starpu_data_release, memobj->handle, 0);
+   ev = task_event(task);
+
+   DEBUG_MSG("Submitting UnmapBuffer task (event %d)\n", task->tag_id);
+
+   err = command_queue_enqueue(cq, task, 0, num_events, events);
+
+   RETURN_EVENT(ev, event);
+
+   return err;
+}

socl/src/cl_enqueuewaitforevents.c.inc

@@ -0,0 +1,29 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueWaitForEvents(cl_command_queue cq,
+                       cl_uint          num_events,
+                       const cl_event * events) CL_API_SUFFIX__VERSION_1_0
+{
+
+   //CL_COMMAND_MARKER has been chosen as CL_COMMAND_WAIT_FOR_EVENTS doesn't exist
+   starpu_task * task = task_create(CL_COMMAND_MARKER);
+
+   command_queue_enqueue(cq, task, 0, num_events, events);
+
+   return CL_SUCCESS;
+}

socl/src/cl_enqueuewritebuffer.c.inc

@@ -0,0 +1,117 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+struct arg_writebuffer {
+   size_t offset;
+   size_t cb;
+   const void * ptr;
+   cl_mem buffer;
+};
+
+static void soclEnqueueWriteBuffer_cpu_task(void *descr[], void *args) {
+   struct arg_writebuffer *arg;
+   arg = (struct arg_writebuffer*)args;
+   void * ptr = (void*)STARPU_VARIABLE_GET_PTR(descr[0]);
+   DEBUG_MSG("[Buffer %d] Writing %ld bytes from %p to %p\n", arg->buffer->id, arg->cb, arg->ptr, ptr+arg->offset);
+
+   //FIXME: Fix for people who use USE_HOST_PTR, modify data at host_ptr and use WriteBuffer to commit the change.
+   // StarPU may have erased host mem at host_ptr (for instance by retrieving current buffer data at host_ptr)
+   // Buffer mapping facilities should be used instead
+   // Maybe we should report the bug here... for now, we just avoid memcpy crash due to overlapping regions...
+   if (ptr+arg->offset != arg->ptr)
+      memcpy(ptr+arg->offset, arg->ptr, arg->cb);
+
+   gc_entity_unstore(&arg->buffer);
+   free(args);
+}
+
+static void soclEnqueueWriteBuffer_opencl_task(void *descr[], void *args) {
+   struct arg_writebuffer *arg;
+   arg = (struct arg_writebuffer*)args;
+
+   cl_mem mem = (cl_mem)STARPU_VARIABLE_GET_PTR(descr[0]);
+
+   DEBUG_MSG("[Buffer %d] Writing %ld bytes to offset %ld from %p\n", arg->buffer->id, arg->cb, arg->offset, arg->ptr);
+
+   int wid = starpu_worker_get_id();
+   cl_command_queue cq;
+   starpu_opencl_get_queue(wid, &cq);
+
+   cl_int err = clEnqueueWriteBuffer(cq, mem, CL_TRUE, arg->offset, arg->cb, arg->ptr, 0, NULL, NULL);
+   if (err != CL_SUCCESS)
+      DEBUG_CL("clEnqueueWriteBuffer", err);
+
+   gc_entity_unstore(&arg->buffer);
+   free(args);
+}
+
+static starpu_codelet codelet_writebuffer = {
+   .where = STARPU_OPENCL,
+   .model = NULL,
+   .cpu_func = &soclEnqueueWriteBuffer_cpu_task,
+   .opencl_func = &soclEnqueueWriteBuffer_opencl_task,
+   .nbuffers = 1
+};
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueWriteBuffer(cl_command_queue cq, 
+                     cl_mem             buffer, 
+                     cl_bool            blocking, 
+                     size_t             offset, 
+                     size_t             cb, 
+                     const void *       ptr, 
+                     cl_uint            num_events, 
+                     const cl_event *   events, 
+                     cl_event *         event) CL_API_SUFFIX__VERSION_1_0
+{ 
+   struct starpu_task *task;
+   struct arg_writebuffer *arg;
+   cl_event ev;
+
+   task = task_create(CL_COMMAND_WRITE_BUFFER);
+   ev = task_event(task);
+
+   task->buffers[0].handle = buffer->handle;
+   //If only a subpart of the buffer is written, RW access mode is required
+   if (cb != buffer->size)
+      task->buffers[0].mode = STARPU_RW;
+   else 
+      task->buffers[0].mode = STARPU_W;
+   task->cl = &codelet_writebuffer;
+
+   arg = (struct arg_writebuffer*)malloc(sizeof(struct arg_writebuffer));
+   arg->offset = offset;
+   arg->cb = cb;
+   arg->ptr = ptr;
+   task->cl_arg = arg;
+   task->cl_arg_size = sizeof(struct arg_writebuffer);
+
+   gc_entity_store(&arg->buffer, buffer);
+
+   //The buffer now contains meaningful data
+   arg->buffer->scratch = 0;
+
+   task->synchronous = (blocking == CL_TRUE);
+
+   DEBUG_MSG("Submitting EnqueueRWBuffer task (event %d)\n", ev->id);
+
+   cl_int ret = command_queue_enqueue(cq, task, 0, num_events,events);
+
+   /* Return retained event if required by user */
+   RETURN_EVENT(ev,event);
+
+   return ret;
+}

socl/src/cl_enqueuewriteimage.c.inc

@@ -0,0 +1,31 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclEnqueueWriteImage(cl_command_queue  UNUSED(command_queue),
+                    cl_mem              UNUSED(image),
+                    cl_bool             UNUSED(blocking_write), 
+                    const size_t *      UNUSED(origin),
+                    const size_t *      UNUSED(region),
+                    size_t              UNUSED(input_row_pitch),
+                    size_t              UNUSED(input_slice_pitch), 
+                    const void *        UNUSED(ptr),
+                    cl_uint             UNUSED(num_events_in_wait_list),
+                    const cl_event *    UNUSED(event_wait_list),
+                    cl_event *          UNUSED(event)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_finish.c.inc

@@ -0,0 +1,25 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclFinish(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
+{
+   cl_event ev = enqueueBarrier(cq);
+   soclWaitForEvents(1, &ev);
+   gc_entity_release(ev);
+
+   return CL_SUCCESS;
+}

socl/src/cl_flush.c.inc

@@ -0,0 +1,21 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclFlush(cl_command_queue UNUSED(command_queue)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_SUCCESS;
+}

socl/src/cl_getcommandqueueinfo.c.inc

@@ -0,0 +1,37 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetCommandQueueInfo(cl_command_queue    cq,
+                      cl_command_queue_info param_name,
+                      size_t                param_value_size,
+                      void *                param_value,
+                      size_t *              param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (cq == NULL)
+      return CL_INVALID_COMMAND_QUEUE;
+
+   switch (param_name) {
+      INFO_CASE(CL_QUEUE_CONTEXT, cq->context);
+      INFO_CASE(CL_QUEUE_DEVICE, cq->device);
+      INFO_CASE(CL_QUEUE_REFERENCE_COUNT, cq->_entity.refs);
+      INFO_CASE(CL_QUEUE_PROPERTIES, cq->properties);
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   return CL_SUCCESS; 
+}

socl/src/cl_getcontextinfo.c.inc

@@ -0,0 +1,36 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetContextInfo(cl_context       context, 
+                 cl_context_info    param_name, 
+                 size_t             param_value_size, 
+                 void *             param_value, 
+                 size_t *           param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (context == NULL)
+      return CL_INVALID_CONTEXT;
+   
+   switch (param_name) {
+      INFO_CASE(CL_CONTEXT_REFERENCE_COUNT, context->_entity.refs);
+      INFO_CASE_EX(CL_CONTEXT_DEVICES, context->devices, context->num_devices * sizeof(cl_device_id));
+      INFO_CASE_EX(CL_CONTEXT_PROPERTIES, context->properties, context->num_properties * sizeof(cl_device_id));
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_getdeviceids.c.inc

@@ -0,0 +1,56 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+/**
+ * \brief Return one device of each kind
+ *
+ * \param[in] platform Must be StarPU platform ID or NULL
+ */
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetDeviceIDs(cl_platform_id   platform,
+               cl_device_type   device_type, 
+               cl_uint          num_entries, 
+               cl_device_id *   devices, 
+               cl_uint *        num_devices) CL_API_SUFFIX__VERSION_1_0
+{
+   if (platform != NULL && platform != &socl_platform)
+      return CL_INVALID_PLATFORM;
+
+   if ((devices != NULL && num_entries == 0)
+      || (devices == NULL && num_devices == NULL))
+      return CL_INVALID_VALUE;
+
+   if (!(device_type & (CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU | CL_DEVICE_TYPE_ACCELERATOR | CL_DEVICE_TYPE_DEFAULT))
+      && (device_type != CL_DEVICE_TYPE_ALL))
+      return CL_INVALID_DEVICE_TYPE;
+
+   {
+      int i;
+      unsigned int num = 0;
+      for (i=0; i<socl_device_count; i++) {
+         if (socl_devices[i].type & device_type) {
+            if (devices != NULL && num < num_entries)
+               devices[num] = (cl_device_id)&socl_devices[i];
+            num++;
+         }
+      }
+      if (num_devices != NULL)
+         *num_devices = num;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_getdeviceinfo.c.inc

@@ -0,0 +1,90 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+/**
+ * \brief Return dummy infos
+ */
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetDeviceInfo(cl_device_id    device,
+                cl_device_info  param_name, 
+                size_t          param_value_size, 
+                void *          param_value,
+                size_t *        param_value_size_ret) CL_API_SUFFIX__VERSION_1_0 
+{
+   int i, found=0;
+   for (i=0; i<socl_device_count; i++) {
+      if (device == &socl_devices[i]) {
+        found = 1;
+        break;
+      }
+   }
+
+   if (!found)
+      return CL_INVALID_DEVICE;
+
+   switch (param_name) {
+      INFO_CASE(CL_DEVICE_TYPE, device->type)
+      INFO_CASE(CL_DEVICE_VENDOR_ID, SOCL_DEVICE_VENDOR_ID)
+      INFO_CASE(CL_DEVICE_MAX_COMPUTE_UNITS, device->max_compute_units)
+      INFO_CASE(CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, device->max_work_item_dimensions)
+      INFO_CASE(CL_DEVICE_MAX_WORK_ITEM_SIZES, device->max_work_item_sizes)
+      INFO_CASE(CL_DEVICE_MAX_WORK_GROUP_SIZE, device->max_work_group_size)
+      INFO_CASE(CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR, device->preferred_vector_widths[0])
+      INFO_CASE(CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, device->preferred_vector_widths[1])
+      INFO_CASE(CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, device->preferred_vector_widths[2])
+      INFO_CASE(CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG, device->preferred_vector_widths[3])
+      INFO_CASE(CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, device->preferred_vector_widths[4])
+      INFO_CASE(CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, device->preferred_vector_widths[5])
+      INFO_CASE(CL_DEVICE_MAX_CLOCK_FREQUENCY, device->max_clock_frequency)
+      INFO_CASE(CL_DEVICE_ADDRESS_BITS, device->address_bits)
+      INFO_CASE(CL_DEVICE_MAX_MEM_ALLOC_SIZE, device->max_mem_alloc_size)
+      INFO_CASE(CL_DEVICE_IMAGE_SUPPORT, device->image_support)
+      INFO_CASE(CL_DEVICE_MAX_PARAMETER_SIZE, device->max_parameter_size)
+      INFO_CASE(CL_DEVICE_MEM_BASE_ADDR_ALIGN, device->mem_base_addr_align)
+      INFO_CASE(CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE, device->min_data_type_align_size)
+      INFO_CASE(CL_DEVICE_SINGLE_FP_CONFIG, device->single_fp_config)
+      INFO_CASE(CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, device->global_mem_cache_type)
+      INFO_CASE(CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, device->global_mem_cacheline_size)
+      INFO_CASE(CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, device->global_mem_cache_size)
+      INFO_CASE(CL_DEVICE_GLOBAL_MEM_SIZE, device->global_mem_size)
+      INFO_CASE(CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, device->max_constant_buffer_size)
+      INFO_CASE(CL_DEVICE_MAX_CONSTANT_ARGS, device->max_constant_args)
+      INFO_CASE(CL_DEVICE_LOCAL_MEM_TYPE, device->local_mem_type)
+      INFO_CASE(CL_DEVICE_LOCAL_MEM_SIZE, device->local_mem_size)
+      INFO_CASE(CL_DEVICE_ERROR_CORRECTION_SUPPORT, device->error_correction_support)
+      INFO_CASE(CL_DEVICE_PROFILING_TIMER_RESOLUTION, device->profiling_timer_resolution)
+      INFO_CASE(CL_DEVICE_ENDIAN_LITTLE, device->endian_little)
+      INFO_CASE(CL_DEVICE_AVAILABLE, device->available)
+      INFO_CASE(CL_DEVICE_COMPILER_AVAILABLE, device->compiler_available)
+      INFO_CASE(CL_DEVICE_EXECUTION_CAPABILITIES, device->execution_capabilities)
+      INFO_CASE(CL_DEVICE_QUEUE_PROPERTIES, device->queue_properties)
+      case CL_DEVICE_PLATFORM: {
+         cl_platform_id p = &socl_platform;
+         INFO_CASE_EX2(p);
+      }
+      INFO_CASE(CL_DEVICE_NAME, device->name)
+      INFO_CASE(CL_DEVICE_VENDOR, SOCL_VENDOR)
+      INFO_CASE(CL_DRIVER_VERSION, SOCL_DRIVER_VERSION)
+      INFO_CASE(CL_DEVICE_PROFILE, SOCL_PROFILE)
+      INFO_CASE(CL_DEVICE_VERSION, SOCL_VERSION)
+      INFO_CASE(CL_DEVICE_EXTENSIONS, device->extensions)
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_geteventinfo.c.inc

@@ -0,0 +1,41 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetEventInfo(cl_event       event,
+               cl_event_info    param_name,
+               size_t           param_value_size,
+               void *           param_value,
+               size_t *         param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (event == NULL)
+      return CL_INVALID_EVENT;
+
+   #define STAT_CASE(starpu,opencl) case starpu: \
+      status = opencl; \
+      break;
+
+   switch (param_name) {
+      INFO_CASE(CL_EVENT_COMMAND_QUEUE, event->cq);
+      INFO_CASE(CL_EVENT_COMMAND_TYPE, event->type);
+      INFO_CASE(CL_EVENT_COMMAND_EXECUTION_STATUS, event->status);
+      INFO_CASE(CL_EVENT_REFERENCE_COUNT, event->_entity.refs);
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   return CL_SUCCESS; 
+}

socl/src/cl_geteventprofilinginfo.c.inc

@@ -0,0 +1,43 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetEventProfilingInfo(cl_event          event,
+                        cl_profiling_info   param_name,
+                        size_t              param_value_size,
+                        void *              param_value,
+                        size_t *            param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   struct starpu_task_profiling_info * prof = event->profiling_info;
+
+   if (prof == NULL)
+      return CL_PROFILING_INFO_NOT_AVAILABLE;
+
+   #define TONANO(t) (t.tv_nsec + t.tv_sec*1e9)
+
+   switch (param_name) {
+      case CL_PROFILING_COMMAND_QUEUED:
+      INFO_CASE_VALUE(CL_PROFILING_COMMAND_SUBMIT, cl_ulong, TONANO(prof->submit_time));
+      INFO_CASE_VALUE(CL_PROFILING_COMMAND_START, cl_ulong, TONANO(prof->start_time));
+      INFO_CASE_VALUE(CL_PROFILING_COMMAND_END, cl_ulong, TONANO(prof->end_time));
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   #undef TONANO
+
+   return CL_SUCCESS;
+}

socl/src/cl_getextensionfunctionaddress.c.inc

@@ -0,0 +1,22 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY void * CL_API_CALL
+soclGetExtensionFunctionAddress(const char * UNUSED(func_name)) CL_API_SUFFIX__VERSION_1_0
+{
+   //TODO
+   return NULL;
+}

socl/src/cl_getimageinfo.c.inc

@@ -0,0 +1,25 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetImageInfo(cl_mem           UNUSED(image),
+               cl_image_info    UNUSED(param_name), 
+               size_t           UNUSED(param_value_size),
+               void *           UNUSED(param_value),
+               size_t *         UNUSED(param_value_size_ret)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_getkernelinfo.c.inc

@@ -0,0 +1,38 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetKernelInfo(cl_kernel       kernel,
+                cl_kernel_info  param_name,
+                size_t          param_value_size,
+                void *          param_value,
+                size_t *        param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (kernel == NULL)
+      return CL_INVALID_KERNEL;
+
+   switch (param_name) {
+      INFO_CASE_EX(CL_KERNEL_FUNCTION_NAME, kernel->kernel_name, strlen(kernel->kernel_name)+1)
+      INFO_CASE(CL_KERNEL_NUM_ARGS, kernel->arg_count)
+      INFO_CASE(CL_KERNEL_REFERENCE_COUNT, kernel->_entity.refs)
+      INFO_CASE(CL_KERNEL_PROGRAM, kernel->program)
+      INFO_CASE(CL_KERNEL_CONTEXT, kernel->program->context)
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_getkernelworkgroupinfo.c.inc

@@ -0,0 +1,132 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+struct gkwgi_data {
+   cl_kernel_work_group_info param_name;
+   cl_kernel kernel;
+   union {
+      size_t work_group_size;
+      size_t compile_work_group_size[3];
+      cl_ulong local_mem_size;
+   };
+};
+
+static void gkwgi_task(void *data) {
+   cl_int err;
+   struct gkwgi_data *d = (struct gkwgi_data*)data;
+
+   int wid = starpu_worker_get_id();
+   int range = starpu_worker_get_range();
+
+   cl_device_id device;
+   starpu_opencl_get_device(wid, &device);
+
+   size_t value;
+   size_t oldval;
+   err = clGetKernelWorkGroupInfo(d->kernel->cl_kernels[range], device, d->param_name, sizeof(value), &value, NULL);
+   if (err != CL_SUCCESS) {
+      DEBUG_MSG("Worker [%d] failed\n", wid);
+      DEBUG_CL("clGetKernelWorkGroupInfo", err);
+   }
+
+   switch (d->param_name) {
+      case CL_KERNEL_WORK_GROUP_SIZE: {
+         //Get the smallest work group size
+         do {
+            oldval = d->work_group_size;
+         } while (value < oldval && !(__sync_bool_compare_and_swap(&d->work_group_size, oldval, value)));
+      }
+      break;
+      case CL_KERNEL_LOCAL_MEM_SIZE: {
+         //Get the biggest local mem size
+         do {
+            oldval = d->local_mem_size;
+         } while (value > oldval && !(__sync_bool_compare_and_swap(&d->local_mem_size, oldval, value)));
+      }
+      break;
+   }
+
+}
+
+static void gkwgi_task2(void **UNUSED(desc), void *data) {
+   cl_int err;
+   struct gkwgi_data *d = (struct gkwgi_data*)data;
+
+   int wid = starpu_worker_get_id();
+   int range = starpu_worker_get_range();
+
+   cl_device_id device;
+   starpu_opencl_get_device(wid, &device);
+
+   err = clGetKernelWorkGroupInfo(d->kernel->cl_kernels[range], device, d->param_name, sizeof(d->compile_work_group_size), &d->compile_work_group_size, NULL);
+   if (err != CL_SUCCESS) {
+      DEBUG_MSG("Worker [%d] failed\n", wid);
+      DEBUG_CL("clGetKernelWorkGroupInfo", err);
+   }
+}
+
+static starpu_codelet gkwgi_codelet = {
+   .where = STARPU_OPENCL,
+   .opencl_func = gkwgi_task2,
+   .nbuffers = 0,
+   .model = NULL
+};
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetKernelWorkGroupInfo(cl_kernel                kernel,
+                         cl_device_id               UNUSED(device),
+                         cl_kernel_work_group_info  param_name,
+                         size_t                     param_value_size,
+                         void *                     param_value,
+                         size_t *                   param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (kernel == NULL)
+      return CL_INVALID_KERNEL;
+
+
+   struct gkwgi_data data;
+   data.param_name = param_name;
+   data.kernel = kernel;
+
+   switch (param_name) {
+      case CL_KERNEL_WORK_GROUP_SIZE:
+         /* We take the smallest value to be sure the kernel can be executed on any available device */
+         data.work_group_size = SIZE_MAX;
+         starpu_execute_on_each_worker(gkwgi_task, &data, STARPU_OPENCL);
+         INFO_CASE_EX2(data.work_group_size);
+      case CL_KERNEL_COMPILE_WORK_GROUP_SIZE: {
+         struct starpu_task *task;
+         task = starpu_task_create();
+         task->cl = &gkwgi_codelet;
+         task->cl_arg = &data;
+         task->cl_arg_size = sizeof(data);
+         task->synchronous = 1;
+         starpu_task_submit(task);
+         INFO_CASE_EX2(data.compile_work_group_size);
+         }
+      case CL_KERNEL_LOCAL_MEM_SIZE:
+         /* We take the biggest value to be sure the kernel can be executed on any available device */
+         data.local_mem_size = 0;
+         starpu_execute_on_each_worker(gkwgi_task, &data, STARPU_OPENCL);
+         INFO_CASE_EX2(data.local_mem_size);
+      default:
+         return CL_INVALID_OPERATION;
+   }
+
+   return CL_SUCCESS;
+}
+
+

socl/src/cl_getmemobjectinfo.c.inc

@@ -0,0 +1,39 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetMemObjectInfo(cl_mem           mem,
+                   cl_mem_info      param_name, 
+                   size_t           param_value_size,
+                   void *           param_value,
+                   size_t *         param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   static cl_mem_object_type mot = CL_MEM_OBJECT_BUFFER;
+
+   switch (param_name) {
+      INFO_CASE(CL_MEM_TYPE, mot)
+      INFO_CASE(CL_MEM_FLAGS, mem->flags)
+      INFO_CASE(CL_MEM_SIZE, mem->size)
+      INFO_CASE(CL_MEM_HOST_PTR, mem->host_ptr)
+      INFO_CASE(CL_MEM_MAP_COUNT, mem->map_count)
+      INFO_CASE(CL_MEM_REFERENCE_COUNT, mem->_entity.refs)
+      INFO_CASE(CL_MEM_CONTEXT, mem->context)
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_getplatformids.c.inc

@@ -0,0 +1,46 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+/**
+ * \brief Get StarPU platform ID
+ */
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetPlatformIDs(cl_uint          num_entries,
+                 cl_platform_id * platforms,
+                 cl_uint *        num_platforms) CL_API_SUFFIX__VERSION_1_0
+{
+   if ((num_entries == 0 && platforms != NULL)
+      || (num_platforms == NULL && platforms == NULL))
+      return CL_INVALID_VALUE;
+
+   if (starpu_opencl_worker_get_count() == 0) {
+      DEBUG_ERROR("StarPU didn't find any OpenCL device. Try disabling CUDA support in StarPU (export STARPU_NCUDA=0).")
+
+      if (num_platforms != NULL)
+         *num_platforms = 0;
+   }
+   else {
+
+      if (platforms != NULL)
+         platforms[0] = &socl_platform;
+
+      if (num_platforms != NULL)
+         *num_platforms = 1;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_getplatforminfo.c.inc

@@ -0,0 +1,44 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+/**
+ * \brief Get information about StarPU platform
+ *
+ * \param[in] platform StarPU platform ID or NULL
+ */
+CL_API_ENTRY cl_int CL_API_CALL 
+soclGetPlatformInfo(cl_platform_id   platform, 
+                  cl_platform_info param_name,
+                  size_t           param_value_size, 
+                  void *           param_value,
+                  size_t *         param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (platform != NULL && platform != &socl_platform)
+      return CL_INVALID_PLATFORM;
+
+   switch (param_name) {
+      INFO_CASE(CL_PLATFORM_PROFILE, SOCL_PROFILE)
+      INFO_CASE(CL_PLATFORM_VERSION, SOCL_VERSION)
+      INFO_CASE(CL_PLATFORM_NAME,    SOCL_PLATFORM_NAME)
+      INFO_CASE(CL_PLATFORM_VENDOR,  SOCL_VENDOR)
+      INFO_CASE(CL_PLATFORM_EXTENSIONS, SOCL_PLATFORM_EXTENSIONS)
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_getprogrambuildinfo.c.inc

@@ -0,0 +1,40 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetProgramBuildInfo(cl_program          program,
+                      cl_device_id          UNUSED(device),
+                      cl_program_build_info param_name,
+                      size_t                param_value_size,
+                      void *                param_value,
+                      size_t *              param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (program == NULL)
+      return CL_INVALID_PROGRAM;
+
+
+   switch (param_name) {
+      //TODO
+      //INFO_CASE(CL_PROGRAM_BUILD_STATUS, program->build_status);
+      INFO_CASE_EX(CL_PROGRAM_BUILD_OPTIONS, program->options, program->options_size);
+      //TODO
+      //INFO_CASE(CL_PROGRAM_BUILD_LOG, program->build_log);
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_getprograminfo.c.inc

@@ -0,0 +1,41 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetProgramInfo(cl_program       program,
+                 cl_program_info    param_name,
+                 size_t             param_value_size,
+                 void *             param_value,
+                 size_t *           param_value_size_ret) CL_API_SUFFIX__VERSION_1_0
+{
+   if (program == NULL)
+      return CL_INVALID_PROGRAM;
+
+   switch (param_name) {
+      INFO_CASE(CL_PROGRAM_REFERENCE_COUNT, program->_entity.refs);
+      INFO_CASE(CL_PROGRAM_CONTEXT, program->context);
+      INFO_CASE(CL_PROGRAM_NUM_DEVICES, program->context->num_devices);
+      INFO_CASE_EX(CL_PROGRAM_DEVICES, program->context->devices, sizeof(cl_device_id)*program->context->num_devices);
+      //TODO
+      /*INFO_CASE(CL_PROGRAM_SOURCE, program->source);
+      INFO_CASE(CL_PROGRAM_BINARY_SIZE, program->binary_sizes);
+      INFO_CASE(CL_PROGRAM_BINARIES, program->binaries);*/
+      default:
+         return CL_INVALID_VALUE;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_getsamplerinfo.c.inc

@@ -0,0 +1,25 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetSamplerInfo(cl_sampler       UNUSED(sampler),
+                 cl_sampler_info    UNUSED(param_name),
+                 size_t             UNUSED(param_value_size),
+                 void *             UNUSED(param_value),
+                 size_t *           UNUSED(param_value_size_ret)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_getsupportedimageformats.c.inc

@@ -0,0 +1,26 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclGetSupportedImageFormats(cl_context           UNUSED(context),
+                           cl_mem_flags         UNUSED(flags),
+                           cl_mem_object_type   UNUSED(image_type),
+                           cl_uint              UNUSED(num_entries),
+                           cl_image_format *    UNUSED(image_formats),
+                           cl_uint *            UNUSED(num_image_formats)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_releasecommandqueue.c.inc

@@ -0,0 +1,43 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+static void release_callback_command_queue(void * e) {
+  cl_command_queue cq = (cl_command_queue)e;
+
+  //Disable StarPU profiling if necessary
+  if (cq->properties & CL_QUEUE_PROFILING_ENABLE) {
+    profiling_queue_count -= 1;
+    if (profiling_queue_count == 0)
+      starpu_profiling_status_set(STARPU_PROFILING_DISABLE);
+  }
+
+  /* Release references */
+  gc_entity_unstore(&cq->context);
+
+  /* Destruct object */
+  pthread_spin_destroy(&cq->spin);
+  free(cq->events);
+}
+
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseCommandQueue(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
+{
+
+  gc_entity_release(cq);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releasecontext.c.inc

@@ -0,0 +1,38 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+static void release_callback_context(void * e) {
+  cl_context context = (cl_context)e;
+
+  /* Destruct object */
+  if (context->properties != NULL)
+    free(context->properties);
+
+  free(context->devices);
+}
+
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseContext(cl_context context) CL_API_SUFFIX__VERSION_1_0
+{
+
+  if (context == NULL)
+    return CL_INVALID_CONTEXT;
+
+  gc_entity_release(context);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releaseevent.c.inc

@@ -0,0 +1,59 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+static void release_callback_event(void * e) {
+  cl_event event = (cl_event)e;
+
+  cl_command_queue cq = event->cq;
+
+  /* Remove from command queue */
+  if (cq != NULL) {
+    /* Lock command queue */
+    pthread_spin_lock(&cq->spin);
+
+    /* Remove barrier if applicable */
+    if (cq->barrier == event)
+      cq->barrier = NULL;
+
+    /* Remove from the list of out-of-order events */
+    if (event->prev != NULL)
+      event->prev->next = event->next;
+    if (event->next != NULL)
+      event->next->prev = event->prev;
+    if (cq->events == event)
+      cq->events = event->next;
+
+    /* Unlock command queue */
+    pthread_spin_unlock(&cq->spin);
+
+    gc_entity_unstore(&cq);
+  }
+
+  /* Destruct object */
+  //FIXME: we cannot release tag because it makes StarPU crash
+  //starpu_tag_remove(event->id);
+}
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseEvent(cl_event event) CL_API_SUFFIX__VERSION_1_0
+{
+  if (event == NULL)
+    return CL_INVALID_EVENT;
+
+  gc_entity_release(event);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releasekernel.c.inc

@@ -0,0 +1,70 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+static void rk_task(void *data) {
+   cl_kernel k = (cl_kernel)data;
+
+   int range = starpu_worker_get_range();
+
+   cl_int err = clReleaseKernel(k->cl_kernels[range]);
+   if (err != CL_SUCCESS)
+      DEBUG_CL("clReleaseKernel", err);
+}
+
+static void release_callback_kernel(void * e) {
+  cl_kernel kernel = (cl_kernel)e;
+
+  //Free args
+  unsigned int j;
+  for (j=0; j<kernel->arg_count; j++) {
+    switch (kernel->arg_type[j]) {
+      case Null:
+        break;
+      case Buffer:
+        gc_entity_unstore((cl_mem*)&kernel->arg_value[j]);
+        break;
+      case Immediate:
+        free(kernel->arg_value[j]);
+        break;
+    }
+  }
+  if (kernel->arg_size != NULL)
+    free(kernel->arg_size);
+  if (kernel->arg_value != NULL)
+    free(kernel->arg_value);
+  if (kernel->arg_type != NULL)
+    free(kernel->arg_type);
+
+  //Release real kernels...
+  starpu_execute_on_each_worker(rk_task, kernel, STARPU_OPENCL);
+
+  gc_entity_unstore(&kernel->program);
+
+  free(kernel->kernel_name);
+  free(kernel->cl_kernels);
+  free(kernel->errcodes);
+}
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseKernel(cl_kernel kernel) CL_API_SUFFIX__VERSION_1_0
+{
+  if (kernel == NULL)
+    return CL_INVALID_KERNEL;
+
+  gc_entity_release(kernel);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releasememobject.c.inc

@@ -0,0 +1,39 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+static void release_callback_memobject(void * e) {
+  cl_mem mem = (cl_mem)e;
+
+  /* Release references */
+  gc_entity_unstore(&mem->context);
+
+  //Delete this mem_object from the mem_object list
+  mem_object_release(mem);
+
+  /* Destruct object */
+  starpu_data_unregister_no_coherency(mem->handle);
+
+  if (!(mem->flags & CL_MEM_USE_HOST_PTR))
+    free(mem->ptr);
+}
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseMemObject(cl_mem mem) CL_API_SUFFIX__VERSION_1_0
+{
+  gc_entity_release(mem);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releaseprogram.c.inc

@@ -0,0 +1,51 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+static void rp_task(void *data) {
+   struct _cl_program *d = (struct _cl_program*)data;
+
+   int range = starpu_worker_get_range();
+
+   cl_int err = clReleaseProgram(d->cl_programs[range]);
+   if (err != CL_SUCCESS)
+      DEBUG_CL("clReleaseProgram", err);
+}
+
+static void release_callback_program(void * e) {
+  cl_program program = (cl_program)e;
+
+  /* Destruct object */
+  starpu_execute_on_each_worker(rp_task, program, STARPU_OPENCL);
+
+  /* Release references */
+  gc_entity_unstore(&program->context);
+
+  free(program->cl_programs);
+
+  if (program->options != NULL)
+    free(program->options);
+}
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseProgram(cl_program program) CL_API_SUFFIX__VERSION_1_0
+{
+  if (program == NULL)
+    return CL_INVALID_PROGRAM;
+
+  gc_entity_release(program);
+
+  return CL_SUCCESS;
+}

socl/src/cl_releasesampler.c.inc

@@ -0,0 +1,21 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclReleaseSampler(cl_sampler UNUSED(sampler)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_retaincommandqueue.c.inc

@@ -0,0 +1,26 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclRetainCommandQueue(cl_command_queue cq) CL_API_SUFFIX__VERSION_1_0
+{
+   if (cq == NULL)
+      return CL_INVALID_COMMAND_QUEUE;
+
+   gc_entity_retain(cq);
+
+   return CL_SUCCESS;
+}

socl/src/cl_retaincontext.c.inc

@@ -0,0 +1,26 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclRetainContext(cl_context context) CL_API_SUFFIX__VERSION_1_0
+{
+   if (context == NULL)
+      return CL_INVALID_CONTEXT;
+
+   gc_entity_retain(context);
+
+   return CL_SUCCESS;
+}

socl/src/cl_retainevent.c.inc

@@ -0,0 +1,26 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclRetainEvent(cl_event event) CL_API_SUFFIX__VERSION_1_0
+{
+   if (event == NULL)
+      return CL_INVALID_EVENT;
+
+   gc_entity_retain(event);
+
+   return CL_SUCCESS;
+}

socl/src/cl_retainkernel.c.inc

@@ -0,0 +1,26 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclRetainKernel(cl_kernel    kernel) CL_API_SUFFIX__VERSION_1_0
+{
+   if (kernel == NULL)
+      return CL_INVALID_KERNEL;
+
+   gc_entity_retain(kernel);
+
+   return CL_SUCCESS;
+}

socl/src/cl_retainmemobject.c.inc

@@ -0,0 +1,26 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclRetainMemObject(cl_mem mem) CL_API_SUFFIX__VERSION_1_0
+{
+   if (mem == NULL)
+      return CL_INVALID_MEM_OBJECT;
+
+   gc_entity_retain(mem);
+   
+   return CL_SUCCESS;
+}

socl/src/cl_retainprogram.c.inc

@@ -0,0 +1,26 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclRetainProgram(cl_program program) CL_API_SUFFIX__VERSION_1_0
+{
+   if (program == NULL)
+      return CL_INVALID_PROGRAM;
+
+   gc_entity_retain(program);
+
+   return CL_SUCCESS;
+}

socl/src/cl_retainsampler.c.inc

@@ -0,0 +1,21 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclRetainSampler(cl_sampler UNUSED(sampler)) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_INVALID_OPERATION;
+}

socl/src/cl_setcommandqueueproperty.c.inc

@@ -0,0 +1,50 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclSetCommandQueueProperty(cl_command_queue            command_queue,
+                          cl_command_queue_properties   properties, 
+                          cl_bool                       enable,
+                          cl_command_queue_properties * old_properties) CL_API_SUFFIX__VERSION_1_0
+{
+   if (command_queue == NULL)
+      return CL_INVALID_COMMAND_QUEUE;
+
+   if (old_properties != NULL)
+      *old_properties = command_queue->properties;
+
+   if (enable) {
+      //Enable StarPU profiling if necessary
+      if (properties & (~command_queue->properties) & CL_QUEUE_PROFILING_ENABLE) {
+         if (profiling_queue_count == 0)
+            starpu_profiling_status_set(STARPU_PROFILING_ENABLE);
+         profiling_queue_count += 1;
+      }  
+      //Set new properties
+      command_queue->properties |= properties;
+   } else {
+      //Disable StarPU profiling if necessary
+      if ((~properties) & command_queue->properties & CL_QUEUE_PROFILING_ENABLE) {
+         profiling_queue_count -= 1;
+         if (profiling_queue_count == 0)
+            starpu_profiling_status_set(STARPU_PROFILING_DISABLE);
+      }  
+      //Set new properties
+      command_queue->properties &= ~properties;
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_setkernelarg.c.inc

@@ -0,0 +1,73 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclSetKernelArg(cl_kernel  kernel,
+               cl_uint      arg_index,
+               size_t       arg_size,
+               const void * arg_value) CL_API_SUFFIX__VERSION_1_0
+{
+   if (kernel == NULL)
+      return CL_INVALID_KERNEL;
+
+   if (arg_index >= kernel->arg_count)
+      return CL_INVALID_ARG_INDEX;
+
+   //FIXME: we don't return CL_INVALID_ARG_VALUE if "arg_value is NULL for an argument that is not declared with __local qualifier or vice-versa"
+   //FIXME: we don't return CL_INVALID_MEM_OBJECT
+   //FIXME: we don't return CL_INVALID_ARG_SIZE
+
+   /* Free previous argument (set to NULL) */
+   switch (kernel->arg_type[arg_index]) {
+      case Null:
+         break;
+      case Buffer:
+         kernel->arg_type[arg_index] = Null;
+         gc_entity_unstore((cl_mem*)&kernel->arg_value[arg_index]);
+         kernel->arg_value[arg_index] = NULL;
+         break;
+      case Immediate:
+         free(kernel->arg_value[arg_index]);
+         kernel->arg_type[arg_index] = Null;
+         kernel->arg_value[arg_index] = NULL;
+         break;
+   }
+
+   kernel->arg_type[arg_index] = Null;
+   kernel->arg_size[arg_index] = arg_size;
+
+   DEBUG_MSG("[Kernel %d] Set argument %d: argsize %ld argvalue %p\n", kernel->id, arg_index, arg_size, arg_value);
+
+   /* Argument is not Null */
+   if (arg_value != NULL) {
+      cl_mem buf = NULL;
+      /* Check if argument is a memory object */
+      if ((arg_size == sizeof(cl_mem)) && ((buf = mem_object_fetch(arg_value)) != NULL)) {
+         DEBUG_MSG("Found buffer %d \n", buf->id);
+         kernel->arg_type[arg_index] = Buffer;
+         gc_entity_store(&kernel->arg_value[arg_index], buf);
+      }
+      else {
+         /* Argument must be an immediate buffer  */
+         DEBUG_MSG("Immediate data\n");
+         kernel->arg_type[arg_index] = Immediate;
+         kernel->arg_value[arg_index] = malloc(arg_size);
+         memcpy(kernel->arg_value[arg_index], arg_value, arg_size);
+      }
+   }
+
+   return CL_SUCCESS;
+}

socl/src/cl_unloadcompiler.c.inc

@@ -0,0 +1,21 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclUnloadCompiler(void) CL_API_SUFFIX__VERSION_1_0
+{
+   return CL_SUCCESS;
+}

socl/src/cl_waitforevents.c.inc

@@ -0,0 +1,33 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+CL_API_ENTRY cl_int CL_API_CALL
+soclWaitForEvents(cl_uint           num_events,
+                const cl_event *    event_list) CL_API_SUFFIX__VERSION_1_0
+{
+   unsigned int i;
+   DEBUG_MSG("Waiting for events: ");
+   for (i=0; i<num_events; i++) {
+      DEBUG_MSG_NOHEAD("%d ", event_list[i]->id);
+   }
+   DEBUG_MSG_NOHEAD("\n");
+
+   for (i=0; i<num_events; i++)
+      starpu_tag_wait(event_list[i]->id);
+
+   DEBUG_MSG("Stop waiting :)\n");
+   return CL_SUCCESS;
+}

socl/src/device_descriptions.c.inc

@@ -0,0 +1,321 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+// OpenCL 1.0 : Mandatory format: major_number.minor_number
+static const char SOCL_DRIVER_VERSION[] = "0.1";
+
+static const cl_uint SOCL_DEVICE_VENDOR_ID = 666;
+
+struct _cl_device_id {
+   cl_device_type    type;
+   cl_uint           max_compute_units;        //OpenCL 1.0: minimum value is 1
+   cl_uint           max_work_item_dimensions; //OpenCL 1.0: minimum value is 3
+   size_t            max_work_item_sizes[3];   //array size should be set accordingly to the maximum max_work_item_dimensions
+   size_t            max_work_group_size;      //OpenCL 1.0: minimum value is 1
+   cl_uint           preferred_vector_widths[6]; //Char, Short, Int, Long, Float, Double
+   cl_uint           max_clock_frequency;
+   cl_uint           address_bits;             //OpenCL 1.0: 32 or 64
+   cl_ulong          max_mem_alloc_size;       //OpenCL 1.0: minimum value is max(CL_DEVICE_GLOBAL_MEM_SIZE/4, 128*1024*1024)
+   cl_bool           image_support;
+   //image fields not present
+   size_t            max_parameter_size;       //OpenCL 1.0: minimum is 256
+   cl_uint           mem_base_addr_align;
+   cl_uint           min_data_type_align_size;
+   cl_device_fp_config single_fp_config;       //OpenCL 1.0: CL_FP_ROUND_TO_NEAREST and CL_FP_INF_NAN are mandatory
+   cl_device_mem_cache_type global_mem_cache_type;
+   cl_uint           global_mem_cacheline_size;
+   cl_ulong          global_mem_cache_size;
+   cl_ulong          global_mem_size;
+   cl_ulong          max_constant_buffer_size; //OpenCL 1.0: minimum value is 64KB
+   cl_uint           max_constant_args;        //OpenCL 1.0: minimum value is 8
+   cl_device_local_mem_type local_mem_type;
+   cl_ulong          local_mem_size;           //OpenCL 1.0: minimum value is 16KB
+   cl_bool           error_correction_support;
+   size_t            profiling_timer_resolution;
+   cl_bool           endian_little;
+   cl_bool           available;
+   cl_bool           compiler_available;       //OpenCL 1.0: mandatory for FULL_PROFILE platforms
+   cl_device_exec_capabilities execution_capabilities; //OpenCL 1.0: CL_EXEC_KERNEL is mandatory
+   cl_command_queue_properties queue_properties; //OpenCL 1.0: CL_QUEUE_PROFILING_ENABLE is mandatory
+   char              name[40];                 //Array size has been arbitrarily defined
+   //versions, profile and vendor are statically defined for all devices
+   char              extensions[100];          //Array size has been arbitrarily defined
+};
+
+static const struct _cl_device_id socl_devices[] = {
+   { 
+      .type = CL_DEVICE_TYPE_CPU,
+      .max_compute_units = 1,
+      .max_work_item_dimensions = 3,
+      .max_work_item_sizes = {1,1,1},
+      .max_work_group_size = 1,
+      .preferred_vector_widths = {16,8,4,2,4,2},
+      .max_clock_frequency = 3000,
+      .address_bits = 64,
+      .max_mem_alloc_size = 1024*1024*1024,
+      .image_support = CL_FALSE,
+      .max_parameter_size = 256,
+      .mem_base_addr_align = 0,
+      .min_data_type_align_size = 0,
+      .single_fp_config = CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN,
+      .global_mem_cache_type = CL_READ_WRITE_CACHE,
+      .global_mem_cacheline_size = 128,
+      .global_mem_cache_size = 16*1024,
+      .global_mem_size = (cl_ulong)4*1024*1024*1024,
+      .max_constant_args = 8,
+      .local_mem_type = CL_GLOBAL,
+      .local_mem_size = 16*1024,
+      .error_correction_support = CL_FALSE,
+      .profiling_timer_resolution = 100,
+      .endian_little = CL_TRUE,
+      .available = CL_TRUE,
+      .compiler_available = CL_TRUE,
+      .execution_capabilities = CL_EXEC_KERNEL,
+      .queue_properties = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE,
+      .name = "StarPU virtual CPU 1",
+      .extensions = ""
+   },
+   {
+      .type = CL_DEVICE_TYPE_GPU,
+      .max_compute_units = 12,
+      .max_work_item_dimensions = 3,
+      .max_work_item_sizes = {512,512,64},
+      .max_work_group_size = 512,
+      .preferred_vector_widths = {16,8,4,2,4,2},
+      .max_clock_frequency = 1600,
+      .address_bits = 32,
+      .max_mem_alloc_size = 100*1024*1024,
+      .image_support = CL_FALSE,
+      .max_parameter_size = 256,
+      .mem_base_addr_align = 0,
+      .min_data_type_align_size = 0,
+      .single_fp_config = CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN,
+      .global_mem_cache_type = CL_NONE,
+      .global_mem_cacheline_size = 0,
+      .global_mem_cache_size = 0,
+      .global_mem_size = (cl_ulong)500*1024*1024,
+      .max_constant_args = 8,
+      .local_mem_type = CL_LOCAL,
+      .local_mem_size = 16*1024,
+      .error_correction_support = CL_FALSE,
+      .profiling_timer_resolution = 10,
+      .endian_little = CL_TRUE,
+      .available = CL_TRUE,
+      .compiler_available = CL_TRUE,
+      .execution_capabilities = CL_EXEC_KERNEL,
+      .queue_properties = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE,
+      .name = "StarPU virtual GPU 1",
+      .extensions = ""
+   },
+   {
+      .type = CL_DEVICE_TYPE_GPU,
+      .max_compute_units = 12,
+      .max_work_item_dimensions = 3,
+      .max_work_item_sizes = {512,512,64},
+      .max_work_group_size = 512,
+      .preferred_vector_widths = {16,8,4,2,4,2},
+      .max_clock_frequency = 1600,
+      .address_bits = 32,
+      .max_mem_alloc_size = 100*1024*1024,
+      .image_support = CL_FALSE,
+      .max_parameter_size = 256,
+      .mem_base_addr_align = 0,
+      .min_data_type_align_size = 0,
+      .single_fp_config = CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN,
+      .global_mem_cache_type = CL_NONE,
+      .global_mem_cacheline_size = 0,
+      .global_mem_cache_size = 0,
+      .global_mem_size = (cl_ulong)500*1024*1024,
+      .max_constant_args = 8,
+      .local_mem_type = CL_LOCAL,
+      .local_mem_size = 16*1024,
+      .error_correction_support = CL_FALSE,
+      .profiling_timer_resolution = 10,
+      .endian_little = CL_TRUE,
+      .available = CL_TRUE,
+      .compiler_available = CL_TRUE,
+      .execution_capabilities = CL_EXEC_KERNEL,
+      .queue_properties = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE,
+      .name = "StarPU virtual GPU 2",
+      .extensions = ""
+   },
+   {
+      .type = CL_DEVICE_TYPE_GPU,
+      .max_compute_units = 12,
+      .max_work_item_dimensions = 3,
+      .max_work_item_sizes = {512,512,64},
+      .max_work_group_size = 512,
+      .preferred_vector_widths = {16,8,4,2,4,2},
+      .max_clock_frequency = 1600,
+      .address_bits = 32,
+      .max_mem_alloc_size = 100*1024*1024,
+      .image_support = CL_FALSE,
+      .max_parameter_size = 256,
+      .mem_base_addr_align = 0,
+      .min_data_type_align_size = 0,
+      .single_fp_config = CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN,
+      .global_mem_cache_type = CL_NONE,
+      .global_mem_cacheline_size = 0,
+      .global_mem_cache_size = 0,
+      .global_mem_size = (cl_ulong)500*1024*1024,
+      .max_constant_args = 8,
+      .local_mem_type = CL_LOCAL,
+      .local_mem_size = 16*1024,
+      .error_correction_support = CL_FALSE,
+      .profiling_timer_resolution = 10,
+      .endian_little = CL_TRUE,
+      .available = CL_TRUE,
+      .compiler_available = CL_TRUE,
+      .execution_capabilities = CL_EXEC_KERNEL,
+      .queue_properties = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE,
+      .name = "StarPU virtual GPU 3",
+      .extensions = ""
+   },
+   {
+      .type = CL_DEVICE_TYPE_GPU,
+      .max_compute_units = 12,
+      .max_work_item_dimensions = 3,
+      .max_work_item_sizes = {512,512,64},
+      .max_work_group_size = 512,
+      .preferred_vector_widths = {16,8,4,2,4,2},
+      .max_clock_frequency = 1600,
+      .address_bits = 32,
+      .max_mem_alloc_size = 100*1024*1024,
+      .image_support = CL_FALSE,
+      .max_parameter_size = 256,
+      .mem_base_addr_align = 0,
+      .min_data_type_align_size = 0,
+      .single_fp_config = CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN,
+      .global_mem_cache_type = CL_NONE,
+      .global_mem_cacheline_size = 0,
+      .global_mem_cache_size = 0,
+      .global_mem_size = (cl_ulong)500*1024*1024,
+      .max_constant_args = 8,
+      .local_mem_type = CL_LOCAL,
+      .local_mem_size = 16*1024,
+      .error_correction_support = CL_FALSE,
+      .profiling_timer_resolution = 10,
+      .endian_little = CL_TRUE,
+      .available = CL_TRUE,
+      .compiler_available = CL_TRUE,
+      .execution_capabilities = CL_EXEC_KERNEL,
+      .queue_properties = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE,
+      .name = "StarPU virtual GPU 4",
+      .extensions = ""
+   },
+   {
+      .type = CL_DEVICE_TYPE_GPU,
+      .max_compute_units = 12,
+      .max_work_item_dimensions = 3,
+      .max_work_item_sizes = {512,512,64},
+      .max_work_group_size = 512,
+      .preferred_vector_widths = {16,8,4,2,4,2},
+      .max_clock_frequency = 1600,
+      .address_bits = 32,
+      .max_mem_alloc_size = 100*1024*1024,
+      .image_support = CL_FALSE,
+      .max_parameter_size = 256,
+      .mem_base_addr_align = 0,
+      .min_data_type_align_size = 0,
+      .single_fp_config = CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN,
+      .global_mem_cache_type = CL_NONE,
+      .global_mem_cacheline_size = 0,
+      .global_mem_cache_size = 0,
+      .global_mem_size = (cl_ulong)500*1024*1024,
+      .max_constant_args = 8,
+      .local_mem_type = CL_LOCAL,
+      .local_mem_size = 16*1024,
+      .error_correction_support = CL_FALSE,
+      .profiling_timer_resolution = 10,
+      .endian_little = CL_TRUE,
+      .available = CL_TRUE,
+      .compiler_available = CL_TRUE,
+      .execution_capabilities = CL_EXEC_KERNEL,
+      .queue_properties = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE,
+      .name = "StarPU virtual GPU 5",
+      .extensions = ""
+   },
+   {
+      .type = CL_DEVICE_TYPE_GPU,
+      .max_compute_units = 12,
+      .max_work_item_dimensions = 3,
+      .max_work_item_sizes = {512,512,64},
+      .max_work_group_size = 512,
+      .preferred_vector_widths = {16,8,4,2,4,2},
+      .max_clock_frequency = 1600,
+      .address_bits = 32,
+      .max_mem_alloc_size = 100*1024*1024,
+      .image_support = CL_FALSE,
+      .max_parameter_size = 256,
+      .mem_base_addr_align = 0,
+      .min_data_type_align_size = 0,
+      .single_fp_config = CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN,
+      .global_mem_cache_type = CL_NONE,
+      .global_mem_cacheline_size = 0,
+      .global_mem_cache_size = 0,
+      .global_mem_size = (cl_ulong)500*1024*1024,
+      .max_constant_args = 8,
+      .local_mem_type = CL_LOCAL,
+      .local_mem_size = 16*1024,
+      .error_correction_support = CL_FALSE,
+      .profiling_timer_resolution = 10,
+      .endian_little = CL_TRUE,
+      .available = CL_TRUE,
+      .compiler_available = CL_TRUE,
+      .execution_capabilities = CL_EXEC_KERNEL,
+      .queue_properties = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE,
+      .name = "StarPU virtual GPU 6",
+      .extensions = ""
+   },
+   {
+      .type = CL_DEVICE_TYPE_GPU,
+      .max_compute_units = 12,
+      .max_work_item_dimensions = 3,
+      .max_work_item_sizes = {512,512,64},
+      .max_work_group_size = 512,
+      .preferred_vector_widths = {16,8,4,2,4,2},
+      .max_clock_frequency = 1600,
+      .address_bits = 32,
+      .max_mem_alloc_size = 100*1024*1024,
+      .image_support = CL_FALSE,
+      .max_parameter_size = 256,
+      .mem_base_addr_align = 0,
+      .min_data_type_align_size = 0,
+      .single_fp_config = CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN,
+      .global_mem_cache_type = CL_NONE,
+      .global_mem_cacheline_size = 0,
+      .global_mem_cache_size = 0,
+      .global_mem_size = (cl_ulong)500*1024*1024,
+      .max_constant_args = 8,
+      .local_mem_type = CL_LOCAL,
+      .local_mem_size = 16*1024,
+      .error_correction_support = CL_FALSE,
+      .profiling_timer_resolution = 10,
+      .endian_little = CL_TRUE,
+      .available = CL_TRUE,
+      .compiler_available = CL_TRUE,
+      .execution_capabilities = CL_EXEC_KERNEL,
+      .queue_properties = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE,
+      .name = "StarPU virtual GPU 7",
+      .extensions = ""
+   }
+
+};
+
+static const int socl_device_count = sizeof(socl_devices) / sizeof(struct _cl_device_id);
+
+
+

socl/src/gc.c.inc

@@ -0,0 +1,187 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+/**
+ * Garbage collection thread
+ */
+
+/* List of entities to be released */
+static volatile entity gc_list = NULL;
+static volatile entity entities = NULL;
+
+/* Mutex and cond for release */
+static pthread_mutex_t gc_mutex = PTHREAD_MUTEX_INITIALIZER;
+static pthread_cond_t  gc_cond = PTHREAD_COND_INITIALIZER;
+
+/* Set to 1 to stop release thread execution */
+static volatile int gc_stop_required = 0;
+
+#define GC_LOCK pthread_mutex_lock(&gc_mutex)
+#define GC_UNLOCK { pthread_cond_signal(&gc_cond); \
+                    pthread_mutex_unlock(&gc_mutex);}
+#define GC_UNLOCK_NO_SIGNAL pthread_mutex_unlock(&gc_mutex)
+
+/* Thread routine */
+static void * gc_thread_routine(void *UNUSED(arg)) {
+  
+  GC_LOCK;
+
+  do {
+    /* Make a copy of the gc_list to allow callbacks to add things into it */
+    entity rs = gc_list;
+    gc_list = NULL;
+
+    GC_UNLOCK_NO_SIGNAL;
+
+    entity r = rs;
+    while (r != NULL) {
+      /* Call entity release callback */
+      if (r->release_callback != NULL) {
+        r->release_callback(r);
+      }
+
+      /* Release entity */
+      entity next = r->next;
+      free(r);
+
+      r = next;
+    }
+
+    GC_LOCK;
+
+    /* Check if new entities have been added */
+    if (gc_list != NULL)
+      continue;
+
+    /* Stop if required */
+    if (gc_stop_required) {
+      GC_UNLOCK_NO_SIGNAL;
+      break;
+    }
+
+    /* Otherwise we sleep */
+    pthread_cond_wait(&gc_cond, &gc_mutex);
+
+  } while (1);
+
+  pthread_exit(NULL);
+}
+
+static pthread_t gc_thread;
+
+/* Start garbage collection */
+static void gc_start(void) {
+  pthread_create(&gc_thread, NULL, gc_thread_routine, NULL);
+}
+
+/* Stop garbage collection */
+static void gc_stop(void) {
+  GC_LOCK;
+
+  gc_stop_required = 1;
+
+  GC_UNLOCK;
+
+  pthread_join(gc_thread, NULL);
+}
+
+/**
+ * Decrement reference counter and release entity if applicable
+ */
+#define gc_entity_release(a) gc_entity_release_ex(&(a)->_entity)
+
+static void gc_entity_release_ex(entity e) {
+
+  /* Decrement reference count */
+  int refs = __sync_sub_and_fetch(&e->refs, 1);
+
+  if (refs != 0)
+    return;
+
+  GC_LOCK;
+
+  /* Remove entity from the entities list */
+  if (e->prev != NULL)
+    e->prev->next = e->next;
+  if (e->next != NULL)
+    e->next->prev = e->prev;
+  if (entities == e)
+    entities = e->next;
+
+  /* Put entity in the release queue */
+  e->next = gc_list;
+  gc_list = e;
+
+  GC_UNLOCK;
+}
+
+
+
+/**
+ * Initialize entity
+ */
+static void gc_entity_init(void *arg, void (*release_callback)(void*)) {
+  struct entity * e = (entity)arg;
+
+  e->refs = 1;
+  e->release_callback = release_callback;
+  e->prev = NULL;
+
+  GC_LOCK;
+
+  e->next = entities;
+  if (entities != NULL)
+    entities->prev = e;
+  entities = e;
+
+  GC_UNLOCK_NO_SIGNAL;
+}
+
+/**
+ * Allocate and initialize entity
+ */
+static void * gc_entity_alloc(unsigned int size, void (*release_callback)(void*)) {
+  void * e = malloc(size);
+  gc_entity_init(e, release_callback);
+  return e;
+}
+
+/**
+ * Retain entity
+ */
+static void gc_entity_retain(void *arg) {
+
+  struct entity * e = (entity)arg;
+
+  __sync_fetch_and_add(&e->refs, 1);
+}
+
+
+#define gc_entity_store(dest,e) \
+  do {\
+    gc_entity_retain(e); \
+    *dest = e;\
+  } while(0);
+
+#define gc_entity_unstore(dest) \
+  do {\
+    gc_entity_release(*dest); \
+    *dest = NULL;\
+  } while(0);
+
+#undef GC_LOCK
+#undef GC_UNLOCK
+#undef GC_UNLOCK_NO_SIGNAL

socl/src/helper_command_queue.c.inc

@@ -0,0 +1,94 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+/**
+ * WARNING: command queues do NOT hold references on events. Only events hold references
+ * on command queues. This way, event release will automatically remove the event from
+ * its command queue.
+ */
+
+/**
+ * Enqueue the given task but put fake_event into the command queue.
+ * This is used when a tag notified by application is used (cf clEnqueueMapBuffer, etc.)
+ */
+static cl_int command_queue_enqueue_fakeevent(cl_command_queue cq, starpu_task *task, cl_int barrier, cl_int num_events, const cl_event * events, cl_event fake_event) {
+
+  int in_order = !(cq->properties & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE);
+
+  /* Set explicit task dependencies */
+  task_dependency_add(task, num_events, events);
+
+  /* Lock command queue */
+  pthread_spin_lock(&cq->spin);
+
+  /* Add dependency to last barrier if applicable */
+  if (cq->barrier != NULL)
+    task_dependency_add(task, 1, &cq->barrier);
+
+  /* Add dependencies to out-of-order events (if any) */
+  if (barrier) {
+    while (cq->events != NULL) {
+      task_dependency_add(task, 1, &cq->events);
+      cq->events = cq->events->next;
+    }
+  }
+
+  cl_event ev = (fake_event == NULL ? task_event(task) : fake_event);
+
+  /* Add event to the list of out-of-order events */
+  if (!in_order) {
+    ev->next = cq->events;
+    ev->prev = NULL;
+    if (cq->events != NULL)
+      cq->events->prev = ev;
+    cq->events = ev;
+  }
+
+  /* Register this event as last barrier */
+  if (barrier || in_order)
+    cq->barrier = ev;
+
+   /* Unlock command queue */
+   pthread_spin_unlock(&cq->spin);
+
+   /* Add reference to the command queue */
+   gc_entity_store(&ev->cq, cq);
+
+   /* Submit task */
+   gc_entity_retain(task_event(task));
+   int ret = starpu_task_submit(task);
+   if (ret != 0)
+      DEBUG_ERROR("Unable to submit a task. Error %d\n", ret);
+
+   return CL_SUCCESS;
+}
+
+static cl_int command_queue_enqueue(cl_command_queue cq, starpu_task *task, cl_int barrier, cl_int num_events, const cl_event * events) {
+  return command_queue_enqueue_fakeevent(cq, task, barrier, num_events, events, NULL);
+}
+
+
+static cl_event enqueueBarrier(cl_command_queue cq) {
+
+   //CL_COMMAND_MARKER has been chosen as CL_COMMAND_BARRIER doesn't exist
+   starpu_task * task = task_create(CL_COMMAND_MARKER);
+
+   DEBUG_MSG("Submitting barrier task (event %d)\n", task->tag_id);
+   command_queue_enqueue(cq, task, 1, 0, NULL);
+
+   return task_event(task);
+}

socl/src/helper_debug.c.inc

@@ -0,0 +1,93 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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>
+
+#ifdef DEBUG
+   #define DEBUG_MSG(...) do { fprintf(stderr, "[SOCL] [%s] ", __func__); fprintf(stderr, __VA_ARGS__); } while (0);
+   #define DEBUG_MSG_NOHEAD(...) fprintf(stderr, __VA_ARGS__)
+   #define DEBUG_ERROR(...) do { fprintf(stderr, "[SOCL] ERROR: "__VA_ARGS__); exit(1); } while (0);
+#else
+   #define DEBUG_MSG(...) while(0);
+   #define DEBUG_MSG_NOHEAD(...) while(0);
+   #define DEBUG_ERROR(...) while(0);
+#endif
+
+
+#define ERROR_MSG(...) do { fprintf(stderr, "[SOCL] [%s] ERROR: ", __func__); fprintf(stderr, __VA_ARGS__); } while (0);
+#define ERROR_MSG_NOHEAD(...) fprintf(stderr, __VA_ARGS__)
+#define ERROR_STOP(...) do { ERROR_MSG(__VA_ARGS__); exit(1); } while(0);
+
+#ifdef DEBUG
+void DEBUG_CL(char *s, cl_int err) {
+   #define ERR_CASE(a) case a: DEBUG_MSG("[OpenCL] %s CL error: %s\n", s, #a); break;
+   switch(err) {
+      case CL_SUCCESS:
+         DEBUG_MSG("[OpenCL] %s SUCCESS.\n", s);
+         break;
+      ERR_CASE(CL_DEVICE_NOT_FOUND)
+      ERR_CASE(CL_DEVICE_NOT_AVAILABLE)
+      ERR_CASE(CL_COMPILER_NOT_AVAILABLE)
+      ERR_CASE(CL_MEM_OBJECT_ALLOCATION_FAILURE)
+      ERR_CASE(CL_OUT_OF_RESOURCES)
+      ERR_CASE(CL_OUT_OF_HOST_MEMORY)
+      ERR_CASE(CL_PROFILING_INFO_NOT_AVAILABLE)
+      ERR_CASE(CL_MEM_COPY_OVERLAP)
+      ERR_CASE(CL_IMAGE_FORMAT_MISMATCH)
+      ERR_CASE(CL_IMAGE_FORMAT_NOT_SUPPORTED)
+      ERR_CASE(CL_BUILD_PROGRAM_FAILURE)
+      ERR_CASE(CL_MAP_FAILURE)
+      ERR_CASE(CL_INVALID_VALUE)
+      ERR_CASE(CL_INVALID_DEVICE_TYPE)
+      ERR_CASE(CL_INVALID_PLATFORM)
+      ERR_CASE(CL_INVALID_DEVICE)
+      ERR_CASE(CL_INVALID_CONTEXT)
+      ERR_CASE(CL_INVALID_QUEUE_PROPERTIES)
+      ERR_CASE(CL_INVALID_COMMAND_QUEUE)
+      ERR_CASE(CL_INVALID_HOST_PTR)
+      ERR_CASE(CL_INVALID_MEM_OBJECT)
+      ERR_CASE(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR)
+      ERR_CASE(CL_INVALID_IMAGE_SIZE)
+      ERR_CASE(CL_INVALID_SAMPLER)
+      ERR_CASE(CL_INVALID_BINARY)
+      ERR_CASE(CL_INVALID_BUILD_OPTIONS)
+      ERR_CASE(CL_INVALID_PROGRAM)
+      ERR_CASE(CL_INVALID_PROGRAM_EXECUTABLE)
+      ERR_CASE(CL_INVALID_KERNEL_NAME)
+      ERR_CASE(CL_INVALID_KERNEL_DEFINITION)
+      ERR_CASE(CL_INVALID_KERNEL)
+      ERR_CASE(CL_INVALID_ARG_INDEX)
+      ERR_CASE(CL_INVALID_ARG_VALUE)
+      ERR_CASE(CL_INVALID_ARG_SIZE)
+      ERR_CASE(CL_INVALID_KERNEL_ARGS)
+      ERR_CASE(CL_INVALID_WORK_DIMENSION)
+      ERR_CASE(CL_INVALID_WORK_GROUP_SIZE)
+      ERR_CASE(CL_INVALID_WORK_ITEM_SIZE)
+      ERR_CASE(CL_INVALID_GLOBAL_OFFSET)
+      ERR_CASE(CL_INVALID_EVENT_WAIT_LIST)
+      ERR_CASE(CL_INVALID_EVENT)
+      ERR_CASE(CL_INVALID_OPERATION)
+      ERR_CASE(CL_INVALID_GL_OBJECT)
+      ERR_CASE(CL_INVALID_BUFFER_SIZE)
+      ERR_CASE(CL_INVALID_MIP_LEVEL)
+      ERR_CASE(CL_INVALID_GLOBAL_WORK_SIZE)
+      default:
+         DEBUG_MSG("%s CL error: Error message not supported by print_cl_error (%d).\n", s, err);
+   }
+}
+#else
+   #define DEBUG_CL(...) while(0);
+#endif

socl/src/helper_event.c.inc

@@ -0,0 +1,37 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+/**
+ * Create a new event
+ *
+ * Events have one-to-one relation with tag. Tag number is event ID
+ */
+static cl_event event_create(void) {
+   static int id = 1;
+   cl_event ev;
+   ev = gc_entity_alloc(sizeof(struct _cl_event), release_callback_event);
+
+   ev->next = NULL;
+   ev->prev = NULL;
+   ev->id = __sync_fetch_and_add(&id,1);
+   ev->status = CL_SUBMITTED;
+   ev->type = 0;
+   ev->profiling_info = NULL;
+   ev->cq = NULL;
+
+   return ev;
+}
+

socl/src/helper_getinfo.c.inc

@@ -0,0 +1,44 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+
+#define INFO_CASE_EX2(var) if (param_value != NULL) { \
+      if (param_value_size < sizeof(var)) \
+         return CL_INVALID_VALUE; \
+      memcpy(param_value, &var, sizeof(var)); \
+   } \
+   if (param_value_size_ret != NULL) \
+      *param_value_size_ret = sizeof(var); \
+   break;
+
+#define INFO_CASE(param, var) case param: \
+   INFO_CASE_EX2(var)
+
+#define INFO_CASE_VALUE(param, type, value) case param: {\
+      type tmp = (value);\
+      INFO_CASE_EX2(tmp);\
+   }
+
+//warning: var is a reference
+#define INFO_CASE_EX(param, var, size) case param: \
+   if (param_value != NULL) { \
+      if (param_value_size < size) \
+         return CL_INVALID_VALUE; \
+      memcpy(param_value, var, size); \
+   } \
+   if (param_value_size_ret != NULL) \
+      *param_value_size_ret = size; \
+   break;

socl/src/helper_mem_objects.c.inc

@@ -0,0 +1,90 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+#define mem_object_hash_key 257
+
+static cl_mem p_mem_objects[mem_object_hash_key] = {NULL};
+static volatile pthread_spinlock_t p_mem_objects_spinlock[mem_object_hash_key];
+
+#define LOCK(i) pthread_spin_lock(&p_mem_objects_spinlock[i]);
+#define UNLOCK(i) pthread_spin_unlock(&p_mem_objects_spinlock[i]);
+
+static void mem_object_init(void) {
+  int i;
+  for (i=0; i<mem_object_hash_key; i++) {
+    pthread_spin_init(&p_mem_objects_spinlock[i], 0);
+  }
+}
+
+static int mem_object_hash(const void * addr) {
+  uintptr_t t = (uintptr_t)addr;
+  uintptr_t t2 = t >> 4;
+  uintptr_t t3 = t2 % mem_object_hash_key;
+  return (int)t3;
+}
+
+static void mem_object_store(cl_mem m) {
+   int hash = mem_object_hash(m);
+
+   LOCK(hash);
+
+   m->prev = NULL;
+   m->next = p_mem_objects[hash];
+   if (p_mem_objects[hash] != NULL)
+      p_mem_objects[hash]->prev = m;
+   p_mem_objects[hash] = m;
+
+   UNLOCK(hash);
+}
+
+static void mem_object_release(cl_mem m) {
+
+   int hash = mem_object_hash(m);
+
+   LOCK(hash);
+
+   if (m->prev != NULL)
+     m->prev->next = m->next;
+   if (m->next != NULL)
+     m->next->prev = m->prev;
+
+   if (p_mem_objects[hash] == m) {
+      p_mem_objects[hash] = m->next;
+   }
+
+   UNLOCK(hash)
+}
+
+static cl_mem mem_object_fetch(const void * addr) {
+   int hash = mem_object_hash(*(cl_mem*)addr);
+
+   LOCK(hash);
+
+   cl_mem buf;
+   for (buf = p_mem_objects[hash]; buf != NULL; buf = buf->next) {
+      if (*(cl_mem*)addr == buf) {
+         UNLOCK(hash);
+         return buf;
+      }
+   }
+
+   UNLOCK(hash);
+   return NULL;
+}
+
+#undef LOCK
+#undef UNLOCK
+#undef mem_object_hash_key

socl/src/helper_task.c.inc

@@ -0,0 +1,118 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+static inline cl_event task_event(starpu_task *task) {
+  return (cl_event)task->callback_arg;
+}
+
+static void task_release_callback(void *arg) {
+  starpu_task *task = starpu_get_current_task();
+  cl_event ev = (cl_event)arg;
+  
+  ev->status = CL_COMPLETE;
+
+  if (task->profiling_info != NULL && (intptr_t)task->profiling_info != -ENOSYS) {
+    ev->profiling_info = malloc(sizeof(*task->profiling_info));
+    memcpy(ev->profiling_info, task->profiling_info, sizeof(*task->profiling_info));
+  }
+
+  gc_entity_release(ev);
+}
+
+
+/*
+ * Create a StarPU task
+ *
+ * Task's callback_arg is event
+ * Task's tag is set to event ID
+ */
+static starpu_task * task_create(cl_command_type type) {
+   cl_event event;
+   struct starpu_task * task;
+
+   /* Create event */
+   event = event_create();
+   event->type = type;
+
+   /* Create StarPU task */
+   task = starpu_task_create();
+
+   /* Task tag is set to event id */
+   task->use_tag = 1;
+   task->tag_id = event->id;
+
+   /* Set task common settings */
+   task->destroy = 1;
+   task->detach = 1;
+   task->callback_func = task_release_callback;
+   task->callback_arg = event;
+
+   return task;
+}
+
+
+static void task_dependency_add(starpu_task * task, cl_uint num, const cl_event *events) {
+   unsigned int i;
+
+   for (i=0; i<num; i++) {
+      starpu_tag_t tag = events[i]->id;
+      DEBUG_MSG("Event %d depends on event %d\n", task->tag_id, events[i]->id);
+      starpu_tag_declare_deps_array(task->tag_id, 1, &tag);
+   }
+}
+
+
+/*********************************
+ * CPU task helper
+ *********************************/
+
+struct cputask_arg {
+  void (*callback)(void*);
+  void * arg;
+  int free_arg;
+};
+
+static void cputask_task(__attribute__((unused)) void *descr[], void *args) {
+  struct cputask_arg * arg = (struct cputask_arg*)args;
+
+  arg->callback(arg->arg);
+
+  if (arg->free_arg)
+    free(arg->arg);
+
+  free(arg);
+}
+
+static starpu_codelet cputask_codelet = {
+   .where = STARPU_CPU,
+   .model = NULL,
+   .cpu_func = &cputask_task
+};
+
+static starpu_task * task_create_cpu(cl_command_type type, void (*callback)(void*), void *arg, int free_arg) {
+  
+  struct cputask_arg * a = malloc(sizeof(struct cputask_arg));
+  a->callback = callback;
+  a->arg = arg;
+  a->free_arg = free_arg;
+
+  starpu_task *task = task_create(type);
+  task->cl = &cputask_codelet;
+  task->cl_arg = a;
+
+  return task;
+}
+

socl/src/helper_workerid.c.inc

@@ -0,0 +1,24 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+int starpu_worker_get_range() {
+   int id = starpu_worker_get_id();
+   int i, oid = 0;
+   for (i=0; i<id; i++)
+      if (starpu_worker_get_type(i) == STARPU_OPENCL_WORKER) oid++;
+
+   return oid;
+}

socl/src/init.c.inc

@@ -0,0 +1,53 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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.
+ */
+
+/**
+ * Initialize SOCL
+ */
+__attribute__((constructor)) static void socl_init() {
+  
+  mem_object_init();
+
+  starpu_init(NULL);
+  
+  //Disable dataflow implicit dependencies
+  starpu_data_set_default_sequential_consistency_flag(0);
+
+  gc_start();
+}
+
+/**
+ * Shutdown SOCL
+ */
+__attribute__((destructor)) static void socl_shutdown() {
+
+  starpu_task_wait_for_all();
+
+  gc_stop();
+
+  starpu_task_wait_for_all();
+
+  int i = 0;
+  while (entities != NULL) {
+    i++;
+    entities = entities->next;
+  }
+
+  if (i != 0)
+    fprintf(stderr, "Unreleased entities: %d\n", i);
+
+  starpu_shutdown();
+}

socl/src/opencl.c

@@ -0,0 +1,343 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010,2011 University of Bordeaux
+ *
+ * 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 CL_HEADERS
+#include "CL/cl.h"
+#else
+#include CL_HEADERS "CL/cl.h"
+#endif
+
+
+#include <string.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <unistd.h>
+#include <pthread.h>
+
+#include <starpu.h>
+#include <starpu_opencl.h>
+#include <starpu_data_interfaces.h>
+#include <starpu_profiling.h>
+#include <starpu_task.h>
+
+typedef struct starpu_task starpu_task;
+
+#ifdef UNUSED
+#elif defined(__GNUC__)
+   #define UNUSED(x) UNUSED_ ## x __attribute__((unused))
+#else
+   #define UNUSED(x) x
+#endif
+
+#define RETURN_EVENT(ev, event) \
+   if (event != NULL) \
+      *event = ev; \
+   else\
+      gc_entity_release(ev);
+
+#include "helper_debug.c.inc"
+#include "helper_getinfo.c.inc"
+
+/**
+ * Entity that can be managed by the garbage collector
+ */
+typedef struct entity * entity;
+
+struct entity {
+  /* Reference count */
+  size_t refs;
+
+  /* Callback called on release */
+  void (*release_callback)(void*entity);
+
+  /* Next entity in garbage collector queue */
+  entity prev;
+  entity next;
+};
+
+/* OpenCL entities (context, command queues, buffers...) must use
+ * this macro as their first field */
+#define CL_ENTITY struct entity _entity;
+
+
+struct _cl_platform_id {};
+
+static struct _cl_platform_id socl_platform = {};
+
+static const char SOCL_PROFILE[] = "FULL_PROFILE";
+static const char SOCL_VERSION[] = "OpenCL 1.0 StarPU Edition (0.0.1)";
+static const char SOCL_PLATFORM_NAME[]    = "StarPU Platform";
+static const char SOCL_VENDOR[]  = "INRIA";
+static const char SOCL_PLATFORM_EXTENSIONS[] = "";
+
+struct _cl_context {
+  CL_ENTITY;
+
+  void (*pfn_notify)(const char *, const void *, size_t, void *);
+  void *user_data;
+
+  /* Associated devices */
+  cl_device_id * devices;
+  cl_uint num_devices;
+
+  /* Properties */
+  cl_context_properties * properties;
+  cl_uint num_properties;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+
+struct _cl_command_queue {
+  CL_ENTITY;
+
+  cl_command_queue_properties properties;
+  cl_device_id device;
+  cl_context context;
+
+  /* Stored command events */
+  cl_event events;
+
+  /* Last enqueued barrier-like event */
+  cl_event barrier;
+
+  /* Mutex */
+  pthread_spinlock_t spin;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+struct _cl_event {
+  CL_ENTITY;
+
+  /* Command queue */
+  cl_command_queue cq;
+
+  /* Command type */
+  cl_command_type type;
+
+  /* Command queue list */
+  cl_event prev;
+  cl_event next;
+
+  /* Event status */
+  cl_int status;
+
+  /* ID  
+   * This ID is used as a tag for StarPU dependencies
+   */
+  int id;
+
+  /* Profiling info are copied here */
+  struct starpu_task_profiling_info *profiling_info;
+};
+
+struct _cl_mem {
+  CL_ENTITY;
+
+  /* StarPU handle */
+  starpu_data_handle handle;
+
+  /* Pointer to data in host memory */
+  void *ptr;    
+
+  /* Buffer size */
+  size_t size;
+
+  /* Indicates how many references (mapping, MEM_USE_HOST_PTR...) require
+   * coherence in host memory. If set to zero, no coherency is maintained
+   * (this is the most efficient) */
+  int map_count; 
+
+  /* Creation flags */
+  cl_mem_flags flags;
+
+  /* Creation context */
+  cl_context context;
+
+  /* Access mode */
+  int mode;
+
+  /* Host ptr */
+  void * host_ptr;
+
+  /* Fields used to store cl_mems in mem_objects list */
+  cl_mem prev;
+  cl_mem next;
+
+  /* Indicates if a buffer may contain meaningful data. Otherwise
+     we don't have to transfer it */
+  int scratch;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+struct _cl_program {
+  CL_ENTITY;
+
+  /* Real OpenCL Programs
+   * There is one entry for each device (even non OpenCL ones)
+   * in order to index this array with dev_id
+   */
+  cl_program *cl_programs;
+
+  /* Context used to create this program */
+  cl_context context;
+
+  /* Options  */
+  char * options;
+  unsigned int options_size;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+enum kernel_arg_type { Null, Buffer, Immediate };
+
+struct _cl_kernel {
+  CL_ENTITY;
+
+  /* Associated program */
+  cl_program program;
+
+  /* Kernel name */
+  char * kernel_name;
+
+  /* Real OpenCL kernels */
+  cl_kernel *cl_kernels;
+
+  /* clCreateKernel return codes */
+  cl_int *errcodes;
+
+  /* Arguments */
+  unsigned int arg_count;
+  size_t *arg_size;
+  enum kernel_arg_type  *arg_type;
+  void  **arg_value;
+
+  /* ID  */
+#ifdef DEBUG
+  int id;
+#endif
+};
+
+/* Command queues with profiling enabled
+ * This allows us to disable StarPU profiling it
+ * is equal to 0
+ */
+static int profiling_queue_count = 0;
+
+#include "helper_workerid.c.inc"
+
+#include "gc.c.inc"
+
+#include "cl_getplatformids.c.inc"
+#include "cl_getplatforminfo.c.inc"
+
+#include "device_descriptions.c.inc"
+#include "cl_getdeviceids.c.inc"
+#include "cl_getdeviceinfo.c.inc"
+
+#include "cl_releasecontext.c.inc"
+#include "cl_createcontext.c.inc"
+#include "cl_createcontextfromtype.c.inc"
+#include "cl_retaincontext.c.inc"
+#include "cl_getcontextinfo.c.inc"
+
+#include "cl_releasecommandqueue.c.inc"
+#include "cl_createcommandqueue.c.inc"
+#include "cl_retaincommandqueue.c.inc"
+#include "cl_getcommandqueueinfo.c.inc"
+#include "cl_setcommandqueueproperty.c.inc"
+
+#include "cl_releaseevent.c.inc"
+#include "helper_event.c.inc"
+#include "helper_task.c.inc"
+#include "cl_waitforevents.c.inc"
+#include "cl_geteventinfo.c.inc"
+#include "cl_retainevent.c.inc"
+
+#include "helper_command_queue.c.inc"
+
+#include "cl_enqueuemarker.c.inc"
+#include "cl_enqueuewaitforevents.c.inc"
+#include "cl_enqueuebarrier.c.inc"
+#include "cl_flush.c.inc"
+#include "cl_finish.c.inc"
+
+#include "helper_mem_objects.c.inc"
+#include "cl_releasememobject.c.inc"
+#include "cl_createbuffer.c.inc"
+#include "cl_createimage2d.c.inc"
+#include "cl_createimage3d.c.inc"
+#include "cl_retainmemobject.c.inc"
+#include "cl_getsupportedimageformats.c.inc"
+#include "cl_getmemobjectinfo.c.inc"
+#include "cl_getimageinfo.c.inc"
+
+#include "cl_createsampler.c.inc"
+#include "cl_retainsampler.c.inc"
+#include "cl_releasesampler.c.inc"
+#include "cl_getsamplerinfo.c.inc"
+
+#include "cl_releaseprogram.c.inc"
+#include "cl_createprogramwithsource.c.inc"
+#include "cl_createprogramwithbinary.c.inc"
+#include "cl_retainprogram.c.inc"
+#include "cl_buildprogram.c.inc"
+#include "cl_unloadcompiler.c.inc"
+#include "cl_getprograminfo.c.inc"
+#include "cl_getprogrambuildinfo.c.inc"
+
+#include "cl_releasekernel.c.inc"
+#include "cl_createkernel.c.inc"
+#include "cl_createkernelsinprogram.c.inc"
+#include "cl_retainkernel.c.inc"
+#include "cl_setkernelarg.c.inc"
+#include "cl_getkernelinfo.c.inc"
+#include "cl_getkernelworkgroupinfo.c.inc"
+
+#include "cl_enqueuereadbuffer.c.inc"
+#include "cl_enqueuewritebuffer.c.inc"
+#include "cl_enqueuecopybuffer.c.inc"
+#include "cl_enqueuereadimage.c.inc"
+#include "cl_enqueuewriteimage.c.inc"
+#include "cl_enqueuecopyimage.c.inc"
+#include "cl_enqueuecopyimagetobuffer.c.inc"
+#include "cl_enqueuecopybuffertoimage.c.inc"
+#include "cl_enqueuemapbuffer.c.inc"
+#include "cl_enqueuemapimage.c.inc"
+#include "cl_enqueueunmapmemobject.c.inc"
+#include "cl_enqueuetask.c.inc"
+#include "cl_enqueuendrangekernel.c.inc"
+#include "cl_enqueuenativekernel.c.inc"
+
+#include "cl_geteventprofilinginfo.c.inc"
+#include "cl_getextensionfunctionaddress.c.inc"
+
+#include "init.c.inc"