summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorSerge Cohen <serge@chocolatnoir.net>2017-01-04 23:17:00 +0100
committerMatthias Vogelgesang <matthias.vogelgesang@kit.edu>2017-01-10 15:43:52 +0100
commit20e0efa57c172521397438bcf30081f2a1577019 (patch)
treefa6c4fa2afe2de1735311875a47fa1e960d5ed76
parent0e3cca9cb63656c88bb336de42c54bfadc16d827 (diff)
downloadufo-filters-20e0efa57c172521397438bcf30081f2a1577019.tar.gz
ufo-filters-20e0efa57c172521397438bcf30081f2a1577019.tar.bz2
ufo-filters-20e0efa57c172521397438bcf30081f2a1577019.tar.xz
ufo-filters-20e0efa57c172521397438bcf30081f2a1577019.zip
Include limited contrib/ directory
Cmake integration included a first contributed filter(s), building contrib requires the CMake option : -DWITH_CONTRIBS:BOOL=ON (this one is OFF by default). Added 4 new contributed filters : * Median-MAD rejection in 3D 3x3x3 box size * Median-MAD rejection in 2D but variable box size * OpenCL 1-liner and multi-input filter * Monitoring statistics of an image stream (min, max, mean, var)
-rw-r--r--CMakeLists.txt3
-rw-r--r--contrib/CMakeLists.txt26
-rw-r--r--contrib/sxc/CMakeLists.txt28
-rw-r--r--contrib/sxc/COPYING674
-rw-r--r--contrib/sxc/src/CMakeLists.txt102
-rw-r--r--contrib/sxc/src/kernels/med-mad-reject-2d.cl95
-rw-r--r--contrib/sxc/src/kernels/med-mad-reject.cl207
-rw-r--r--contrib/sxc/src/kernels/ocl-1liner-skel.cl56
-rw-r--r--contrib/sxc/src/kernels/stat-monitor.cl233
-rw-r--r--contrib/sxc/src/ufo-med-mad-reject-2d-task.c254
-rw-r--r--contrib/sxc/src/ufo-med-mad-reject-2d-task.h55
-rw-r--r--contrib/sxc/src/ufo-med-mad-reject-task.c340
-rw-r--r--contrib/sxc/src/ufo-med-mad-reject-task.h54
-rw-r--r--contrib/sxc/src/ufo-ocl-1liner-task.c294
-rw-r--r--contrib/sxc/src/ufo-ocl-1liner-task.h56
-rw-r--r--contrib/sxc/src/ufo-stat-monitor-task.c581
-rw-r--r--contrib/sxc/src/ufo-stat-monitor-task.h54
-rw-r--r--contrib/sxc/src/ufo-sxc-common.c158
-rw-r--r--contrib/sxc/src/ufo-sxc-common.h44
19 files changed, 3314 insertions, 0 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 15bd539..367fb83 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -69,4 +69,7 @@ endif ()
add_subdirectory(docs)
add_subdirectory(deps)
add_subdirectory(src)
+if ("${WITH_CONTRIBS}")
+ add_subdirectory(contribs)
+endif ()
#}}}
diff --git a/contrib/CMakeLists.txt b/contrib/CMakeLists.txt
new file mode 100644
index 0000000..76e2bb0
--- /dev/null
+++ b/contrib/CMakeLists.txt
@@ -0,0 +1,26 @@
+
+## This file is part of ufo-serge filter set.
+## Copyright (C) 2016 Serge Cohen
+##
+## 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/>.
+##
+## Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+##
+
+cmake_minimum_required(VERSION 2.6)
+project(ufo C)
+
+list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
+
+add_subdirectory(sxc)
diff --git a/contrib/sxc/CMakeLists.txt b/contrib/sxc/CMakeLists.txt
new file mode 100644
index 0000000..d7b6970
--- /dev/null
+++ b/contrib/sxc/CMakeLists.txt
@@ -0,0 +1,28 @@
+
+## This file is part of ufo-serge filter set.
+## Copyright (C) 2016 Serge Cohen
+##
+## 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/>.
+##
+## Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+##
+
+cmake_minimum_required(VERSION 2.6)
+project(ufo C)
+
+set(PKG_UFO_CORE_MIN_REQUIRED "0.6")
+
+list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
+
+add_subdirectory(src)
diff --git a/contrib/sxc/COPYING b/contrib/sxc/COPYING
new file mode 100644
index 0000000..94a9ed0
--- /dev/null
+++ b/contrib/sxc/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>.
diff --git a/contrib/sxc/src/CMakeLists.txt b/contrib/sxc/src/CMakeLists.txt
new file mode 100644
index 0000000..8f096b6
--- /dev/null
+++ b/contrib/sxc/src/CMakeLists.txt
@@ -0,0 +1,102 @@
+##
+## This file is part of ufo-serge filter set.
+## Copyright (C) 2016 Serge Cohen
+##
+## 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/>.
+##
+## Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+##
+
+
+cmake_minimum_required(VERSION 2.6)
+
+set(contribs_sxc_filter_SRCS
+ ufo-med-mad-reject-2d-task.c
+ ufo-med-mad-reject-task.c
+ ufo-ocl-1liner-task.c
+ ufo-stat-monitor-task.c
+ )
+
+set(contribs_sxc_aux_SRCS
+ ufo-sxc-common.c)
+
+file(GLOB contribs_sxc_filter_KERNELS "kernels/*.cl")
+
+include(ConfigurePaths)
+include(PkgConfigVars)
+include(CheckCSourceCompiles)
+
+configure_paths(CONTRIBS_SXC_FILTERS)
+
+find_package(OpenCL REQUIRED)
+find_package(PkgConfig REQUIRED)
+
+pkg_check_modules(UFO ufo>=${PKG_UFO_CORE_MIN_REQUIRED} REQUIRED)
+pkg_check_variable(ufo plugindir)
+pkg_check_variable(ufo kerneldir)
+
+add_definitions("-std=c99 -Wall -fPIC")
+add_definitions(-DG_LOG_DOMAIN="Ufo")
+
+include_directories(${CMAKE_CURRENT_SOURCE_DIR}
+ ${OPENCL_INCLUDE_DIRS}
+ ${UFO_INCLUDE_DIRS})
+
+link_directories(${UFO_LIBRARY_DIRS} ${OPENCL_LIBRARY_DIRS})
+
+set(CMAKE_REQUIRED_INCLUDES ${OPENCL_INCLUDE_DIRS})
+
+set(ufofilter_LIBS
+ m
+ ${UFO_LIBRARIES}
+ ${OpenCL_LIBRARIES})
+
+# build static auxiliary library first
+add_library(contribs_sxc_aux STATIC ${contribs_sxc_aux_SRCS})
+
+foreach(_src ${contribs_sxc_filter_SRCS})
+ # find plugin suffix
+ string(REGEX REPLACE "ufo-([^ \\.]+)-task.*" "\\1" task "${_src}")
+
+ # build string to get miscalleanous sources
+ string(REPLACE "-" "_" _misc ${task})
+ string(TOUPPER ${_misc} _misc_upper)
+
+ # create an option name and add this to disable filters
+ set(_misc "${_misc}_misc_SRCS")
+
+ string(REPLACE "-" "" _targetname ${task})
+ set(target "ufofilter${_targetname}")
+
+ # build single shared library per filter
+ if (${CMAKE_SYSTEM_NAME} MATCHES "Darwin")
+ add_library(${target} MODULE ${_src} ${${_misc}})
+ else()
+ add_library(${target} SHARED ${_src} ${${_misc}})
+ endif()
+
+ target_link_libraries(${target} ${ufofilter_LIBS} ${${_aux_libs}} contribs_sxc_aux)
+
+ list(APPEND all_targets ${target})
+
+ install(TARGETS ${target}
+ ARCHIVE DESTINATION ${UFO_PLUGINDIR}
+ LIBRARY DESTINATION ${UFO_PLUGINDIR})
+endforeach()
+
+
+# copy kernels
+foreach(_kernel ${contribs_sxc_filter_KERNELS})
+ install(FILES ${_kernel} DESTINATION ${UFO_KERNELDIR})
+endforeach()
diff --git a/contrib/sxc/src/kernels/med-mad-reject-2d.cl b/contrib/sxc/src/kernels/med-mad-reject-2d.cl
new file mode 100644
index 0000000..ce1e26d
--- /dev/null
+++ b/contrib/sxc/src/kernels/med-mad-reject-2d.cl
@@ -0,0 +1,95 @@
+/*
+ * Replacing pixels by local median value when detected as outliers based on MAD (2D box, variable size)
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+// The size of the box in which to perform the filtering
+// This one should be provided at compile time through a -D directive.
+// #define BOXSIZE 15
+
+// The threshold to apply on the MAD to reject (or not) the pixel's value
+
+typedef int ssize_t;
+
+constant const ssize_t half_box = (BOXSIZE-1)>>1;
+
+__kernel void
+med_mad_rej_2D (
+ __global float *iImage,
+ __global float *oFilteredImage,
+ const float iRejThresh
+ )
+{
+ size_t sizeX = get_global_size(0);
+ size_t sizeY = get_global_size(1);
+ size_t globalX = get_global_id(0);
+ size_t globalY = get_global_id(1);
+
+ // Getting the element of an image (both lvalue and rvalue) from index :
+#define FromImage(x,y,image) image[(x) + (y)*sizeX]
+
+ /* Getting all the value within the box and in the image : */
+ float v[BOXSIZE*BOXSIZE];
+
+ ssize_t cX, cY;
+ size_t index=0;
+ for ( cX=-half_box + globalX; (half_box + globalX + 1) != cX; ++cX) {
+ for ( cY=-half_box + globalY; (half_box + globalY + 1) != cY; ++cY) {
+ if ( (0 <= cX) && (sizeX > cX) && (0 <= cY) && (sizeY > cY) ) {
+ v[index] = FromImage(cX, cY, iImage);
+ ++index;
+ }
+ }
+ }
+
+ size_t num_px = index;
+ float swapper;
+
+ /* Computing the median : */
+ for (index = 0; num_px != (1+index); ++index) {
+ for (size_t j = 0; j!= (num_px-index); ++j ) {
+ swapper = v[j];
+ v[j] = min(v[j], v[j+1]);
+ v[j+1] = max(swapper, v[j+1]);
+ }
+ }
+ float med = v[num_px>>1];
+
+ /* Computing the MAD */
+ for ( index = 0; (num_px>>1) != index; ++index) {
+ v[index] = med - v[index];
+ }
+ for ( index = (num_px>>1) ; num_px != index; ++index) {
+ v[index] = v[index] - med;
+ }
+ for (index = 0; num_px != (1+index); ++index) {
+ for (size_t j = 0; j!= (num_px-index); ++j ) {
+ swapper = v[j];
+ v[j] = min(v[j], v[j+1]);
+ v[j+1] = max(swapper, v[j+1]);
+ }
+ }
+ float mad = v[num_px>>1];
+
+ /* Testing the current value : */
+ if ( fabs(FromImage(globalX, globalY, iImage) - med) < (mad*iRejThresh) )
+ FromImage(globalX, globalY, oFilteredImage) = FromImage(globalX, globalY, iImage);
+ else
+ FromImage(globalX, globalY, oFilteredImage) = med;
+}
diff --git a/contrib/sxc/src/kernels/med-mad-reject.cl b/contrib/sxc/src/kernels/med-mad-reject.cl
new file mode 100644
index 0000000..bcaa986
--- /dev/null
+++ b/contrib/sxc/src/kernels/med-mad-reject.cl
@@ -0,0 +1,207 @@
+/*
+ * Replacing pixels by local median value when detected as outliers based on MAD (3x3x3 box)
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+
+/* Kernel to reject ouliers based on distance to median larger than ... times mad. */
+__kernel void
+outliersRej_MedMad_3x3x3_f32(
+ __global float *iPreviousImage,
+ __global float *iThisImage,
+ __global float *iNextImage,
+ __global float *oFilteredImage,
+ const float iRejThresh
+ )
+{
+ bool Zfront = (0 == iPreviousImage);
+ bool Zrear = (0 == iNextImage);
+
+ size_t ioStride = get_global_size(0);
+ size_t globalX = get_global_id(0);
+ size_t globalY = get_global_id(1);
+
+ // Pixel in image buffers : globalX + globalY*ioStride
+ // Pixel in work item : localX, localY
+ // Pixel in __local buffer : [globalX][YinBuf]
+
+ // __local float groupBox[get_local_size(0)+2][get_local_size(1)+2][3];
+ // Conversion using the following RegExp (Perl style) : groupBox\[([-[:alnum:] .+*]*)\]\[([-[:alnum:] .+*]*)\]\[([-[:alnum:] .+*]*)\] => GB(\1, \2, \3)
+ //#define GB(a,b,c) groupBox[a + lSizeX * (b + lSizeY * c)]
+#define FromImage(x,y,image) image[(x) + (y)*ioStride]
+
+ // We can finally start the computation itself :
+ float v[27];
+
+ bool X_m1 = ( globalX ), X_p1 = ( (globalX + 1) != ioStride );
+ bool Y_m1 = ( globalY ), Y_p1 = ( (globalY + 1) != get_global_size(1) );
+
+ if ( Zfront ) {
+ v[0] = v[1] = v[2] = v[3] = v[4] = v[5] = v[6] = v[7] = v[8] = 0.0f;
+ }
+ else {
+ v[0] = ( X_m1 && Y_m1 ) ? FromImage(globalX-1, globalY-1, iPreviousImage) : 0.0f;
+ v[1] = ( X_m1 ) ? FromImage(globalX-1, globalY , iPreviousImage) : 0.0f;
+ v[2] = ( X_m1 && Y_p1 ) ? FromImage(globalX-1, globalY+1, iPreviousImage) : 0.0f;
+
+ v[3] = ( Y_m1 ) ? FromImage(globalX , globalY-1, iPreviousImage) : 0.0f;
+ v[4] = FromImage(globalX , globalY , iPreviousImage);
+ v[5] = ( Y_p1 ) ? FromImage(globalX , globalY+1, iPreviousImage) : 0.0f;
+
+ v[6] = ( X_p1 && Y_m1 ) ? FromImage(globalX+1, globalY-1, iPreviousImage) : 0.0f;
+ v[7] = ( X_p1 ) ? FromImage(globalX+1, globalY , iPreviousImage) : 0.0f;
+ v[8] = ( X_p1 && Y_p1 ) ? FromImage(globalX+1, globalY+1, iPreviousImage) : 0.0f;
+ }
+
+ v[9] = ( X_m1 && Y_m1 ) ? FromImage(globalX-1, globalY-1, iThisImage) : 0.0f;
+ v[10] = ( X_m1 ) ? FromImage(globalX-1, globalY , iThisImage) : 0.0f;
+ v[11] = ( X_m1 && Y_p1 ) ? FromImage(globalX-1, globalY+1, iThisImage) : 0.0f;
+
+ v[12] = ( Y_m1 ) ? FromImage(globalX , globalY-1, iThisImage) : 0.0f;
+ v[13] = FromImage(globalX , globalY , iThisImage);
+ v[14] = ( Y_p1 ) ? FromImage(globalX , globalY+1, iThisImage) : 0.0f;
+
+ v[15] = ( X_p1 && Y_m1 ) ? FromImage(globalX+1, globalY-1, iThisImage) : 0.0f;
+ v[16] = ( X_p1 ) ? FromImage(globalX+1, globalY , iThisImage) : 0.0f;
+ v[17] = ( X_p1 && Y_p1 ) ? FromImage(globalX+1, globalY+1, iThisImage) : 0.0f;
+
+ if ( Zrear ) {
+ v[18] = v[19] = v[20] = v[21] = v[22] = v[23] = v[24] = v[25] = v[26] = 0.0f;
+ }
+ else {
+ v[18] = ( X_m1 && Y_m1 ) ? FromImage(globalX-1, globalY-1, iNextImage) : 0.0f;
+ v[19] = ( X_m1 ) ? FromImage(globalX-1, globalY , iNextImage) : 0.0f;
+ v[20] = ( X_m1 && Y_p1 ) ? FromImage(globalX-1, globalY+1, iNextImage) : 0.0f;
+
+ v[21] = ( Y_m1 ) ? FromImage(globalX , globalY-1, iNextImage) : 0.0f;
+ v[22] = FromImage(globalX , globalY , iNextImage);
+ v[23] = ( Y_p1 ) ? FromImage(globalX , globalY+1, iNextImage) : 0.0f;
+
+ v[24] = ( X_p1 && Y_m1 ) ? FromImage(globalX+1, globalY-1, iNextImage) : 0.0f;
+ v[25] = ( X_p1 ) ? FromImage(globalX+1, globalY , iNextImage) : 0.0f;
+ v[26] = ( X_p1 && Y_p1 ) ? FromImage(globalX+1, globalY+1, iNextImage) : 0.0f;
+ }
+
+#undef b
+#undef bub
+#undef rabbit_up
+#undef rabbit_down
+#define b(a, b) {float swap=a; a=min(a, b); b=max(swap, b);}
+#define bub(a) b(v[a], v[a+1])
+#define rabbit_up bub(0) bub(1) bub(2) bub(3) bub(4) bub(5) bub(6) bub(7) bub(8) bub(9) bub(10) bub(11) bub(12) bub(13)
+#define rabbit_down bub(25) bub(24) bub(23) bub(22) bub(21) bub(20) bub(19) bub(18) bub(17) bub(16) bub(15) bub(14) bub(13) bub(12)
+
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+
+ float median = v[13];
+ // Now computing the MAD :
+#undef abs_diff_med
+#define abs_diff_med(a) v[a] = fabs(v[a] - median);
+ abs_diff_med(0)
+ abs_diff_med(1)
+ abs_diff_med(2)
+ abs_diff_med(3)
+ abs_diff_med(4)
+ abs_diff_med(5)
+ abs_diff_med(6)
+ abs_diff_med(7)
+ abs_diff_med(8)
+ abs_diff_med(9)
+ abs_diff_med(10)
+ abs_diff_med(11)
+ abs_diff_med(12)
+ abs_diff_med(13)
+ abs_diff_med(14)
+ abs_diff_med(15)
+ abs_diff_med(16)
+ abs_diff_med(17)
+ abs_diff_med(18)
+ abs_diff_med(19)
+ abs_diff_med(20)
+ abs_diff_med(21)
+ abs_diff_med(22)
+ abs_diff_med(23)
+ abs_diff_med(24)
+ abs_diff_med(25)
+ abs_diff_med(26)
+ // And the median of that again :
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+ rabbit_up
+ rabbit_down
+
+ float mad = v[13];
+
+ // If the distance to the median is larger than iRejThresh*mad, then replace pixel's value by the median :
+
+ if ( fabs(FromImage(globalX , globalY , iThisImage) - median) < (iRejThresh*mad) )
+ oFilteredImage[globalX + globalY*ioStride] = FromImage(globalX , globalY , iThisImage);
+ else
+ oFilteredImage[globalX + globalY*ioStride] = median;
+
+ // oFilteredImage[globalX + globalY*ioStride] = mad;
+
+}
diff --git a/contrib/sxc/src/kernels/ocl-1liner-skel.cl b/contrib/sxc/src/kernels/ocl-1liner-skel.cl
new file mode 100644
index 0000000..55bfa88
--- /dev/null
+++ b/contrib/sxc/src/kernels/ocl-1liner-skel.cl
@@ -0,0 +1,56 @@
+/*
+ * Instant compiling a one liner OpenCL filter
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+/*
+ * In this kernel, the one line imported can use multiple ways to address the
+ * modified pixel.
+ * In general the model is that there is exactly one work-item per pixel and
+ * the work is performed on a 2D grid matching exactly the image diemensions.
+ *
+ * in0 .. inN are input 1D array(s).
+ * out is 1D output array.
+ *
+ * Those can be addressed using px_index which is the /current/ pixel for the workitem.
+ *
+ * One can also access a /random/ pixel value using the macro : IMG_VAL which takes three
+ * arguments : the (x, y) pixel coordinate, and the array pointer for the image.
+ * NB : this macro acn be used also for lvalue (but it is risky that a work-item
+ * modifies a pixel value in =out= that is NOT the one indexed by px_index.
+ */
+
+#define IMG_VAL(hx,hy,image) image[(hx) + (hy)*sizeX]
+
+__kernel void
+ocl_1liner (
+ %s // The inputs...
+ __global float *out,
+ )
+{
+ size_t sizeX = get_global_size(0);
+ size_t sizeY = get_global_size(1);
+ size_t x = get_global_id(0);
+ size_t y = get_global_id(1);
+
+ size_t px_index = x + y*sizeX;
+
+ // And here the one line of /variable/ code :
+ %s;
+}
diff --git a/contrib/sxc/src/kernels/stat-monitor.cl b/contrib/sxc/src/kernels/stat-monitor.cl
new file mode 100644
index 0000000..21b71ac
--- /dev/null
+++ b/contrib/sxc/src/kernels/stat-monitor.cl
@@ -0,0 +1,233 @@
+/*
+ * Gathering statistics on a image stream, copying input to output
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+__kernel void
+stat_monitor_f32(__global float *in_img, // Input image to reduce (will compute min, max, sum, sum of sq per workgroup)
+ __global float *out_stat, // Output : one 4-tupple per workgroup ordered (min, max, sum, sum of sq)
+ __const uint img_size, // The total number of elements in the image to reduce
+ __local float *local_scr // A memory scrathpad for within workgroup final reduction
+ )
+{
+ size_t gi = get_global_id(0);
+ float wi_min = in_img[gi];
+ float wi_max = in_img[gi];
+ float wi_sum = 0.0f;
+ float wi_sum_sq = 0.0f;
+
+ // First part of the reduction : doing it sequentially as much as possible in each work-item :
+ while (gi < img_size) {
+ float here_val = in_img[gi];
+ wi_min = fmin(wi_min, here_val);
+ wi_max = fmax(wi_max, here_val);
+ wi_sum += here_val;
+ wi_sum_sq += here_val * here_val;
+ gi += get_global_size(0);
+ }
+
+ // Further reduce all the work-items of the current work-group using parallel reduce :
+ size_t li = get_local_id(0);
+ size_t li_b = li << 2; // there is 4 elements per work-items
+
+ // Initing the parallel reduce
+ local_scr[li_b ] = wi_min;
+ local_scr[li_b+1] = wi_max;
+ local_scr[li_b+2] = wi_sum;
+ local_scr[li_b+3] = wi_sum_sq;
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ // Reducing recursively by a factor of 2 :
+ for(size_t offset = get_local_size(0) >> 1; offset != 0; offset >>= 1) {
+ if (li < offset) {
+ size_t sof = offset<<2;
+ local_scr[li_b ] = fmin(local_scr[li_b ], local_scr[li_b + sof]);
+ local_scr[li_b+1] = fmax(local_scr[li_b+1], local_scr[li_b+1 + sof]);
+ local_scr[li_b+2] += local_scr[li_b+2 + sof];
+ local_scr[li_b+3] += local_scr[li_b+3 + sof];
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+ if (li == 0) {
+ size_t gr_id = get_group_id(0) << 2; // 4 datas per work-group
+ out_stat[gr_id ] = local_scr[0];
+ out_stat[gr_id+1] = local_scr[1];
+ out_stat[gr_id+2] = local_scr[2];
+ out_stat[gr_id+3] = local_scr[3];
+ }
+
+}
+
+/***************************************************************/
+/* Performing the last round of reduction in workgroup 0 alone */
+/***************************************************************/
+
+/*
+ * Unfortunately the only way to synchronize all worgroups is to
+ * launch another kernel :-(
+ */
+
+/* Performing purely parallel reduce based on the workgroup size and number of elements in the
+ * input.
+ *
+ * NB : This kernel is supposed to be run in a single workgroup to fiish up the reduction
+ * mostly computed by the above kernel.
+ */
+__kernel void
+stat_monitor_f32_fin(__global float *red_in,
+ __global float *red_out,
+ __const uint num_elts,
+ __local float *local_scr // A memory scrathpad for within workgroup final reduction
+ )
+{
+ // Since there is a single workgroup, locol_id and global_id are identical.
+ // The total number of workitems is (pproximately) half the num_elts
+ size_t gi = get_global_id(0);
+ size_t grs = get_local_size(0);
+
+ // Copying all data to the local memory, for speed during the reduction
+ event_t copied_to_scr = async_work_group_copy(local_scr, red_in, num_elts<<2, 0);
+ // Synchronising on the copy being done
+ wait_group_events(1, &copied_to_scr);
+
+ // Performing the reduction
+ for ( size_t offset = grs; 0 != offset; offset >>=1 ) {
+ if ( (gi < offset) && (gi+offset < num_elts) ) {
+ size_t gi_4B = gi << 2;
+ size_t gi_of_4B = (gi + offset) << 2;
+ local_scr[gi_4B ] = fmin(local_scr[gi_4B ], local_scr[gi_of_4B ]);
+ local_scr[gi_4B+1] = fmax(local_scr[gi_4B+1], local_scr[gi_of_4B+1]);
+ local_scr[gi_4B+2] += local_scr[gi_4B+2];
+ local_scr[gi_4B+3] += local_scr[gi_4B+3];
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+
+ // Returning the results
+ event_t copied_from_scr = async_work_group_copy(red_out, local_scr, 4, 0);
+ wait_group_events(1, &copied_from_scr);
+}
+
+
+
+#ifdef cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+__kernel void
+stat_monitor_f64(__global float *in_img, // Input image to reduce (will compute min, max, sum, sum of sq per workgroup)
+ __global double *out_stat, // Output : one 4-tupple per workgroup ordered (min, max, sum, sum of sq)
+ __const uint img_size, // The total number of elements in the image to reduce
+ __local double *local_scr // A memory scrathpad for within workgroup final reduction
+ )
+{
+ size_t gi = get_global_id(0);
+ double wi_min = convert_double(in_img[gi]);
+ double wi_max = convert_double(in_img[gi]);
+ double wi_sum = 0.0;
+ double wi_sum_sq = 0.0;
+
+ // First part of the reduction : doing it sequentially as much as possible in each work-item :
+ while (gi < img_size) {
+ double here_val = convert_double(in_img[gi]);
+ wi_min = fmin(wi_min, here_val);
+ wi_max = fmax(wi_max, here_val);
+ wi_sum += here_val;
+ wi_sum_sq += here_val * here_val;
+ gi += get_global_size(0);
+ }
+
+ // Further reduce all the work-items of the current work-group using parallel reduce :
+ size_t li = get_local_id(0);
+ size_t li_b = li << 2; // there is 4 elements per work-items
+
+ // Initing the parallel reduce
+ local_scr[li_b ] = wi_min;
+ local_scr[li_b+1] = wi_max;
+ local_scr[li_b+2] = wi_sum;
+ local_scr[li_b+3] = wi_sum_sq;
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ // Reducing recursively by a factor of 2 :
+ for(size_t offset = get_local_size(0) >> 1; offset != 0; offset >>= 1) {
+ if (li < offset) {
+ size_t sof = offset<<2;
+ local_scr[li_b ] = fmin(local_scr[li_b ], local_scr[li_b + sof]);
+ local_scr[li_b+1] = fmax(local_scr[li_b+1], local_scr[li_b+1 + sof]);
+ local_scr[li_b+2] += local_scr[li_b+2 + sof];
+ local_scr[li_b+3] += local_scr[li_b+3 + sof];
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+ if (li == 0) {
+ size_t gr_id = get_group_id(0) << 2; // 4 datas per work-group
+ out_stat[gr_id ] = local_scr[0];
+ out_stat[gr_id+1] = local_scr[1];
+ out_stat[gr_id+2] = local_scr[2];
+ out_stat[gr_id+3] = local_scr[3];
+ }
+
+}
+
+
+/***************************************************************/
+/* Performing the last round of reduction in workgroup 0 alone */
+/***************************************************************/
+
+/* Performing purely parallel reduce based on the workgroup size and number of elements in the
+ * input.
+ *
+ * NB : This kernel is supposed to be run in a single workgroup to fiish up the reduction
+ * mostly computed by the above kernel.
+ */
+__kernel void
+stat_monitor_f64_fin(__global double *red_in,
+ __global double *red_out,
+ __const uint num_elts,
+ __local double *local_scr // A memory scrathpad for within workgroup final reduction
+ )
+{
+ // Since there is a single workgroup, locol_id and global_id are identical.
+ // The total number of workitems is (pproximately) half the num_elts
+ size_t gi = get_global_id(0);
+ size_t grs = get_local_size(0);
+
+ // Copying all data to the local memory, for speed during the reduction
+ event_t copied_to_scr = async_work_group_copy(local_scr, red_in, num_elts<<2, 0);
+ // Synchronising on the copy being done
+ wait_group_events(1, &copied_to_scr);
+
+ // Performing the reduction
+ for ( size_t offset = grs; 0 != offset; offset >>=1 ) {
+ if ( (gi < offset) && (gi+offset < num_elts) ) {
+ size_t gi_4B = gi << 2;
+ size_t gi_of_4B = (gi + offset) << 2;
+ local_scr[gi_4B ] = fmin(local_scr[gi_4B ], local_scr[gi_of_4B ]);
+ local_scr[gi_4B+1] = fmax(local_scr[gi_4B+1], local_scr[gi_of_4B+1]);
+ local_scr[gi_4B+2] += local_scr[gi_4B+2];
+ local_scr[gi_4B+3] += local_scr[gi_4B+3];
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+
+ // Returning the results
+ event_t copied_from_scr = async_work_group_copy(red_out, local_scr, 4, 0);
+ wait_group_events(1, &copied_from_scr);
+}
+
+#endif // cl_khr_fp64
diff --git a/contrib/sxc/src/ufo-med-mad-reject-2d-task.c b/contrib/sxc/src/ufo-med-mad-reject-2d-task.c
new file mode 100644
index 0000000..14e7dc0
--- /dev/null
+++ b/contrib/sxc/src/ufo-med-mad-reject-2d-task.c
@@ -0,0 +1,254 @@
+/*
+ * Replacing pixels by local median value when detected as outliers based on MAD (2D box, variable size)
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+#include <stdio.h>
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include "ufo-med-mad-reject-2d-task.h"
+
+
+struct _UfoMedMadReject2DTaskPrivate {
+ gfloat threshold;
+ guint32 box_size;
+ cl_kernel kernel;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoMedMadReject2DTask, ufo_med_mad_reject_2d_task, UFO_TYPE_TASK_NODE,
+ G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+ ufo_task_interface_init))
+
+#define UFO_MED_MAD_REJECT_2D_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_MED_MAD_REJECT_2D_TASK, UfoMedMadReject2DTaskPrivate))
+
+enum {
+ PROP_0,
+ PROP_THRESHOLD,
+ PROP_BOX_SIZE,
+ N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_med_mad_reject_2d_task_new (void)
+{
+ return UFO_NODE (g_object_new (UFO_TYPE_MED_MAD_REJECT_2D_TASK, NULL));
+}
+
+static void
+ufo_med_mad_reject_2d_task_setup (UfoTask *task,
+ UfoResources *resources,
+ GError **error)
+{
+ UfoMedMadReject2DTaskPrivate *priv;
+
+ priv = UFO_MED_MAD_REJECT_2D_TASK_GET_PRIVATE (task);
+
+ // Checking that the threshold is positive (being 0, makes it a median filter)
+ if ( priv->threshold <= 0.0f ) {
+ g_set_error (error, UFO_TASK_ERROR, UFO_TASK_ERROR_SETUP,
+ "Threshold value %f is not positive",
+ priv->threshold);
+ return;
+ }
+
+ // Checking that the box size is odd :
+ if ( ! (0x1 & priv->box_size) ) {
+ g_set_error (error, UFO_TASK_ERROR, UFO_TASK_ERROR_SETUP,
+ "Boxsize value %u is not odd",
+ priv->box_size);
+ return;
+ }
+
+ char kernel_opts[1024];
+ snprintf(kernel_opts, 1023, "-DBOXSIZE=%u", priv->box_size);
+ priv->kernel = ufo_resources_get_kernel_with_opts (resources, "med-mad-reject-2d.cl", "med_mad_rej_2D", kernel_opts, error);
+
+ if (priv->kernel != NULL)
+ UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel));
+}
+
+static void
+ufo_med_mad_reject_2d_task_get_requisition (UfoTask *task,
+ UfoBuffer **inputs,
+ UfoRequisition *requisition)
+{
+ // requisition->n_dims = 0;
+ // Transferring the buffer «size» (size, data type and dimension) to the output
+ ufo_buffer_get_requisition(inputs[0], requisition);
+}
+
+static guint
+ufo_med_mad_reject_2d_task_get_num_inputs (UfoTask *task)
+{
+ return 1;
+}
+
+static guint
+ufo_med_mad_reject_2d_task_get_num_dimensions (UfoTask *task,
+ guint input)
+{
+ return 2;
+}
+
+static UfoTaskMode
+ufo_med_mad_reject_2d_task_get_mode (UfoTask *task)
+{
+ return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU; // This is a processor type task, using GPU.
+}
+
+static gboolean
+ufo_med_mad_reject_2d_task_process (UfoTask *task,
+ UfoBuffer **inputs,
+ UfoBuffer *output,
+ UfoRequisition *requisition)
+{
+ UfoMedMadReject2DTaskPrivate *priv;
+ priv = UFO_MED_MAD_REJECT_2D_TASK_GET_PRIVATE (task);
+
+ UfoGpuNode *node;
+ UfoProfiler *profiler;
+ cl_command_queue cmd_queue;
+ cl_mem in_mem;
+ cl_mem out_mem;
+
+ node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
+ cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+
+ // making sure the input image is in GPU memory :
+ in_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue);
+ out_mem = ufo_buffer_get_device_array (output, cmd_queue);
+
+ // Setting up the kernel to run :
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof (cl_mem), &in_mem));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof (cl_mem), &out_mem));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof (cl_float), &priv->threshold));
+
+ profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task));
+ ufo_profiler_call (profiler, cmd_queue, priv->kernel, 2, requisition->dims, NULL);
+
+ return TRUE;
+}
+
+
+static void
+ufo_med_mad_reject_2d_task_set_property (GObject *object,
+ guint property_id,
+ const GValue *value,
+ GParamSpec *pspec)
+{
+ UfoMedMadReject2DTaskPrivate *priv = UFO_MED_MAD_REJECT_2D_TASK_GET_PRIVATE (object);
+
+ switch (property_id) {
+ case PROP_THRESHOLD:
+ priv->threshold = g_value_get_float (value);
+ break;
+ case PROP_BOX_SIZE:
+ priv->box_size = g_value_get_uint (value);
+ break;
+ default:
+ G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+ break;
+ }
+}
+
+static void
+ufo_med_mad_reject_2d_task_get_property (GObject *object,
+ guint property_id,
+ GValue *value,
+ GParamSpec *pspec)
+{
+ UfoMedMadReject2DTaskPrivate *priv = UFO_MED_MAD_REJECT_2D_TASK_GET_PRIVATE (object);
+
+ switch (property_id) {
+ case PROP_THRESHOLD:
+ g_value_set_float (value, priv->threshold);
+ break;
+ case PROP_BOX_SIZE:
+ g_value_set_uint (value, priv->box_size);
+ break;
+ default:
+ G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+ break;
+ }
+}
+
+static void
+ufo_med_mad_reject_2d_task_finalize (GObject *object)
+{
+ UfoMedMadReject2DTaskPrivate *priv = UFO_MED_MAD_REJECT_2D_TASK_GET_PRIVATE (object);
+
+ if ( priv->kernel )
+ UFO_RESOURCES_CHECK_CLERR(clReleaseKernel(priv->kernel));
+
+ G_OBJECT_CLASS (ufo_med_mad_reject_2d_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+ iface->setup = ufo_med_mad_reject_2d_task_setup;
+ iface->get_num_inputs = ufo_med_mad_reject_2d_task_get_num_inputs;
+ iface->get_num_dimensions = ufo_med_mad_reject_2d_task_get_num_dimensions;
+ iface->get_mode = ufo_med_mad_reject_2d_task_get_mode;
+ iface->get_requisition = ufo_med_mad_reject_2d_task_get_requisition;
+ iface->process = ufo_med_mad_reject_2d_task_process;
+}
+
+static void
+ufo_med_mad_reject_2d_task_class_init (UfoMedMadReject2DTaskClass *klass)
+{
+ GObjectClass *oclass = G_OBJECT_CLASS (klass);
+
+ oclass->set_property = ufo_med_mad_reject_2d_task_set_property;
+ oclass->get_property = ufo_med_mad_reject_2d_task_get_property;
+ oclass->finalize = ufo_med_mad_reject_2d_task_finalize;
+
+ properties[PROP_THRESHOLD] =
+ g_param_spec_float ("threshold",
+ "Rejection threshold",
+ "Pixel value replaced by median when more than threshold away from it (in term of MAD).",
+ 0.0f, G_MAXFLOAT, 3.0f,
+ G_PARAM_READWRITE);
+
+ properties[PROP_BOX_SIZE] =
+ g_param_spec_float ("box-size",
+ "Size of the box in which the median and mad are computed",
+ "Should be an odd number so that current pixel is the exact center of the box.",
+ 1, 1023, 3,
+ G_PARAM_READWRITE);
+
+ for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+ g_object_class_install_property (oclass, i, properties[i]);
+ g_type_class_add_private (oclass, sizeof(UfoMedMadReject2DTaskPrivate));
+}
+
+static void
+ufo_med_mad_reject_2d_task_init(UfoMedMadReject2DTask *self)
+{
+ self->priv = UFO_MED_MAD_REJECT_2D_TASK_GET_PRIVATE(self);
+}
diff --git a/contrib/sxc/src/ufo-med-mad-reject-2d-task.h b/contrib/sxc/src/ufo-med-mad-reject-2d-task.h
new file mode 100644
index 0000000..a16a864
--- /dev/null
+++ b/contrib/sxc/src/ufo-med-mad-reject-2d-task.h
@@ -0,0 +1,55 @@
+/*
+ * Replacing pixels by local median value when detected as outliers based on MAD (2D box, variable size)
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+#ifndef __UFO_MED_MAD_REJECT_2D_TASK_H
+#define __UFO_MED_MAD_REJECT_2D_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_MED_MAD_REJECT_2D_TASK (ufo_med_mad_reject_2d_task_get_type())
+#define UFO_MED_MAD_REJECT_2D_TASK(obj) (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_MED_MAD_REJECT_2D_TASK, UfoMedMadReject2DTask))
+#define UFO_IS_MED_MAD_REJECT_2D_TASK(obj) (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_MED_MAD_REJECT_2D_TASK))
+#define UFO_MED_MAD_REJECT_2D_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_MED_MAD_REJECT_2D_TASK, UfoMedMadReject2DTaskClass))
+#define UFO_IS_MED_MAD_REJECT_2D_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_MED_MAD_REJECT_2D_TASK))
+#define UFO_MED_MAD_REJECT_2D_TASK_GET_CLASS(obj) (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_MED_MAD_REJECT_2D_TASK, UfoMedMadReject2DTaskClass))
+
+typedef struct _UfoMedMadReject2DTask UfoMedMadReject2DTask;
+typedef struct _UfoMedMadReject2DTaskClass UfoMedMadReject2DTaskClass;
+typedef struct _UfoMedMadReject2DTaskPrivate UfoMedMadReject2DTaskPrivate;
+
+struct _UfoMedMadReject2DTask {
+ UfoTaskNode parent_instance;
+
+ UfoMedMadReject2DTaskPrivate *priv;
+};
+
+struct _UfoMedMadReject2DTaskClass {
+ UfoTaskNodeClass parent_class;
+};
+
+UfoNode *ufo_med_mad_reject_2d_task_new (void);
+GType ufo_med_mad_reject_2d_task_get_type (void);
+
+G_END_DECLS
+
+#endif
diff --git a/contrib/sxc/src/ufo-med-mad-reject-task.c b/contrib/sxc/src/ufo-med-mad-reject-task.c
new file mode 100644
index 0000000..5a1a09c
--- /dev/null
+++ b/contrib/sxc/src/ufo-med-mad-reject-task.c
@@ -0,0 +1,340 @@
+/*
+ * Replacing pixels by local median value when detected as outliers based on MAD
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+#include "ufo-med-mad-reject-task.h"
+#include "ufo-sxc-common.h"
+
+#include <stdio.h>
+
+/*
+ In general for a type REDUCTOR : process is called to input data and generate to generate the output.
+ For the current reject : the first call to process should return TRUE then FALSE (hence we can generate
+ the output for first frame).
+
+ Then we should return FALSE to each new call to process (so that we have a chance to generate the filtered
+ frame).
+
+ When generate is called WITHOUT a call to process before : it means we are dealing with last frame … act accordingly
+*/
+
+struct _UfoMedMadRejectTaskPrivate {
+ gfloat threshold;
+ cl_kernel kernel;
+ UfoBuffer *in0;
+ UfoBuffer *in1;
+ UfoBuffer *in2;
+ gboolean proc_was_called;
+ gboolean gene_shortcut;
+ gint32 proc_call_number;
+ gint32 gene_call_number;
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoMedMadRejectTask, ufo_med_mad_reject_task, UFO_TYPE_TASK_NODE,
+ G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+ ufo_task_interface_init))
+
+#define UFO_MED_MAD_REJECT_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_MED_MAD_REJECT_TASK, UfoMedMadRejectTaskPrivate))
+
+enum {
+ PROP_0,
+ PROP_THRESHOLD,
+ N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_med_mad_reject_task_new (void)
+{
+ return UFO_NODE (g_object_new (UFO_TYPE_MED_MAD_REJECT_TASK, NULL));
+}
+
+// Only called at the start : should be used to setup the reject.
+static void
+ufo_med_mad_reject_task_setup (UfoTask *task,
+ UfoResources *resources,
+ GError **error)
+{
+ UfoMedMadRejectTaskPrivate *priv;
+
+ priv = UFO_MED_MAD_REJECT_TASK_GET_PRIVATE (task);
+
+ if ( priv->threshold <= 0.0f ) {
+ g_set_error (error, UFO_TASK_ERROR, UFO_TASK_ERROR_SETUP,
+ "Threshold value %f is not positive",
+ priv->threshold);
+ return;
+ }
+
+ priv->kernel = ufo_resources_get_kernel (resources, "med-mad-reject.cl", "outliersRej_MedMad_3x3x3_f32", error);
+
+ if (priv->kernel != NULL)
+ UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel));
+}
+
+// This one is called for each frame, so that it can adapt from one to the other frame
+static void
+ufo_med_mad_reject_task_get_requisition (UfoTask *task,
+ UfoBuffer **inputs,
+ UfoRequisition *requisition)
+{
+ // requisition->n_dims = 0;
+ // Transferring the buffer «size» (size, data type and dimension) to the output
+ ufo_buffer_get_requisition(inputs[0], requisition);
+}
+
+// Only called at the start
+static guint
+ufo_med_mad_reject_task_get_num_inputs (UfoTask *task)
+{
+ return 1;
+}
+
+// Only called at the start (returns the dimension of data returned when processing each frame)
+static guint
+ufo_med_mad_reject_task_get_num_dimensions (UfoTask *task,
+ guint input)
+{
+ return 2;
+}
+
+// Only called at the start
+static UfoTaskMode
+ufo_med_mad_reject_task_get_mode (UfoTask *task)
+{
+ // return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU; // Requesting also the GPU for computation
+ return UFO_TASK_MODE_REDUCTOR | UFO_TASK_MODE_GPU; // Requesting also the GPU for computation
+}
+
+// Actual process function :
+static gboolean
+ufo_med_mad_reject_task_process (UfoTask *task,
+ UfoBuffer **inputs,
+ UfoBuffer *output,
+ UfoRequisition *requisition)
+{
+ UfoMedMadRejectTaskPrivate *priv;
+ UfoBuffer *swapper;
+
+ priv = UFO_MED_MAD_REJECT_TASK_GET_PRIVATE (task);
+ g_debug("%s (l%d) : entering process with index %d, generate index %d, proc_was_called %s", __func__, __LINE__, priv->proc_call_number, priv->gene_call_number, (priv->proc_was_called)? "TRUE" : "FALSE");
+
+ switch ( priv->proc_call_number ) {
+ case 0:
+ priv->in0 = ufo_buffer_dup(inputs[0]);
+ priv->in1 = ufo_buffer_dup(inputs[0]);
+ priv->in2 = ufo_buffer_dup(inputs[0]);
+
+ ufo_buffer_copy(inputs[0], priv->in0);
+ ufo_buffer_copy(inputs[0], priv->in1);
+ priv->proc_call_number += 1;
+ priv->proc_was_called = TRUE;
+ g_debug ("%s (l%d) : returning TRUE from process with process index %d, generate index %d, proc_was_called %s", __func__, __LINE__, priv->proc_call_number, priv->gene_call_number, (priv->proc_was_called)? "TRUE" : "FALSE");
+ return TRUE;
+
+ case 1:
+ ufo_buffer_copy(inputs[0], priv->in2);
+ break;
+
+ default :
+ swapper = priv->in0;
+ priv->in0 = priv->in1;
+ priv->in1 = priv->in2;
+ priv->in2 = swapper;
+
+ ufo_buffer_copy(inputs[0], priv->in2);
+ break;
+ }
+
+ priv->proc_call_number += 1;
+ priv->proc_was_called = TRUE;
+
+ g_debug ("%s (l%d) : returning FALSE from process with process index %d, generate index %d, proc_was_called %s", __func__, __LINE__, priv->proc_call_number, priv->gene_call_number, (priv->proc_was_called)? "TRUE" : "FALSE");
+ return FALSE; // returning TRUE means ready to process another frame (kind of)
+ // On REDUCTOR FALSE means scheduler has to call generate
+}
+
+static gboolean
+ufo_med_mad_reject_task_generate (UfoTask *task,
+ UfoBuffer *output,
+ UfoRequisition *requisition)
+{
+ UfoMedMadRejectTaskPrivate *priv;
+ UfoGpuNode *node;
+ UfoProfiler *profiler;
+ cl_command_queue cmd_queue;
+ UfoBuffer *swapper=NULL;
+
+ cl_mem in0_mem, in1_mem, in2_mem;
+ cl_mem out_mem;
+
+ priv = UFO_MED_MAD_REJECT_TASK_GET_PRIVATE (task);
+
+ if ( priv->gene_shortcut ) {
+ g_debug ("%s (l%d) : shortcutting generate with process index %d, generate index %d, proc_was_called %s", __func__, __LINE__, priv->proc_call_number, priv->gene_call_number, (priv->proc_was_called)? "TRUE" : "FALSE");
+ priv->gene_shortcut = FALSE;
+ return FALSE;
+ }
+
+ node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
+ cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+ out_mem = ufo_buffer_get_device_array (output, cmd_queue);
+
+ g_debug ("%s (l%d) : entering generate with process index %d, generate index %d, proc_was_called %s", __func__, __LINE__, priv->proc_call_number, priv->gene_call_number, (priv->proc_was_called)? "TRUE" : "FALSE");
+
+ if ( !priv->proc_was_called ) { // Special case, we are handling the last view
+ swapper = priv->in0;
+ priv->in0 = priv->in1;
+ priv->in1 = priv->in2;
+ }
+
+ in0_mem = ufo_buffer_get_device_array (priv->in0, cmd_queue);
+ in1_mem = ufo_buffer_get_device_array (priv->in1, cmd_queue);
+ in2_mem = ufo_buffer_get_device_array (priv->in2, cmd_queue);
+
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof (cl_mem), &in0_mem));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof (cl_mem), &in1_mem));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof (cl_mem), &in2_mem));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 3, sizeof (cl_mem), &out_mem));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 4, sizeof (cl_float), &priv->threshold));
+
+ profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task));
+ ufo_profiler_call (profiler, cmd_queue, priv->kernel, 2, requisition->dims, NULL);
+
+ if ( !priv->proc_was_called ) { // Special case, we are handling the last view
+ priv->in2 = swapper;
+ // so that all buffers are released at object deallocation.
+ }
+
+ priv->proc_was_called = FALSE;
+ priv->gene_shortcut = ( 0 != priv->gene_call_number );;
+
+ priv->gene_call_number += 1;
+
+ g_debug ("%s (l%d) : shortcut %s from generate with process index %d, generate index %d, proc_was_called %s", __func__, __LINE__, (priv->gene_shortcut)? "TRUE":"FALSE", priv->proc_call_number, priv->gene_call_number, (priv->proc_was_called)? "TRUE" : "FALSE");
+ return TRUE;
+}
+
+static void
+ufo_med_mad_reject_task_set_property (GObject *object,
+ guint property_id,
+ const GValue *value,
+ GParamSpec *pspec)
+{
+ UfoMedMadRejectTaskPrivate *priv = UFO_MED_MAD_REJECT_TASK_GET_PRIVATE (object);
+
+ switch (property_id) {
+ case PROP_THRESHOLD:
+ priv->threshold = g_value_get_float (value);
+ break;
+ default:
+ G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+ break;
+ }
+}
+
+static void
+ufo_med_mad_reject_task_get_property (GObject *object,
+ guint property_id,
+ GValue *value,
+ GParamSpec *pspec)
+{
+ UfoMedMadRejectTaskPrivate *priv = UFO_MED_MAD_REJECT_TASK_GET_PRIVATE (object);
+
+ switch (property_id) {
+ case PROP_THRESHOLD:
+ g_value_set_float (value, priv->threshold);
+ break;
+ default:
+ G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+ break;
+ }
+}
+
+static void
+ufo_med_mad_reject_task_finalize (GObject *object)
+{
+ // Releasing the resources held in the private structure...
+ UfoMedMadRejectTaskPrivate *priv = UFO_MED_MAD_REJECT_TASK_GET_PRIVATE (object);
+
+ if ( priv->in0 )
+ g_object_unref(priv->in0);
+ if ( priv->in1 )
+ g_object_unref(priv->in1);
+ if ( priv->in2 )
+ g_object_unref(priv->in2);
+
+ if ( priv->kernel )
+ UFO_RESOURCES_CHECK_CLERR(clReleaseKernel(priv->kernel));
+
+ G_OBJECT_CLASS (ufo_med_mad_reject_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+ iface->setup = ufo_med_mad_reject_task_setup;
+ iface->get_num_inputs = ufo_med_mad_reject_task_get_num_inputs;
+ iface->get_num_dimensions = ufo_med_mad_reject_task_get_num_dimensions;
+ iface->get_mode = ufo_med_mad_reject_task_get_mode;
+ iface->get_requisition = ufo_med_mad_reject_task_get_requisition;
+ iface->process = ufo_med_mad_reject_task_process;
+ iface->generate = ufo_med_mad_reject_task_generate; // Added to provide gobject with that additional method (since REDUCTOR).
+}
+
+static void
+ufo_med_mad_reject_task_class_init (UfoMedMadRejectTaskClass *klass)
+{
+ GObjectClass *oclass = G_OBJECT_CLASS (klass);
+
+ oclass->set_property = ufo_med_mad_reject_task_set_property;
+ oclass->get_property = ufo_med_mad_reject_task_get_property;
+ oclass->finalize = ufo_med_mad_reject_task_finalize;
+
+ properties[PROP_THRESHOLD] =
+ g_param_spec_float ("threshold",
+ "Rejection threshold",
+ "Pixel value replaced by median when more than threshold away from it (in term of MAD).",
+ 0.0f, G_MAXFLOAT, 3.0f,
+ G_PARAM_READWRITE);
+
+ for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+ g_object_class_install_property (oclass, i, properties[i]);
+
+ g_type_class_add_private (oclass, sizeof(UfoMedMadRejectTaskPrivate));
+}
+
+static void
+ufo_med_mad_reject_task_init(UfoMedMadRejectTask *self)
+{
+ self->priv = UFO_MED_MAD_REJECT_TASK_GET_PRIVATE(self);
+ self->priv->threshold = 3.0f;
+
+ self->priv->in0 = NULL;
+ self->priv->in1 = NULL;
+ self->priv->in2 = NULL;
+
+ self->priv->proc_was_called = FALSE;
+ self->priv->gene_shortcut = FALSE;
+ self->priv->proc_call_number = 0;
+ self->priv->gene_call_number = 0;
+}
diff --git a/contrib/sxc/src/ufo-med-mad-reject-task.h b/contrib/sxc/src/ufo-med-mad-reject-task.h
new file mode 100644
index 0000000..6de7bd8
--- /dev/null
+++ b/contrib/sxc/src/ufo-med-mad-reject-task.h
@@ -0,0 +1,54 @@
+/*
+ * Replacing pixels by local median value when detected as outliers based on MAD
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+#ifndef __UFO_MED_MAD_REJECT_TASK_H
+#define __UFO_MED_MAD_REJECT_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_MED_MAD_REJECT_TASK (ufo_med_mad_reject_task_get_type())
+#define UFO_MED_MAD_REJECT_TASK(obj) (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_MED_MAD_REJECT_TASK, UfoMedMadRejectTask))
+#define UFO_IS_MED_MAD_REJECT_TASK(obj) (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_MED_MAD_REJECT_TASK))
+#define UFO_MED_MAD_REJECT_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_MED_MAD_REJECT_TASK, UfoMedMadRejectTaskClass))
+#define UFO_IS_MED_MAD_REJECT_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_MED_MAD_REJECT_TASK))
+#define UFO_MED_MAD_REJECT_TASK_GET_CLASS(obj) (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_MED_MAD_REJECT_TASK, UfoMedMadRejectTaskClass))
+
+typedef struct _UfoMedMadRejectTask UfoMedMadRejectTask;
+typedef struct _UfoMedMadRejectTaskClass UfoMedMadRejectTaskClass;
+typedef struct _UfoMedMadRejectTaskPrivate UfoMedMadRejectTaskPrivate;
+
+struct _UfoMedMadRejectTask {
+ UfoTaskNode parent_instance;
+
+ UfoMedMadRejectTaskPrivate *priv;
+};
+
+struct _UfoMedMadRejectTaskClass {
+ UfoTaskNodeClass parent_class;
+};
+
+UfoNode *ufo_med_mad_reject_task_new (void);
+GType ufo_med_mad_reject_task_get_type (void);
+
+G_END_DECLS
+
+#endif
diff --git a/contrib/sxc/src/ufo-ocl-1liner-task.c b/contrib/sxc/src/ufo-ocl-1liner-task.c
new file mode 100644
index 0000000..e641e56
--- /dev/null
+++ b/contrib/sxc/src/ufo-ocl-1liner-task.c
@@ -0,0 +1,294 @@
+/*
+ * Instant compiling a one liner OpenCL filter
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include <stdio.h> // for asprintf at least.
+
+#include "ufo-ocl-1liner-task.h"
+
+
+struct _UfoOCL1LinerTaskPrivate {
+ gchar * one_line;
+ cl_kernel kernel;
+ guint num_inputs;
+};
+
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoOCL1LinerTask, ufo_ocl_1liner_task, UFO_TYPE_TASK_NODE,
+ G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+ ufo_task_interface_init))
+
+#define UFO_OCL_1LINER_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_OCL_1LINER_TASK, UfoOCL1LinerTaskPrivate))
+
+enum {
+ PROP_0,
+ PROP_ONE_LINE,
+ PROP_NUM_INPUTS,
+ N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_ocl_1liner_task_new (void)
+{
+ return UFO_NODE (g_object_new (UFO_TYPE_OCL_1LINER_TASK, NULL));
+}
+
+static void
+ufo_ocl_1liner_task_setup (UfoTask *task,
+ UfoResources *resources,
+ GError **error)
+{
+ UfoOCL1LinerTaskPrivate *priv;
+ priv = UFO_OCL_1LINER_TASK_GET_PRIVATE (task);
+
+ /* Preparing the kernel source to be compiled */
+ gchar * kernel_skel = NULL;
+ char * kernel_src = NULL;
+
+ const gchar * skel_filename = "ocl-1liner-skel.cl";
+
+ kernel_skel = ufo_resources_get_kernel_source (resources, skel_filename, error);
+
+ char skel_in[1024];
+ int one_in_size;
+ char * skel_in_next = skel_in;
+
+ for ( guint i=0; i != priv->num_inputs; ++i ) {
+ one_in_size = snprintf(skel_in, 1024 - (skel_in_next - skel_in), "__global float *in_%d,\n", i);
+ if ( 0 >= one_in_size )
+ goto exit;
+ skel_in_next += one_in_size;
+ }
+
+ asprintf(&kernel_src, kernel_skel, skel_in, priv->one_line);
+ /* Done with the preparation of the OpenCL sources. */
+ fprintf(stdout, "Current version of the one-liner OpenCL source code :\n%s\n", kernel_src);
+
+ /* Copiling the kernel now : */
+ priv->kernel = ufo_resources_get_kernel_from_source(resources,
+ kernel_src,
+ "ocl_1liner",
+ error);
+ /* Done compiling sources into a kernel*/
+
+ exit:
+ /* Releasing resources no longer used */
+ g_free (kernel_skel);
+ g_free (kernel_src);
+}
+
+static void
+ufo_ocl_1liner_task_get_requisition (UfoTask *task,
+ UfoBuffer **inputs,
+ UfoRequisition *requisition)
+{
+ // In the current version of the kernel all the inputs are supposed to have the same dimensions
+ // Or more precisely the output has the same dimension as first input (indexed 0) and one work-item
+ // is run for each pixel of input 0 (hence for each pixel of the output).
+ ufo_buffer_get_requisition (inputs[0], requisition);
+}
+
+static guint
+ufo_ocl_1liner_task_get_num_inputs (UfoTask *task)
+{
+ UfoOCL1LinerTaskPrivate *priv;
+ priv = UFO_OCL_1LINER_TASK_GET_PRIVATE (task);
+
+ return priv->num_inputs;
+}
+
+static guint
+ufo_ocl_1liner_task_get_num_dimensions (UfoTask *task,
+ guint input)
+{
+ return 2;
+}
+
+static UfoTaskMode
+ufo_ocl_1liner_task_get_mode (UfoTask *task)
+{
+ // Running as a single image stream processing and using OpenCL for GPU computation
+ return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU;
+}
+
+static gboolean
+ufo_ocl_1liner_task_process (UfoTask *task,
+ UfoBuffer **inputs,
+ UfoBuffer *output,
+ UfoRequisition *requisition)
+{
+ UfoOCL1LinerTaskPrivate *priv;
+ priv = UFO_OCL_1LINER_TASK_GET_PRIVATE (task);
+
+ UfoGpuNode *node;
+ UfoProfiler *profiler;
+ cl_command_queue cmd_queue;
+ cl_mem in_mem;
+ cl_mem out_mem;
+
+ node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
+ cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+ profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task));
+
+ // Taking care of setting the inputs :
+ for ( guint i=0; i != priv->num_inputs; ++i ) {
+ in_mem = ufo_buffer_get_device_array (inputs[i], cmd_queue);
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, i, sizeof (cl_mem), &in_mem));
+ }
+
+ // Taking care of setting the output :
+ out_mem = ufo_buffer_get_device_array (output, cmd_queue);
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, priv->num_inputs, sizeof (cl_mem), &out_mem));
+
+ // Finally launching the kernel :
+ ufo_profiler_call (profiler, cmd_queue, priv->kernel, 2, requisition->dims, NULL);
+
+ // Done, returning (TRUE):
+ return TRUE;
+}
+
+
+static void
+ufo_ocl_1liner_task_set_property (GObject *object,
+ guint property_id,
+ const GValue *value,
+ GParamSpec *pspec)
+{
+ UfoOCL1LinerTaskPrivate *priv = UFO_OCL_1LINER_TASK_GET_PRIVATE (object);
+
+ switch (property_id) {
+ case PROP_ONE_LINE:
+ g_free (priv->one_line);
+ priv->one_line = g_value_dup_string (value);
+ break;
+ case PROP_NUM_INPUTS:
+ priv->num_inputs = g_value_get_uint (value);
+ break;
+ default:
+ G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+ break;
+ }
+}
+
+static void
+ufo_ocl_1liner_task_get_property (GObject *object,
+ guint property_id,
+ GValue *value,
+ GParamSpec *pspec)
+{
+ UfoOCL1LinerTaskPrivate *priv = UFO_OCL_1LINER_TASK_GET_PRIVATE (object);
+
+ switch (property_id) {
+ case PROP_ONE_LINE:
+ g_value_set_string (value, priv->one_line);
+ break;
+ case PROP_NUM_INPUTS:
+ g_value_set_uint (value, priv->num_inputs);
+ break;
+ default:
+ G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+ break;
+ }
+}
+
+static void
+ufo_ocl_1liner_task_finalize (GObject *object)
+{
+ UfoOCL1LinerTaskPrivate *priv = UFO_OCL_1LINER_TASK_GET_PRIVATE (object);
+ g_free (priv->one_line);
+
+ if ( priv->kernel )
+ UFO_RESOURCES_CHECK_CLERR(clReleaseKernel(priv->kernel));
+
+ G_OBJECT_CLASS (ufo_ocl_1liner_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+ iface->setup = ufo_ocl_1liner_task_setup;
+ iface->get_num_inputs = ufo_ocl_1liner_task_get_num_inputs;
+ iface->get_num_dimensions = ufo_ocl_1liner_task_get_num_dimensions;
+ iface->get_mode = ufo_ocl_1liner_task_get_mode;
+ iface->get_requisition = ufo_ocl_1liner_task_get_requisition;
+ iface->process = ufo_ocl_1liner_task_process;
+}
+
+static void
+ufo_ocl_1liner_task_class_init (UfoOCL1LinerTaskClass *klass)
+{
+ GObjectClass *oclass = G_OBJECT_CLASS (klass);
+
+ oclass->set_property = ufo_ocl_1liner_task_set_property;
+ oclass->get_property = ufo_ocl_1liner_task_get_property;
+ oclass->finalize = ufo_ocl_1liner_task_finalize;
+
+ properties[PROP_ONE_LINE] =
+ g_param_spec_string ("one-line",
+ "The one line C/OpenCL computation to perform",
+ "* in0 .. inN are input array(s).\n"
+ "* out is 1D output array.\n"
+ "\n"
+ "Those can be addressed using px_index which is the /current/ pixel for the workitem.\n"
+ "\n"
+ "One can also access a /random/ pixel value using the macro : IMG_VAL which takes three\n"
+ "arguments : the (x, y) pixel coordinate, and the array pointer for the image.\n"
+ "NB : this macro acn be used also for lvalue (but it is risky that a work-item\n"
+ "modifies a pixel value in =out= that is NOT the one indexed by px_index.\n"
+ "\n"
+ "Examples :\n"
+ "* using the OpenCL sqrt function :\n"
+ "out[px_index] = sqrt(in0[px_index])"
+ "* Using the other adressing methods and the image size to shift one pixel to the left\n",
+ "out[px_index] = IMG_VAL((x<(sizeX-1))?x+1:(sizeX-1),y,in0)",
+ G_PARAM_READWRITE);
+
+ properties[PROP_NUM_INPUTS] =
+ g_param_spec_uint ("num-inputs",
+ "Number of input streams.",
+ "Number of input streams, which will be labelled in0, in1 ... in(n-1).",
+ 0, G_MAXUINT, 0,
+ G_PARAM_READWRITE);
+
+
+ for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+ g_object_class_install_property (oclass, i, properties[i]);
+
+ g_type_class_add_private (oclass, sizeof(UfoOCL1LinerTaskPrivate));
+}
+
+static void
+ufo_ocl_1liner_task_init(UfoOCL1LinerTask *self)
+{
+ self->priv = UFO_OCL_1LINER_TASK_GET_PRIVATE(self);
+ self->priv->one_line = NULL;
+ self->priv->kernel = NULL;
+ self->priv->num_inputs=1;
+}
diff --git a/contrib/sxc/src/ufo-ocl-1liner-task.h b/contrib/sxc/src/ufo-ocl-1liner-task.h
new file mode 100644
index 0000000..117eb53
--- /dev/null
+++ b/contrib/sxc/src/ufo-ocl-1liner-task.h
@@ -0,0 +1,56 @@
+/*
+ * Instant compiling a one liner OpenCL filter
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+
+#ifndef __UFO_OCL_1LINER_TASK_H
+#define __UFO_OCL_1LINER_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_OCL_1LINER_TASK (ufo_ocl_1liner_task_get_type())
+#define UFO_OCL_1LINER_TASK(obj) (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_OCL_1LINER_TASK, UfoOCL1LinerTask))
+#define UFO_IS_OCL_1LINER_TASK(obj) (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_OCL_1LINER_TASK))
+#define UFO_OCL_1LINER_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_OCL_1LINER_TASK, UfoOCL1LinerTaskClass))
+#define UFO_IS_OCL_1LINER_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_OCL_1LINER_TASK))
+#define UFO_OCL_1LINER_TASK_GET_CLASS(obj) (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_OCL_1LINER_TASK, UfoOCL1LinerTaskClass))
+
+typedef struct _UfoOCL1LinerTask UfoOCL1LinerTask;
+typedef struct _UfoOCL1LinerTaskClass UfoOCL1LinerTaskClass;
+typedef struct _UfoOCL1LinerTaskPrivate UfoOCL1LinerTaskPrivate;
+
+struct _UfoOCL1LinerTask {
+ UfoTaskNode parent_instance;
+
+ UfoOCL1LinerTaskPrivate *priv;
+};
+
+struct _UfoOCL1LinerTaskClass {
+ UfoTaskNodeClass parent_class;
+};
+
+UfoNode *ufo_ocl_1liner_task_new (void);
+GType ufo_ocl_1liner_task_get_type (void);
+
+G_END_DECLS
+
+#endif
diff --git a/contrib/sxc/src/ufo-stat-monitor-task.c b/contrib/sxc/src/ufo-stat-monitor-task.c
new file mode 100644
index 0000000..187549e
--- /dev/null
+++ b/contrib/sxc/src/ufo-stat-monitor-task.c
@@ -0,0 +1,581 @@
+/*
+ * Gathering statistics on a image stream, copying input to output
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+// During development, typical testing is :
+// UFO_DEVICES=1 ufo-launch dummy-data width=1024 height=1024 depth=1 init=12.3 number=16 ! monitor print=4 ! stat-monitor print=4 ! null download=TRUE
+
+#include "ufo-stat-monitor-task.h"
+#include "ufo-sxc-common.h"
+#include <stdio.h>
+#include <string.h>
+#include <math.h>
+
+struct _UfoStatMonitorTaskPrivate {
+ FILE * stat_file;
+ gchar * stat_fn;
+ gboolean be_quiet;
+ gboolean node_has_fp64;
+ cl_kernel kernel;
+ cl_kernel kernel_final;
+ gsize im_index;
+ guint n_items;
+ cl_ulong max_local_mem;
+ size_t local_scratch_size;
+ cl_uint wg_size;
+ cl_uint wg_num;
+ cl_mem stat_out_buff; // The buffer used by the kernel to output its results
+ cl_mem stat_out_red; // The buffer used by the final reduction kernel
+};
+
+static void ufo_task_interface_init (UfoTaskIface *iface);
+
+G_DEFINE_TYPE_WITH_CODE (UfoStatMonitorTask, ufo_stat_monitor_task, UFO_TYPE_TASK_NODE,
+ G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
+ ufo_task_interface_init))
+
+#define UFO_STAT_MONITOR_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_STAT_MONITOR_TASK, UfoStatMonitorTaskPrivate))
+
+enum {
+ PROP_0,
+ PROP_NUM_ITEMS,
+ PROP_STAT_FN,
+ PROP_QUIET,
+ N_PROPERTIES
+};
+
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
+
+UfoNode *
+ufo_stat_monitor_task_new (void)
+{
+ return UFO_NODE (g_object_new (UFO_TYPE_STAT_MONITOR_TASK, NULL));
+}
+
+static void
+ufo_stat_monitor_task_setup (UfoTask *task,
+ UfoResources *resources,
+ GError **error)
+{
+ UfoStatMonitorTaskPrivate *priv;
+ UfoGpuNode *node;
+ cl_command_queue cmd_queue;
+ cl_context context_cl;
+ cl_device_id dev_cl;
+
+ cl_uint num_cu;
+ size_t max_wgs, max_wis[3];
+ size_t ker_pref_wgs;
+
+ node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
+ cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+ UFO_RESOURCES_CHECK_CLERR(clGetCommandQueueInfo(cmd_queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &dev_cl, NULL)); // ,
+ UFO_RESOURCES_CHECK_CLERR(clGetCommandQueueInfo(cmd_queue, CL_QUEUE_CONTEXT, sizeof(cl_context), &context_cl, NULL)); // ,
+
+ priv = UFO_STAT_MONITOR_TASK_GET_PRIVATE (task);
+
+ priv->node_has_fp64 = device_has_extension(node, "cl_khr_fp64");
+
+ // Loading the OpenCL kernel :
+ if ( priv->node_has_fp64 ) {
+ priv->kernel = ufo_resources_get_kernel (resources, "stat-monitor.cl", "stat_monitor_f64", error);
+ if (priv->kernel != NULL)
+ UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel));
+ priv->kernel_final = ufo_resources_get_kernel (resources, "stat-monitor.cl", "stat_monitor_f64_fin", error);
+ if (priv->kernel_final != NULL)
+ UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel_final));
+ }
+ else {
+ priv->kernel = ufo_resources_get_kernel (resources, "stat-monitor.cl", "stat_monitor_f32", error);
+ if (priv->kernel != NULL)
+ UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel));
+ priv->kernel_final = ufo_resources_get_kernel (resources, "stat-monitor.cl", "stat_monitor_f32_fin", error);
+ if (priv->kernel_final != NULL)
+ UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel_final));
+ }
+
+ // Getting the CL_DEVICE_LOCAL_MEM_SIZE of the device :
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &(priv->max_local_mem), NULL));
+
+ // CL_DEVICE_MAX_COMPUTE_UNITS -> cl_uint
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &num_cu, NULL));
+ // CL_DEVICE_MAX_WORK_GROUP_SIZE -> size_t
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &max_wgs, NULL));
+ // CL_DEVICE_MAX_WORK_ITEM_SIZES -> size_t[3]
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_wis), max_wis, NULL));
+
+ // CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE -> size_t
+ UFO_RESOURCES_CHECK_CLERR (clGetKernelWorkGroupInfo (priv->kernel, dev_cl, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(size_t), &ker_pref_wgs, NULL));
+
+ priv->wg_num = num_cu<<2; // 4 work-groups per comput unit.
+ if ( priv->wg_num > max_wis[0] ) { // We have to reduce the number of workgroup so that
+ // last reduction can be done in a single workgroup (hence 1-value per workroup should
+ // eventually fit within a single workgroup.
+ priv->wg_num = max_wis[0];
+ }
+ priv->wg_size = ker_pref_wgs;
+ if ( priv->wg_size > max_wis[0] ) {
+ priv->wg_size = max_wis[0];
+ }
+ if ( priv->wg_size < priv->wg_num ) { // To ensure the final reduction step.
+ priv->wg_num = priv->wg_size;
+ }
+
+ priv->local_scratch_size = (size_t)(priv->max_local_mem >>1 ); // dividing by 2 the total local memory.
+ if ( priv->node_has_fp64 ) {
+ // Each workgroup needs 4 items * 8 bytes per work-item, that is (wg_size << 5) B of local memory
+ if ( priv->local_scratch_size > (priv->wg_size << 5) ) {
+ priv->local_scratch_size = (priv->wg_size << 5); // Only request the necessary amount of memory
+ }
+ else {
+ priv->wg_size = (priv->local_scratch_size) >> 5; // Reducing the number of items in a work-group to fit local memory constraints
+ }
+ }
+ else {
+ // Each workgroup needs 4 items * 4 bytes per work-item, that is (wg_size << 4) B of local memory
+ if ( priv->local_scratch_size > (priv->wg_size << 4) ) {
+ priv->local_scratch_size = (priv->wg_size << 4); // Only request the necessary amount of memory
+ }
+ else {
+ priv->wg_size = (priv->local_scratch_size) >> 4; // Reducing the number of items in a work-group to fit local memory constraints
+ }
+ }
+#warning Still have to check that the values here are properly set.
+ // In particular taking into account local memory usage per work-item and work-group to optimsed these values.
+
+ priv->im_index = 0;
+
+ // Opening (if required) the statistic file
+ if ( strcmp("-", priv->stat_fn) ) {
+ priv->stat_file = fopen(priv->stat_fn, "a");
+ fprintf(stdout, "stat-monitor will outputs its results to file '%s'\n", priv->stat_fn);
+ fprintf(priv->stat_file, "# index min max mean var\n");
+ }
+ else {
+ priv->stat_file = stdout;
+ fprintf(stdout, "stat-monitor will outputs its results to stdout\n");
+ }
+
+ // Allocating once for all the output buffer that will be used for statistcis output.
+ cl_int err_code;
+ if ( priv->node_has_fp64 ) {
+ // min, max, mean, sd (one 4-tupple per work-group)
+ priv->stat_out_buff = clCreateBuffer (context_cl, CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, priv->wg_num << 5, NULL, &err_code);
+ UFO_RESOURCES_CHECK_CLERR(err_code);
+ // min, max, mean, sd (one 4-tupple once only)
+ priv->stat_out_red = clCreateBuffer (context_cl, CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY, 1 << 5, NULL, &err_code);
+ UFO_RESOURCES_CHECK_CLERR(err_code);
+ }
+ else {
+ // min, max, mean, sd (one 4-tupple per work-group)
+ priv->stat_out_buff = clCreateBuffer (context_cl, CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, priv->wg_num << 4, NULL, &err_code);
+ UFO_RESOURCES_CHECK_CLERR(err_code);
+ // min, max, mean, sd (one 4-tupple once only)
+ priv->stat_out_red = clCreateBuffer (context_cl, CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY, 1 << 4, NULL, &err_code);
+ UFO_RESOURCES_CHECK_CLERR(err_code);
+ }
+
+ fprintf(stdout, "%s (%s.%d) : instrospected and computed values are as follow :\n", __func__, __FILE__, __LINE__);
+ fprintf(stdout, " * Will print standard monitor message : %s\n", (priv->be_quiet) ? "no" : "yes");
+ fprintf(stdout, " * Will output statistics to %s\n", strcmp("-", priv->stat_fn) ? priv->stat_fn : "stdout");
+ fprintf(stdout, " * OpenCL Device %s fp64/double precision support (cl_khr_fp64)\n", (priv->node_has_fp64) ? "has" : "does not have");
+ fprintf(stdout, " * Maximum local memory is %lluB\n", priv->max_local_mem);
+ fprintf(stdout, " * Number of compute units is %d\n", num_cu);
+ fprintf(stdout, " * Maximum work group size is %zu\n", max_wgs);
+ fprintf(stdout, " * Maximum work items (in WG) [0] is %zu\n", max_wis[0]);
+ fprintf(stdout, " * Preferred multiple of work-group size is %zu\n", ker_pref_wgs);
+ fprintf(stdout, " * Tyring %d work-item per work-group\n", priv->wg_size);
+ fprintf(stdout, " * Tyring %d work-groups\n", priv->wg_num);
+ fprintf(stdout, " * Providing %zuB of local memory to each work-group\n", priv->local_scratch_size);
+}
+
+static void
+ufo_stat_monitor_task_get_requisition (UfoTask *task,
+ UfoBuffer **inputs,
+ UfoRequisition *requisition)
+{
+ // In the current version the statistics are NEVER the output of the filter.
+ // Indeed is behaving as a /pass-through/ filter, doing nothing to the image
+ ufo_buffer_get_requisition (inputs[0], requisition);
+
+ // Rather than the default :
+ // requisition->n_dims = 0;
+}
+
+static guint
+ufo_stat_monitor_task_get_num_inputs (UfoTask *task)
+{
+ return 1;
+}
+
+static guint
+ufo_stat_monitor_task_get_num_dimensions (UfoTask *task,
+ guint input)
+{
+ return 2;
+}
+
+static UfoTaskMode
+ufo_stat_monitor_task_get_mode (UfoTask *task)
+{
+ // We are still needing the GPU (OpenCL device indeed) to perform the statistics computation
+ return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU;
+}
+
+// Copied over from the monitor task, as we are trying to mimick it (plus statistics gathering/printing)
+static gchar *
+join_list (GList *list, const gchar *sep)
+{
+ gchar **array;
+ GList *it;
+ gchar *result;
+ guint i = 0;
+
+ array = g_new0 (gchar *, g_list_length (list) + 1);
+
+ g_list_for (list, it)
+ array[i++] = it->data;
+
+ result = g_strjoinv (sep, array);
+ g_free (array);
+ return result;
+}
+
+static gboolean
+ufo_stat_monitor_task_process (UfoTask *task,
+ UfoBuffer **inputs,
+ UfoBuffer *output,
+ UfoRequisition *requisition)
+{
+ UfoStatMonitorTaskPrivate *priv;
+ priv = UFO_STAT_MONITOR_TASK_GET_PRIVATE (task);
+
+ UfoGpuNode *node;
+ UfoProfiler *profiler;
+ cl_command_queue cmd_queue;
+ cl_mem in_mem;
+ cl_uint img_size;
+ UfoRequisition img_req;
+
+ UfoBufferLocation location;
+ GList *keys;
+ GList *sizes;
+ gchar *keystring;
+ gchar *dimstring;
+
+ // Getting information from the buffer, before computing statistics
+ location = ufo_buffer_get_location (inputs[0]);
+ keys = ufo_buffer_get_metadata_keys (inputs[0]);
+ sizes = NULL;
+
+ // Launching the kernel first, so that it has a bit of extra time while CPU is running
+ node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
+ cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+
+ in_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue);
+ ufo_buffer_get_requisition(inputs[0], &img_req);
+ switch (img_req.n_dims) {
+ case 1:
+ img_size = img_req.dims[0];
+ break;
+ case 2:
+ img_size = img_req.dims[0] * img_req.dims[1];
+ break;
+ case 3:
+ img_size = img_req.dims[0] * img_req.dims[1] * img_req.dims[2];
+ break;
+ default:
+ img_size=0;
+ break;
+ }
+
+ // fprintf(stdout, "img_size is : %d\n", img_size);
+
+ // Now we can start the statistics :
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof (cl_mem), &in_mem));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof (cl_mem), &(priv->stat_out_buff)));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof (cl_uint), &img_size));
+ // The size of the provided local memory is computed /once for all/ during the task's setup call.
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 3, priv->local_scratch_size, NULL));
+
+ gsize total_wi = priv->wg_num * priv->wg_size;
+ gsize wg_size = (gsize)(priv->wg_size);
+
+ total_wi = (total_wi < (gsize)img_size) ? total_wi : (gsize)img_size;
+
+ profiler = ufo_task_node_get_profiler (UFO_TASK_NODE (task));
+ // First reduction step :
+ // fprintf(stdout, "Reduction 1st stage :num pixels %u, num items %lu, workgroup size %lu\n", img_size, total_wi, wg_size);
+ ufo_profiler_call (profiler, cmd_queue, priv->kernel, 1, &total_wi, &wg_size);
+
+#warning Further reduction is required, CPU or GPU ...
+ // At this time, we need to have a second kernel to further reduce the results of the previous results that where
+ // produced at the rate of one per work-group.
+ if ( TRUE ) { // OpenCL version of this last stage of reduction
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel_final, 0, sizeof (cl_mem), &(priv->stat_out_buff)));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel_final, 1, sizeof (cl_mem), &(priv->stat_out_red)));
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel_final, 2, sizeof (cl_uint), &(priv->wg_num)));
+ if ( priv->node_has_fp64 ) {
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel_final, 3, sizeof(cl_double)*4*priv->wg_num, NULL));
+ }
+ else {
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel_final, 3, sizeof(cl_float)*4*priv->wg_num, NULL));
+ }
+
+ cl_uint true_wg_num = ((priv->wg_num * priv->wg_size) < (gsize)img_size) ? priv->wg_num : (1 + (total_wi-1)/wg_size) ;
+
+ total_wi = (true_wg_num & 0x1) + (true_wg_num >> 1);
+ // Second reduction step :
+ // fprintf(stdout, "Reduction 2nd stage : num of 1st stage WG %u, num items %lu, workgroup size %lu\n", true_wg_num, total_wi, total_wi);
+ ufo_profiler_call(profiler, cmd_queue, priv->kernel_final, 1, &total_wi, &total_wi);
+
+ if ( priv->node_has_fp64 ) {
+ double stat_res[4];
+ double img_size_f = (double)img_size;
+
+ UFO_RESOURCES_CHECK_CLERR (clEnqueueReadBuffer(cmd_queue, priv->stat_out_red, CL_TRUE, 0, 4<<3, stat_res, 0, NULL, NULL));
+
+ stat_res[2] = stat_res[2] / img_size_f;
+ stat_res[3] = (stat_res[3] - img_size_f * stat_res[2] * stat_res[2]) / (img_size_f - 1.0);
+
+ fprintf (priv->stat_file, "%zu %le %le %le %le\n", priv->im_index, stat_res[0], stat_res[1], stat_res[2], stat_res[3]);
+ }
+ else {
+ float stat_res[4];
+ float img_size_f = (float)img_size;
+
+ UFO_RESOURCES_CHECK_CLERR (clEnqueueReadBuffer(cmd_queue, priv->stat_out_red, CL_TRUE, 0, 4<<2, stat_res, 0, NULL, NULL));
+
+ stat_res[2] = stat_res[2] / img_size_f;
+ stat_res[3] = (stat_res[3] - img_size_f * stat_res[2] * stat_res[2]) / (img_size_f - 1.0f);
+
+ fprintf (priv->stat_file, "%zu %e %e %e %e\n", priv->im_index, stat_res[0], stat_res[1], stat_res[2], stat_res[3]);
+ }
+ // #error Should continue adding fp32/fp64 branch here ...
+ }
+ else { // CPU version of the last stage of the reduction
+ /*
+ // This is «old code» to perform the last reduction step on the CPU… no more working.
+ float * stat_res = ufo_buffer_get_host_array(priv->stat_out_buff, cmd_queue);
+ float img_size_f = (float)img_size;
+
+ for ( size_t i=0; priv->wg_num != i; i += 1 ) {
+ stat_res[0] = (stat_res[0] < stat_res[ (i<<2)]) ? stat_res[0] : stat_res[ (i<<2)];
+ stat_res[1] = (stat_res[1] > stat_res[1+(i<<2)]) ? stat_res[1] : stat_res[1+(i<<2)];
+ stat_res[2] += stat_res[2+(i<<2)];
+ stat_res[3] += stat_res[3+(i<<2)];
+ }
+ stat_res[2] = stat_res[2] / img_size_f;
+ stat_res[3] = (stat_res[3] - img_size_f * stat_res[2] * stat_res[2]) / (img_size_f - 1.0f);
+
+ fprintf (priv->stat_file, "%e %e %e %e\n", stat_res[0], stat_res[1], stat_res[2], stat_res[3]);
+ */
+ }
+ ++(priv->im_index);
+
+ if ( ! priv->be_quiet ) {
+ for (guint i = 0; i < requisition->n_dims; i++)
+ sizes = g_list_append (sizes, g_strdup_printf ("%zu", requisition->dims[i]));
+
+ dimstring = join_list (sizes, " ");
+ keystring = join_list (keys, ", ");
+
+ g_print ("stat-monitor: dims=[%s] keys=[%s] location=", dimstring, keystring);
+
+ switch (location) {
+ case UFO_BUFFER_LOCATION_HOST:
+ g_print ("host");
+ break;
+ case UFO_BUFFER_LOCATION_DEVICE:
+ g_print ("device");
+ break;
+ case UFO_BUFFER_LOCATION_DEVICE_IMAGE:
+ g_print ("image");
+ break;
+ case UFO_BUFFER_LOCATION_INVALID:
+ g_print ("invalid");
+ break;
+ }
+
+ g_print ("\n");
+
+ g_free (dimstring);
+ g_free (keystring);
+ g_list_free (keys);
+ g_list_free_full (sizes, (GDestroyNotify) g_free);
+ }
+
+ if (priv->n_items > 0) {
+ guint32 *data;
+ gfloat *data_f32;
+
+ data = (guint32 *) ufo_buffer_get_host_array (inputs[0], NULL);
+ data_f32 = (gfloat *) ufo_buffer_get_host_array (inputs[0], NULL);
+
+ g_print (" ");
+
+ for (guint i = 0; i < priv->n_items; i++) {
+ // g_print ("0x%08x ", data[i]);
+ g_print ("%e ", data_f32[i]);
+
+ if ((i != 0) && (((i + 1) % 8) == 0))
+ g_print ("\n ");
+ }
+
+ if ((priv->n_items % 8) != 0)
+ g_print ("\n");
+ }
+
+ ufo_buffer_copy (inputs[0], output);
+
+ return TRUE;
+}
+
+static void
+ufo_stat_monitor_task_set_property (GObject *object,
+ guint property_id,
+ const GValue *value,
+ GParamSpec *pspec)
+{
+ UfoStatMonitorTaskPrivate *priv = UFO_STAT_MONITOR_TASK_GET_PRIVATE (object);
+
+ switch (property_id) {
+ case PROP_NUM_ITEMS:
+ priv->n_items = g_value_get_uint(value);
+ break;
+ case PROP_STAT_FN:
+ g_free (priv->stat_fn);
+ priv->stat_fn = g_value_dup_string (value);
+ break;
+ case PROP_QUIET:
+ priv->be_quiet = g_value_get_boolean (value);
+ break;
+ default:
+ G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+ break;
+ }
+}
+
+static void
+ufo_stat_monitor_task_get_property (GObject *object,
+ guint property_id,
+ GValue *value,
+ GParamSpec *pspec)
+{
+ UfoStatMonitorTaskPrivate *priv = UFO_STAT_MONITOR_TASK_GET_PRIVATE (object);
+
+ switch (property_id) {
+ case PROP_NUM_ITEMS:
+ g_value_set_uint (value, priv->n_items);
+ break;
+ case PROP_STAT_FN:
+ g_value_set_string (value, priv->stat_fn);
+ break;
+ default:
+ G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
+ break;
+ }
+}
+
+static void
+ufo_stat_monitor_task_finalize (GObject *object)
+{
+ UfoStatMonitorTaskPrivate *priv = UFO_STAT_MONITOR_TASK_GET_PRIVATE (object);
+
+ if ( stdout != priv->stat_file ) {
+ fclose(priv->stat_file);
+ priv->stat_file = NULL;
+ }
+
+ g_free (priv->stat_fn);
+ priv->stat_fn = NULL;
+
+ if ( priv->kernel )
+ UFO_RESOURCES_CHECK_CLERR(clReleaseKernel(priv->kernel));
+
+ if ( priv->kernel_final )
+ UFO_RESOURCES_CHECK_CLERR(clReleaseKernel(priv->kernel_final));
+
+ if ( priv->stat_out_buff )
+ UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject(priv->stat_out_buff));
+ // g_object_unref(priv->stat_out_buff);
+
+ if ( priv->stat_out_red )
+ UFO_RESOURCES_CHECK_CLERR (clReleaseMemObject(priv->stat_out_red));
+ // g_object_unref(priv->stat_out_red);
+
+ G_OBJECT_CLASS (ufo_stat_monitor_task_parent_class)->finalize (object);
+}
+
+static void
+ufo_task_interface_init (UfoTaskIface *iface)
+{
+ iface->setup = ufo_stat_monitor_task_setup;
+ iface->get_num_inputs = ufo_stat_monitor_task_get_num_inputs;
+ iface->get_num_dimensions = ufo_stat_monitor_task_get_num_dimensions;
+ iface->get_mode = ufo_stat_monitor_task_get_mode;
+ iface->get_requisition = ufo_stat_monitor_task_get_requisition;
+ iface->process = ufo_stat_monitor_task_process;
+}
+
+static void
+ufo_stat_monitor_task_class_init (UfoStatMonitorTaskClass *klass)
+{
+ GObjectClass *oclass = G_OBJECT_CLASS (klass);
+
+ oclass->set_property = ufo_stat_monitor_task_set_property;
+ oclass->get_property = ufo_stat_monitor_task_get_property;
+ oclass->finalize = ufo_stat_monitor_task_finalize;
+
+ properties[PROP_STAT_FN] =
+ g_param_spec_string("filename",
+ "Filename for the statistics output file.",
+ "If provided with a '-' it will output statistcis to standard output of the process",
+ "-",
+ G_PARAM_READWRITE);
+ properties[PROP_QUIET] =
+ g_param_spec_boolean("quiet",
+ "When turned to true, will not print frame monitoring information on stdout",
+ "Defaulting to 'false', that is mimicking the 'monitor' filter",
+ TRUE,
+ G_PARAM_READWRITE);
+ properties[PROP_NUM_ITEMS] =
+ g_param_spec_uint ("print",
+ "Number of items to print",
+ "Number of items to print",
+ 0, G_MAXUINT, 0,
+ G_PARAM_READWRITE);
+
+ for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
+ g_object_class_install_property (oclass, i, properties[i]);
+
+ g_type_class_add_private (oclass, sizeof(UfoStatMonitorTaskPrivate));
+}
+
+static void
+ufo_stat_monitor_task_init(UfoStatMonitorTask *self)
+{
+ self->priv = UFO_STAT_MONITOR_TASK_GET_PRIVATE(self);
+
+ self->priv->stat_file = stdout;
+ self->priv->stat_fn = g_strdup ("-");
+ self->priv->be_quiet = FALSE;
+ self->priv->n_items = 0;
+}
diff --git a/contrib/sxc/src/ufo-stat-monitor-task.h b/contrib/sxc/src/ufo-stat-monitor-task.h
new file mode 100644
index 0000000..19bbbeb
--- /dev/null
+++ b/contrib/sxc/src/ufo-stat-monitor-task.h
@@ -0,0 +1,54 @@
+/*
+ * Gathering statistics on a image stream, copying input to output
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+#ifndef __UFO_STAT_MONITOR_TASK_H
+#define __UFO_STAT_MONITOR_TASK_H
+
+#include <ufo/ufo.h>
+
+G_BEGIN_DECLS
+
+#define UFO_TYPE_STAT_MONITOR_TASK (ufo_stat_monitor_task_get_type())
+#define UFO_STAT_MONITOR_TASK(obj) (G_TYPE_CHECK_INSTANCE_CAST((obj), UFO_TYPE_STAT_MONITOR_TASK, UfoStatMonitorTask))
+#define UFO_IS_STAT_MONITOR_TASK(obj) (G_TYPE_CHECK_INSTANCE_TYPE((obj), UFO_TYPE_STAT_MONITOR_TASK))
+#define UFO_STAT_MONITOR_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_CAST((klass), UFO_TYPE_STAT_MONITOR_TASK, UfoStatMonitorTaskClass))
+#define UFO_IS_STAT_MONITOR_TASK_CLASS(klass) (G_TYPE_CHECK_CLASS_TYPE((klass), UFO_TYPE_STAT_MONITOR_TASK))
+#define UFO_STAT_MONITOR_TASK_GET_CLASS(obj) (G_TYPE_INSTANCE_GET_CLASS((obj), UFO_TYPE_STAT_MONITOR_TASK, UfoStatMonitorTaskClass))
+
+typedef struct _UfoStatMonitorTask UfoStatMonitorTask;
+typedef struct _UfoStatMonitorTaskClass UfoStatMonitorTaskClass;
+typedef struct _UfoStatMonitorTaskPrivate UfoStatMonitorTaskPrivate;
+
+struct _UfoStatMonitorTask {
+ UfoTaskNode parent_instance;
+
+ UfoStatMonitorTaskPrivate *priv;
+};
+
+struct _UfoStatMonitorTaskClass {
+ UfoTaskNodeClass parent_class;
+};
+
+UfoNode *ufo_stat_monitor_task_new (void);
+GType ufo_stat_monitor_task_get_type (void);
+
+G_END_DECLS
+
+#endif // __UFO_STAT_MONITOR_TASK_H
diff --git a/contrib/sxc/src/ufo-sxc-common.c b/contrib/sxc/src/ufo-sxc-common.c
new file mode 100644
index 0000000..a0c3d6f
--- /dev/null
+++ b/contrib/sxc/src/ufo-sxc-common.c
@@ -0,0 +1,158 @@
+/*
+ * Common functions for this set of filters.
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+
+#include "ufo-sxc-common.h"
+
+// Required headers :
+#include <string.h>
+
+
+GValue *
+get_device_info (UfoGpuNode *node, cl_device_info param_name)
+{
+ cl_command_queue cmd_queue;
+ cl_device_id dev_cl;
+
+ cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+ UFO_RESOURCES_CHECK_CLERR(clGetCommandQueueInfo(cmd_queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &dev_cl, NULL));
+
+ cl_uint uint_val;
+ cl_ulong ulong_val;
+ cl_bool bool_val;
+ size_t sizet_val;
+ size_t sizet_x3_val[3];
+ char string_val[2048];
+
+ // Getting ready to output the value/GValue
+ GValue *value;
+ value = g_new0 (GValue, 1);
+ memset (value, 0, sizeof (GValue));
+
+ switch (param_name) {
+ // All the next ones are cl_uint, so sharing the code :
+ case CL_DEVICE_VENDOR_ID:
+ case CL_DEVICE_MAX_COMPUTE_UNITS:
+ case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:
+ case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
+ case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
+ case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
+ case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
+ case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
+ case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
+ case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF:
+ case CL_DEVICE_MAX_CLOCK_FREQUENCY:
+ case CL_DEVICE_ADDRESS_BITS:
+ case CL_DEVICE_MAX_READ_IMAGE_ARGS:
+ case CL_DEVICE_MAX_WRITE_IMAGE_ARGS:
+ case CL_DEVICE_MAX_SAMPLERS:
+ case CL_DEVICE_MEM_BASE_ADDR_ALIGN:
+ case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE:
+ case CL_DEVICE_MAX_CONSTANT_ARGS:
+ case CL_DEVICE_PARTITION_MAX_SUB_DEVICES:
+ case CL_DEVICE_REFERENCE_COUNT:
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, param_name, sizeof(cl_uint), &uint_val, NULL));
+ g_value_init (value, G_TYPE_UINT);
+ g_value_set_uint (value, uint_val);
+ break;
+ //All the next ones are cl_ulong, so sharing the code :
+ case CL_DEVICE_MAX_MEM_ALLOC_SIZE:
+ case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE:
+ case CL_DEVICE_GLOBAL_MEM_SIZE:
+ case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE:
+ case CL_DEVICE_LOCAL_MEM_SIZE:
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, param_name, sizeof(cl_ulong), &ulong_val, NULL));
+ g_value_init (value, G_TYPE_ULONG);
+ g_value_set_ulong (value, ulong_val);
+ break;
+ //All the next ones are cl_bool, so sharing the code :
+ case CL_DEVICE_IMAGE_SUPPORT:
+ case CL_DEVICE_ERROR_CORRECTION_SUPPORT:
+ case CL_DEVICE_HOST_UNIFIED_MEMORY:
+ case CL_DEVICE_ENDIAN_LITTLE:
+ case CL_DEVICE_AVAILABLE:
+ case CL_DEVICE_COMPILER_AVAILABLE:
+ case CL_DEVICE_LINKER_AVAILABLE:
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, param_name, sizeof(cl_bool), &bool_val, NULL));
+ g_value_init (value, G_TYPE_BOOLEAN);
+ g_value_set_boolean (value, bool_val);
+ break;
+ //All the next ones are size_t, so sharing the code :
+ case CL_DEVICE_MAX_WORK_GROUP_SIZE:
+ case CL_DEVICE_IMAGE2D_MAX_WIDTH:
+ case CL_DEVICE_IMAGE2D_MAX_HEIGHT:
+ case CL_DEVICE_IMAGE3D_MAX_WIDTH:
+ case CL_DEVICE_IMAGE3D_MAX_HEIGHT:
+ case CL_DEVICE_IMAGE3D_MAX_DEPTH:
+ case CL_DEVICE_IMAGE_MAX_BUFFER_SIZE:
+ case CL_DEVICE_IMAGE_MAX_ARRAY_SIZE:
+ case CL_DEVICE_MAX_PARAMETER_SIZE:
+ case CL_DEVICE_PROFILING_TIMER_RESOLUTION:
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, param_name, sizeof(size_t), &sizet_val, NULL));
+ g_value_init (value, G_TYPE_ULONG);
+ g_value_set_ulong (value, sizet_val);
+ break;
+ //All the next ones are size_t[], so sharing the code :
+ case CL_DEVICE_MAX_WORK_ITEM_SIZES:
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, param_name, sizeof(sizet_x3_val), sizet_x3_val, NULL));
+ g_free(value);
+ value = g_new0 (GValue, 3);
+ memset (value, 0, 3*sizeof (GValue));
+ g_value_init (value, G_TYPE_ULONG);
+ g_value_set_ulong (value, sizet_x3_val[0]);
+ g_value_init (value+1, G_TYPE_ULONG);
+ g_value_set_ulong (value+1, sizet_x3_val[1]);
+ g_value_init (value+2, G_TYPE_ULONG);
+ g_value_set_ulong (value+2, sizet_x3_val[2]);
+ break;
+ //All the next ones are char[], so sharing the code :
+ case CL_DEVICE_BUILT_IN_KERNELS:
+ case CL_DEVICE_NAME:
+ case CL_DEVICE_VENDOR:
+ case CL_DRIVER_VERSION:
+ case CL_DEVICE_PROFILE:
+ case CL_DEVICE_VERSION:
+ case CL_DEVICE_OPENCL_C_VERSION:
+ case CL_DEVICE_EXTENSIONS:
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, param_name, sizeof(string_val), string_val, NULL));
+ g_value_init (value, G_TYPE_STRING);
+ g_value_set_string (value, string_val);
+ break;
+ }
+ return value;
+}
+
+gboolean
+device_has_extension (UfoGpuNode *node, const char * i_ext_name)
+{
+ cl_command_queue cmd_queue;
+ cl_device_id dev_cl;
+
+ cmd_queue = ufo_gpu_node_get_cmd_queue (node);
+ UFO_RESOURCES_CHECK_CLERR(clGetCommandQueueInfo(cmd_queue, CL_QUEUE_DEVICE, sizeof(cl_device_id), &dev_cl, NULL));
+
+ char exts_val[2048];
+ char *ret_strstr = NULL;
+
+ UFO_RESOURCES_CHECK_CLERR (clGetDeviceInfo (dev_cl, CL_DEVICE_EXTENSIONS, sizeof(exts_val), exts_val, NULL));
+ ret_strstr = strstr(exts_val, i_ext_name);
+ return(NULL != ret_strstr);
+}
+
diff --git a/contrib/sxc/src/ufo-sxc-common.h b/contrib/sxc/src/ufo-sxc-common.h
new file mode 100644
index 0000000..d7291f1
--- /dev/null
+++ b/contrib/sxc/src/ufo-sxc-common.h
@@ -0,0 +1,44 @@
+/*
+ * Common functions for this set of filters.
+ * This file is part of ufo-serge filter set.
+ * Copyright (C) 2016 Serge Cohen
+ *
+ * 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/>.
+ *
+ * Serge Cohen <serge.cohen@synchrotron-soleil.fr>
+ */
+#ifndef __UFO_SXC_COMMON_H
+#define __UFO_SXC_COMMON_H
+
+#include <ufo/ufo.h>
+
+// Using OpenCL :
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+// This one is copied over from ufo_priv.h, which as implied by its name is private to ufo and hence only accessible when building within ufo-filters.
+#ifndef g_list_for
+#define g_list_for(list, it) \
+ for (it = g_list_first (list); \
+ it != NULL; \
+ it = g_list_next (it))
+#endif
+
+GValue * get_device_info (UfoGpuNode *node, cl_device_info param_name);
+gboolean device_has_extension (UfoGpuNode *node, const char * i_ext_name);
+
+#endif