diff options
author | Serge Cohen <serge@chocolatnoir.net> | 2017-01-04 23:17:00 +0100 |
---|---|---|
committer | Matthias Vogelgesang <matthias.vogelgesang@kit.edu> | 2017-01-10 15:43:52 +0100 |
commit | 20e0efa57c172521397438bcf30081f2a1577019 (patch) | |
tree | fa6c4fa2afe2de1735311875a47fa1e960d5ed76 | |
parent | 0e3cca9cb63656c88bb336de42c54bfadc16d827 (diff) | |
download | ufo-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.txt | 3 | ||||
-rw-r--r-- | contrib/CMakeLists.txt | 26 | ||||
-rw-r--r-- | contrib/sxc/CMakeLists.txt | 28 | ||||
-rw-r--r-- | contrib/sxc/COPYING | 674 | ||||
-rw-r--r-- | contrib/sxc/src/CMakeLists.txt | 102 | ||||
-rw-r--r-- | contrib/sxc/src/kernels/med-mad-reject-2d.cl | 95 | ||||
-rw-r--r-- | contrib/sxc/src/kernels/med-mad-reject.cl | 207 | ||||
-rw-r--r-- | contrib/sxc/src/kernels/ocl-1liner-skel.cl | 56 | ||||
-rw-r--r-- | contrib/sxc/src/kernels/stat-monitor.cl | 233 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-med-mad-reject-2d-task.c | 254 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-med-mad-reject-2d-task.h | 55 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-med-mad-reject-task.c | 340 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-med-mad-reject-task.h | 54 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-ocl-1liner-task.c | 294 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-ocl-1liner-task.h | 56 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-stat-monitor-task.c | 581 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-stat-monitor-task.h | 54 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-sxc-common.c | 158 | ||||
-rw-r--r-- | contrib/sxc/src/ufo-sxc-common.h | 44 |
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 |