diff --git a/corr/coins.bmp b/corr/coins.bmp new file mode 100644 index 0000000..d69308b Binary files /dev/null and b/corr/coins.bmp differ diff --git a/corr/coins.pgm b/corr/coins.pgm new file mode 100644 index 0000000..e6c5750 --- /dev/null +++ b/corr/coins.pgm @@ -0,0 +1,5 @@ +P5 +300 +246 +255 +1201001222211022321102422123212113323343333244344555544531433333345634434454344423455333442433454444556343333445454665566655676688987898779896887898898:9877898888898799999:<::;899:988::9:;:999:9;::97888:;::::9::899999:988:9779:8799887786678898998998989:;:8::9:9789::89;:9::8:98::9:9:::9:::;8889999:77//10212011110010122211112112233424222342344432222233333323334455334554445555443434343444544335455565665544444556535755446745577579888789768799:9988889::9798:8898879998::989;;99888:;:::;:::998:9;;:99999:;:9889:;:9::889898;;:9:9879999866788978898777999:;::;::;:::879:8988:8:98:::;98::8;:::;;:989:988897001111111212221111210022232233321222213323454453343344434435445544344564354465433343444435322443467457546555446454457655566565677778989997799::8:88:9:99678888:8866788999::99;;999:9::9::99::::::;99;:::9999989998899::889:;;;;89;88:988888789:88899878:99;;:;;989999:997899888889::::9989888:8:::88898999890210021220212212121220242123223111222133244535544345344334444444554444555444344431445443244344544554543676654544554589544666656868887678897799989:9::::9788777978879998:::8788989:999989988:;<98:;::;;;::89899:9978998::8:::9989899::888999988877799999:;;:::;9999999979988999999::::::9889:998:9:9:8878889700/0211210111101220212323212111123324442244332333343354324343442355764454453134354454345533565455765644576655456455666566666677888:99878:999899998788877886777978999:989:;9:888:99:::999;;;:::;:;;9;::::;989::;;::::89:::;::8789:9:98888:9978876788989:9;<;99889:9989978:;89:::;;89:9889:98:;::;:9::999:98870012232231000212112312232322221121124433333344232323423323344454343553344553245467644444455445545655665456555676656566656866776677998789::::9:=AEHIJJIFB=7356789:9879:9899:;;989::9;::;::<<;:89::<<;:9999989::::99::99;:::9999:99::99:99888988789988899;;:;::9::9877889989999:9;:88::99:;9989989:8::9898899:01332221310112212131223323321122211133334232334444344344443333344442455544443553456554335445554565656656567665765565557766767887778889888957789889888::9:;<:9:<;::<<:<;<<<<;;<<<<;:;:;:::;<;<;:9789:9999:;98877898989::9898989::9:;:;;;;899776898:::;9:9:;:;:8988::889:;9898:::::88833212222102212211111123233322232223333123333344334334545542344564454555345545444445545656335566567757665766766777777776676678886899:@RpƵV;7877887999:9:;<<<;<<:;<=<;;:;;::;<:9:99:9;;99;<<:999:89;::99;;:99998999::9:9889:8:;;9:;:9<<:9986878::9;:9::9:;;::9:::9::89::9889:9::99::20022220/00221331100112332223324323323444332234545554674445555454444656446567544646656776444445556755566688657667777889988987788989IgôxJ87889:;;::<;::=<:;;:::;::<;:::;<<;:::::;::;9:99998;;999;;::999::;:;;9:;;;:::;:9:;:9:9::;<;::::8:::;<<<;9:;:99:99:;<<;99;;87899888::::;;/0010110011112232322113310113344322232332544445435646766666554555655645555556555666556655545565675555655765667766778879979888:998;:;;;<;<;;;;;:979988::9;::::999::;:999:9;;:;;::;;<;<;=;:::;::;;99;:8998:8:;999989;;:::;::9:::;;:9:9211211101210124321222333234334233323323443333354454565567656555656446756544666677545675567776566667677768776767898787899::;;:::Cdȿ̋H8988::9::;<<;;;9;<<:;:;<;<;:;<;:9:;;::;;;:;;:;;::99::::9;;;;::::;<:;:89:;=<<:9;<;::;:::;::;::;:;9;999889:;;::::;;;9;;9:;;999999:::;2112111211322012344433322332233344444423355445567665566665665556655456666665466766556754676866566666777687777687888::::98:;::9ChþԐI778:;=<;:;:;=<:<=;;===<:;<<;<<::;:9;;999;;:::899:9;98::;:;:9:99:::;;9;;9:99;;:9;;;;;:;;:8:;::9::99;:9;:::988:9:::8:::;;:;:9:98999:21232223311332113334322233323334334333322444455655667765567556665665556567765555577666445588778778775676787778776778:887889::CfļH6889;;;<::<<<<;;;<=<<;<<;;::;;;;::;;:99::9::::988:989:::::9:989:;9;::;;:;;:;<::;;;<::::9;<;;99;;;::9:8999:88999::::;;<;<;;<99999;112312322201222212224343222244343554334354334455544566546776576666665577765667665886774675667786876655788866778987689888;;99?cпֈD88899;;;<<::=<<==;:;<==<:;:;<<<;;;;;:::;::;;;:9:;:8999:;::;;;98:::<;<;<<;;<==;<;:;;9:;::;<<<<<;99;:;9999;:88:;;<<<;:9:::<<:9:::;12120001001101012333333233223354555334545434655656565656677777656656557788656645677577678777668666666788788787889978::9;:;;=<;<<::;<<=<;<::;==<;;<;;9:<=<;;:;;::;<<=<<:;<;;;;;;;;;<<<<;:;::;99889:::=<:99;;;;;:8:;::;<===<<<;;;;:989;<<<<<323543321121212343233324443333343445555544556655677677666577787677678766788567688656756557787776675688887669997898899::<3334322223423112323434532434344456444655445776656666777678787766677898687787788888666777876779987886877866779:79;:98::<<==<;;;;<<=<=<;:;;:;;=<<;;:<;:::<<<<::<;::<<;;<;:;<=;::999;;;;:;;:;::;;<::99::9::;;;<<;9;;;;;<;;:;<;;;;;;<<1323233322332123112333443444355566566667677796688667676666878876678888798755677777877777777677887898778899889:9:::::;:;;ExץL57889:9:;;;<==;<=>==<<<=;<<<<=<<<=<;::<;<;;;;:;;<<<<;:==<;<<:<=<<:;;=>;;<;;;;<:98:;;::;;::9::;;;<=>=<===;<===<;<<<=<=<=>=<;<12223333331332223332223334544555555776677666778985567988776445787786678877556767867777776776889989878897:999899;99;:;;;:Tʻj7698:;;:;;<<<<<>=<;;;;<=<<=;;;<<<<<<:;<;;=<<<;<<<;:::<<<<;;;;;;;==;;;;:;<<;:;;;:9:=<<;:::<;<<;<==<<=<=<<<==>>===<<>=><=>==<=232232212221112333424443224456576446565556676667h׳ŊB7;9:<<:;<<<<<<>>=<<:;;<;;<:;;;;<<==:<<=<=;<=><:;;<9;==;;<;:;;;;;=;;:;;;;;;:99;;;<;9::;::<===:;<<;;:;;;;<<=<;<<>>>;;=;:<;<>>323212333332224422323434332344456556547667876:E[thM92467777667776688877788887787898687878878999:99:;<=;;9F|໧ţT7:88:;;;;;<<;;>>=;;<<<<<::;<=;<==<<<:<=<;====<<=<;<<;;;<;;<<;;==;<=;<=<<;;<;9:::99:::;<;;<<==<=<=<<=<==<<;:;=<===;;<<<=>=>>2333223333243432443333233423465455566676777=>=<=<<=<;<<;<===;;<;;<=><>=32344233422224323434454433434466655578787:Ljǵ^8355866778866779:776798679789899899:98999;;:::<>>=bйȹ{<8::::<;;====>;;===;<<;==<;===><==;=>=;;===:;;:::<;::;<=><=;;;::;<;;<<;::;;<<;<:<<:<=<<;<<>>=<<<<;;;>><=<===<;;;==<;;;=<<=;2323123333433433233365444444345655466877@^ʲK348788668778768776777689:99889879::999:::9;<:<<>pĴĻD4798;;;;;<<=======<==>=<;<==<<<<=;<<<<<<:;===;==;<<=<;;=<;;:<<<=<<=;;<;<<;<;:;<99:<<;;<<<=<===<<<=====<<<;<<<<<==<<==>===<332323332333332234224533554555554556678Hoä_5677755766987788797789899999999889:98::;;;;:;:@}չL3789;;;<;;<<<==<==<<=>>;;<;;==<;<;<<=<<=<<<=?===;<===;:;;;;=><=<:;=<<<<=;<=::;<;:;<<<<==<<;==>=<===<==<=<;<=??=><<=>??>=>=24344334223222222243221346554765688668Myžಪûƽ͵l8556777668878989:97886688::9;;:9:<<;9:::;;;>>=<<=<;;<>=<===<<=;;;;;>=>==>>==><=<;<<<;;;<;;<==<==;<===><;>=<<;:<;=<;<=>><<<;;><>>>=<<===>>;=??>=>334432343354433433454324343354555777:O~˺Ҷ׹q93688:99877699999:997888::8;::::<;;<;9;<===:HûDZÿ[58;:;<;<<<<<=<<<=<=>=;;<<<<;;:<<<===>>>=>>==;;=<=<;;;<=>>>===<<;<=>===<<==>>=<<;;:<;<=====>====>=>>>==>==<>><<>=>>>>>=?>>>44333333554343345434443455565577668:Tɷ¹߳p946788788988889898998988899::;;;99;<=::<>;9K¿°^479:;;<==<<<<<<=>=?=<<<<==;;<><===<<=>>===<<<;=<<==;<=<=>>>=><<==<>>>><<<>>>;;==<;=<<<<<<=?>=>=>==<<=<<<=><=>>==>?>>==>>==53333332353224434433554455544566887PμĿۮd54567669:9887:9879::98:88:;;;;9:;;;=;<<;;:Lι_5889::<;<<===><<=>=;<>==<>>>=>===<;<><>><=>=;<>=<;==<<;;<>==<<>?=>====<;;=<;;=<<<>=>>=<=>>>>>?=>?<<<<>>=<=<<<=<=====<<<<;=3444544334534343466455565455776787Cz¿ƾѺ؟Y5657889898899:877999:9:9999;;9:=;;;;=<;<:Lսø`59:;::;;:;;<==>=<<=====<<=>?==<===>>=<<=<<>>=<=>>;;<<;<;<<>=<=>??>=>=>=<===;=>><<<<=<===>><>>>>=??>==>?==>?==<=>=<>=><<>>@444444345554434455444445666778887:cķȿȎG4667998877:99:98999::989789:;;<;<=<<<;<;IԿ¼¸W3:;;<<;;:<;<========<<<<==>><<<>>>?=>===<;<==<===<;;==>==<====>>>=??>??>=<;:<=<>==>=<<====<=<<>>>@@>>??=<>>>==>???>>?>>?@@333334534455435543354644467766887Oǻɼl6588987799:88:99:98:::::99::;<;<<=;<=;=;Cý¾¹¼N39::;<==<>=>><>>@?><<==><>?=<=;=>==<>>>?><;<>>=<<=>==>>>>>>==>???=<<>>=<;=<;;<=>>??>>>;<>>>>==@?=>>>===??@??>?>>>??<==<>?=44445664434555345454566665566776?rź۔J3777668997888:;988988:;<<:;<<<::;<<<=><>ĴH49::<<<<<=<=>=<=?>=>=>>>==>>??>>==;;==<>>>==>??<<=>=>>==<<=>>>>>?><===<==<=><====<>>==<<=>><>@?>=>>>=<>>>?>==>>?>>===>?==<44335554434565556744665676677787Wôʿαp956577898799888889::9:;;::;<<;:;;<<<<>=9y˸»>7::;<==<<;;<=>>====>>=>?===>?>>?>>=>><=>??>>>=>><<==@><<<<>>>>>=>>?>==?@>===<<==>>==<==>>===??>>>@??=>@>==>>=>>?>?===<>=<<3342434456655476665455465776677?uƻлؽҿĿĿùҐI465889997:978888::;;;;::99:9:><;;<<;==8hӺp78;;;;<==>>>=<;<><<<=>>????==>>>>?@??=;=<=><===?>=<>>=<>==>===>?=>=??>>=>=<<=<>>>=<=<<=???=>??>=<=>>>??>=>>???@=>?@>>>==<;=5653444645554356566777675666556L°ɫĻǼԨa756789899:::99;;9:89;;;<9::;<==;;;<>@?8UƻY59:;=<;=>=>>><<<<<=>=>>>>???>@@>>=?ACFGHGFA?=;;;<>>=>;<>>@?>=???>>>?>>=>=><<>==?>=<<<>>>>?><>=>>==?>==><===>???===>===<<==<5564344554544566656765467555759_ɽҿɾĹĿҺ|=357899:9:9:9::;<;::;::;;<<<;;<;=<=?@@;D֨ǴB599:<=<<==<==<=?>=>>>??==??>@?ADKT]fsxrdTG=:==<<<=<>>=?@?>==>>??=??>>>?>>?>>>==?>>=???>==>?@>@@=>=??>>??><;;;:=;;;=><355653345555676755665457655667@u¼ƺ³ĔJ24699;;9:9;:999;<;;;;;;<=:;;<<;>==??=>9|߳мȺt88:;;<=<;=>>??=>??==??@@??@>?CK[n^F<;;;:;<=???====>=>?>=>=>??=>===>=?>=>>>???>=?>?=>@?>?@@>?>>><;<<<>>>>>>=432465553445554466678655566786HƿüǿŲʦZ3578899:::;:;<9;<::<<==<==<<<;:<=<<<=@9U՟½·¾R59;;:<<=>>>>>==??>>?>??@@>@GUhƿmL:8:;=@?>====>????=?@>=????>?>==?><>?=>>>>?>?>>==>????=>=<==<=>?=>>==<=544345455434434456558767778996QͶʰg5578978:;:::9;:9::::===<<<<<=<<=<;;<=?<;˹ǻy:79;;<<=>?>>>=<=>>?@???@?AHXu¿ʾǯ~L:<===>>==??>?@?>>><=>=>>>>==>==<=>>>?@>?@@=>@?>?>??>?>>=>>>?>=<==;<<445534543444364457645677868:98^ʶƽɿĸv9478899::;;:9;;:::;<>=;<;<;;;<==;<>===?7b騥ֻŷQ49;:;===>=;<=<>=>==>@?@@ETpƼЦeA:;===>?AA??@?=>>??>==>=>>==<=>@>?>>@???>=?@A?@??????@?>==>=>======55455454334335654556578788887;hſϴij<599899;;::;<<<==;;=<<<<==;<<<=<=>>??@@?>ğ·r889:<==<>==<<=<>@?>>>??AKcƾɽɼH<<;<<=??>>>==??@@@>?>?@?AA>>@@>?@??>=>>>>>??@==>>==>>>>>=>>=>>=??65545444333555565566776789889=o˿Ѿɹǯ½@58::9:;;;:;=<:;====<:;<<;;<<:<<=>>>>ABB;UεûH2999;<<=<<=?>@>=???A@@AOlͶŽR:;===<=>=>>@@?=>?>?>@A>?@???@@?@?@@>>>>><=???>>=>=====>==>>><<>>423333344455565577766899989;9=r˿µ@3899999;<;:<=<:;<<;=<<<;<<<=;<>=>==??@@?:xޣٵ_358;;;>??=>?@?@@@>?@BADTuȹ˸Ƥ]=<====??>>@?=>>>==>?>=>==@@@@>>>??>??=??>A@?@==?@?=>==43445433544356666666789::;::9=qѾǶǻɱ>4999::9::;;<=<=<<;;==<<==>>=>====<>??>?@>?ќ׺l439::=>>==?>>@ADX}ջdz¯c>====>>?@@?>AA????@>>>==?>>>?=????>>??AA??>>>>@??>=>>>??>@AA@?555454446656567555777888::::9;mȷ̺̿緰:39:;;;;;;;;:;:;<<<;<><>=<==<>?><=>=@@A=BǷw6.8;<<=<=<>==>>>>@A??@BTzɳĽ·°`==>?@??A?@@?@@AA@@?>?>>>@??@@>>>@?>@>?@>>?>???>=>?>?>>==???>?576546645665556577677789:99;;9dӶ÷˽|759::<<=<=<::;;:;=>>=;:;<;<====<>=?=>>@@??@=====>??>@BAA@BRwƳɿǭ[<=>>??>??@>>@A?>??>>?@>@AA?@A@?@@@AA?>@?=>?>>=>@@@@@==?====@5665565565564666776678:99;:;;8Yϸƪʿƾÿk35789:;;<<;<;=<<=>=<<<;;;<;<>==<==?@?@?????A=Tw:*39<<=>?>??@?@>>>???AAMoϺĿřN;>>?@?@@?>@@@?==>>=>A@>>?>@A@@A@?A@>=?@>@@>>>@?>>@?>?>>???@544545566557676766779999:;::;8PɾñοھX0589::;;9:;<<<<<<;;;<<<<=>>>>>>?>?>>>>?????@A?>?????@??>>??A@Ff綧ƸD>@??>@@>?>>=>>=>>>??>>?>>@?>@@?>>>=>>>@@@>?>?>==?@@@??@@AA65447657678787666889988::9:;<9DŽǵĻɮD199:9:::::<;=;<;::<<>=>=<<=>=>?===<=???@@?@?AC?DԪc/&19<=?>>>@@>>>??>@AAAAC[ߨýa=>>>>?@AA>?>>@???@?>=>@@=??>>??>??@@@>@>>@@?>==<=>?>>?@AA@5656667667877877887:88;98:::<<;lӾϾξx439;<;<=<==<;:;:;;<<;<;=>=<===>?>>=>>>?>??@??@ACA=_~F#%19<@A@???@A@A@@O{ֳ¿ȿŽO=????AB@AA??@?><<>@?=>?=>>?>??@B@@?>??>?@@?=><;=??>?BA?@?6898556777877778878:9:9:99;:<=9T଩űƽ]-689:<<<<<<<=<<:;<<<<<<>>>>>>?>>?@>?==?????>@@@AA@7Eʳ[/&3:@>?>>=??@??>?@@@@Cd˾ƽſqA>>?@?@?A@???>?>=>>@>>=>>?@@A@?@???=??>????@@=>?@BA@@@>?@899888788777677898:;::8:;;:;<<:AƷǫȯ¼@08:;;;<=;::======<;;=>>=??>?>>@>?>???==?=>>@AAA@??@;6Qe9 ,5>?@AA@A@P}Q<=?@A@?>>=>>=>@@?@A@@?@>???BA?>?@?A@??@@@>@@@@BA?@@?>@>@8757887788888779979:999::;:::<=9dݾijÿʵc-799<=;<=<;;;;<<===<<<<=<<;99:;;:;:;=?>=>?>??@?@AA@@A>74Jy~Y5%2:<>??>?>??@@>@@@A@A@AAACdķռľÿg=<>?@>??@>@?@??A????A?@@@@?@?@?>>??>>??@@=?@@@@@>>?????A8766888789898889899999:989:;;;<:CȻλҳȹ7/99;;<===;<<<<<;;<>>==<;;=AFKOOOMHB=99;>>@@?@@>?@AABAAA=1.>?>?>>>=>??@@@B@?>>@>M}Ƽ¶ûG;<<>=>>@?@@@??@@@??@@@?@??@AA??@@>>>>>???@@@?>??@@A@@?>76987888988889988::9:;;;9:;;::;;6[üضȺɻQ'69;=<;;=<<==>><>===>=:;=?@@@@>@@ACBCB@;2*)0;HT^fkjjgaWG:."(18<>@A>??@@A@>?>>AA>>?@>=>?A_ɿĽ¾ÿ]>====>>>A@>?>>>A@A@???@>@@@??>?@AAA@?>@@@AAABBBAA@@AA@>68888777886997898999:;<;;<<;:<<==<~Ĩï¾ͫ帬u02;;<<<<<<=====<<=<<:>Lc|dC8:=>??@???@@BAA??;4/($ !"!$-4:>?????@>>@@BBA@>>>@A@??????EsƾŸ˿¿uA<>>??A@@B?@@AA??B@??@@?@@??@@??@AA@A@??@A?ABBA@@?@?>?A887878899889:879:99:9::9:;:;;<<>=9Aǡɹdz͹ϿӦϥ5,:<<<;<===>>==<;==????@A@>?@A@=:62.)&''&('*-17:=>??@AAAAB??ABA@@?@??@@@?@ABA?MĺտH<>?@AA@?AACBAA>?@?>??@?@@AABC@@@A@@@@@??@@?????@@@A@AA9899699:8899:::::8888:89;;9::<;:;<8UμϷΧ@'8;;<<<==;;;<==<>?>@@AA@???AAA@A@?>>=::::;;=>??@@AAAA@@AABAA@@?@A@@?@@?AA@@@A?Xظ˸V:?>>?AA??@@@?>??>??@@AAAA?@>@BBA@???@BAA?A?@@@?AAAAAA@87889998998998999;;99:;;;:;;;::<<=?6Tװ֯Ѱ?#6<==;;<<;:;<;>?@>@Puwpx|sjc_ZQQVYYVTYhvzwsrD:>?@@???AAAAAABAA?>>>>@??@A@?@@@@BBAAA?????@AAAAA??A@?ACBA@>@dү`:>???@A?A@@A@??>?AA?>?@A@?@>=@A@@?ABCCBA?@@@A@A@AAA@?A978998989::999:::<<<<;;;:;;<=<<;==?<5VἘ|<1:;;<;;;=<;<=<=><@[~jlurjcbddc^^hmmmeZVaptrux|ϐG9=>@??@AA@@AAAA@@>=?>>AABA@>?>@@BBABAAA??@@BBB@??ABAAA@BCC@Dq¿Ŀ¹f:=@@?@ACBBCAA@@@BA@@?@@?>@AAA@@@@@@CB@@@@AA>?@A@@@AAAA9::99789::8::8:<<<;::<;<::<<===<<=7Iˡ˪n.$4:;:::<<;>>=>>=?=BizjnoiehmsrimnpwslaZ[krtwwxЪS;=>@@@A@A?@@AA@?A@AA@@AA@AA@@?BBCBA?@ABABBAAB?@ABBDCC@@BBAF~½ʾm?@?@AAAAAAA@@BBABBCA@>>@BCA@@A@@AA@AA@AA?AABBAABBB@79;;;:9:989::;;;:;:9:;<<;:;::<<<==>>>?7=ưW%%8<<;=<;;;=<<>===>>@A@AA@A@@@@@?@A@@@??@BA@@@AABAAAAA@AABAAAACB?AAAABCB@K¼йǼ½x<=>?@@ABBA?A@@A@AAABAAAA@?@@AAAAABCABB@ABBBA@@@@@BBAAB789:;:9::::9::;;:;;;:<<;<=<<<<<;<>=====81Rže6)8==>=>==<<<=<<<=<@hzuxslgluvrnnyzx{ztqhjf\lxwuutM:=>@@@@@@A?>>?>???@@@?@@@>?>?@?BBCCCBBABA@@AABB@@BBBCCC@KȻºû½}==@??AADB@AAA??ACBBB@@ABBAABAAACBCCABC@@BAA@AABA?AAAAC9888999:;;:;;:;;:;::;<=<==<=<=;:<<<<=;<==42Wj=!2:<<;;=??===>>?>>>=]stmecjsw~}z}yxzwpqtc^u~yrsuA:>@BBBB@@AB@@BBBAAABDBA@AABABCAABA?AA@BA@B@LȾǸ}>>AA@AAAA@BCBBABBBCAAABAB@ABABCDCBB@@AB@@ABABCCBAAA@AC9:988998:;<;;<<<:9::<;:;<;<=<=>>=<=>?=?>>@://Ire9,7<==;<<;==><;==>>?=Kuzrhfjt~w{{tqrxyxvchzxoov~t=<=>?@@@@ABA@@ABBACBBCCCBBCBBAABAABABC@ABBCCB@@AB?ABBCALľѾA=?@AAAAABCCBBBABBCBCCCBBA@@BBCDB@CA@@AAACCABBBABBCCCA;:999;;99;;:;<:;:9::;;<;<=<<===>=<>>>=?>=>>>8/*1D[rydA&-8=><<><=><<<<;;==Ajqrutor||vw}tcrxrmsv޼U9???ABABBCB@?ABBA@ACCBAAAABCCBBAABABA@ABABACCABABBBA?;9:<;99;;;:9;;9:<<<;==><<>=<==>=>>===>==<;=>><;70)%&*-/-.-(!#,49:<<==;<<=>?>>>==<<=>@AA@>>?@ABB@BAABA@AABABAB@@AAABACDCBB@@AA@BB??BBCADƫÿľw<>@??BCCCCDABCCBABB@ADBCCCDDCCC@ACBAACDCBABBBADCBAACBA899::::;:;<<<;<;;;=><<====>>;:<>===<=>==><;;==<>>>950,++**,059:=>=<:<<;<=<==>====<=<>>Am|jrstvy|zvz~uevwsvydzQ9>?>??@>=?AA@@ABCCDBA@A@@ADDCBAAABBAACDDBBCCB@?>ACCED@Գɳj9>?@@BBCBBCBDDCCDEDCBDEDBCCCAAABACDBBABCBBBBBBBB@@BDDA9:99:;;;;;<<<<<=<;<=<<;===>=<<<>=>?????>>>=;==<==<<>=><<;=<=<<=>>=<<=<:<====??>==<>??=Qpmps{zzzl{{utxy|==>>>>BA@ABA@ABABCCCBABBAABCCBBBBBCBACCBBBCCDB>>AACCB>pß¾_9?ABCCBCABCCCFEFFEDEEDECCDCBCCEFDDCCCA@CBBABBBABBBB@@@;<<;:<<:<;<==<;<>>><=><<<=?>>>=<<=<;:;<::;;:<<<==;;===<<;<<=<=>=>===<==>>==>>??<=>==>?lkotv}z}zr|}wrƥK;>?@@ABAA@A@AAAA@@AACDDAAACABDCCBBCDBAACCCCBBA@@@CBB=aҡÿT=BBBCBDEDEEDDEDEDBBBEDABCECDEDEDCCBDCCCACCAACCDBCDCA@A;;:9;:;:;;;<<==<=>><=>?<<>==<<==@EIXd\UTOHA;9889;<<;<=<<=<;<;<==;<<<===>=<>=<=====>>?@A@A@?@B@ACCCBBBBDCCBBBCBBBBDCCBCBBCBBAAAB@>@ABDEATſø¹¿H>BAABCDEDDDECAACCFFGEB???ACCBBA@ACBBBCBBBAAACDCBCBBAA@:89:::;<<::<;<>><<<<>?@?>>?AJYpfPA:89;;<;;:;<=<<=>=;<<;<<;<<<===>==?=>?=^lpgnqrxtw~°~}uw<:=?@ABA@BCCA@ABCCABDBBCBBBCBB@ABBAACCCCBBAA@@?@ABEFCEǺo>=@AABCEECCDFMXi}t_NE?@CCCCDCA@ABCCAACDCABBB@?AA989:;;<;;;:<<;<===>>=<=>@I`yǾtN;79:::99:;<<>>>>?>><<<<>>>=>>><>?>=?mdiny}v}{s{Ĥ}wC9=?Kf»Z=69;;;;<;;=<>>>=>=>=<<<=>>@@???@=C{zbmomos~~vsJ9=?@BBB@AAACCABABBBABCDBCBCBCBABBBABCCBBCBCDEDBCDEDFAYȽH=BCEEFFDMgƷďX?@CCDDCBCCCDBAACCAAAA@@<;<;;;==<==<;=======@LgŻ[<7:9;<<;;<<<<=>=>>><;=>>?@@???>?>?=C^ļε}G4799:;;;<<;<>>?=>>=>>?>?Or˹ɦb<7:::;<=>==<==>>>?>>?@@>?@@=Qugxsqu~{{rq\5=@@?ABABAACDCCCCABCB@AABACBDCCBBCDDDDDDDEECCEEDDDEEEBDҼŹh9@DEEGbQ>CBA@@@@BCBCC@@@?A;;:;:;<==<==<<>>AW}ѷŷʽ|D:;;9:<=><<=>?=>>@?=?>???@?<<<===>><;>?>?>CbͽоȀE89:;<<==>>=?>??>==>?A@@?;Qqjs}zmovz}~znhnZ5=>?ABCAABABABCBACCCBABCCCCCCCBBDDCBCDCCBABDECCCBCDFFGF?ùƷſR;BEEGna@BBCBBCCB@@BAAA@<===;<<;<;<>>=Bf¿¾̊E7::;<=<=<==<=>>>>>???@B=Qts~roloy~uzrsW5=??AAAAAACCCEDCBCCDCCCBCDDEEEDDCCCEEDDDDDCDEDDDBCEDCDFCJᛟ½Ķʷҽd:?CDFj\=BBBBAABBBBB@@A:<<<:;<;=;=>==??>=>?@@=Nxgvvpqppru|ws~|~M5>??@AABBBCCCCEECCCBCCCCCBCCCDDFFDEDDBDDCEFCDDDDCDDCDEEE?Y͚ļܻѽo<=DFFdS>BA@@ACBCCA@AA::<==<=>>>?=AgijʿӿϿξv=9::;<;=>><>>>?@>=ABAA>Gpnornosolvyv}{wB6>@BBBCCBB@AABBBBB@@AABCCBCCCDDEFFDBEDEDDEFDDEFEDEDCFFFGF>aԝĽy?8@DESG=@@@@@@@@BAAB;<<<=>=<=>?=R¼Ѹb<:;;:;<<==<=?@?@@@ABB?@usg`blrsqz|tq~s99>@BAABABBAAABBBCCBCCBCACDCDDDCBCDEDCAABAACCCDDDDCEEEGGGFF=iЛŽ괸ý{@5>CBK{ַÿ@AA@@@A@@@?A@<<==<>>>==;IlɿкR6:;9<>=>>>==>?@A@?AA@>@@<==<=>??>>Agη¿·@7;;;=>??@@>>>??@>>?B=Twnpq|w{{sxz~zsw{~sxC3=@@ABCCDBBEDBBBCDECBDCBCDCDEDCDHS`n}yf[YMCACCBDDFEEFE=Ib51>UƸǻͫȬɶ[7:::;;<<<>>>????>>@A?Dύz~zwtquut~~z~z}}l~28>>@BBACDDCCDCACCDBBDECBDDDBCIVl`IAACDDDFFGE:?˞⭕O//;CFGGpº㼛ǹa?A@AA@BBBAA<;=?>==<6:;<<==<=>>>@@?@@AB@<|tx|yzrqwu|}v~{{~~zjnX.;?ACDDBCCCBACCBDCCDCCDDDCCFNciMAACCEGGFH>0T߸Ɵt@(.;BEHHOؿʸ?A@AA@BBBBA:<==<>??@_ÿƴ`6:<;;=???>?@@@???@BA=Lcln|uoqt{y|yz|~wlj~45?BCDDCCCCCBBCCDCCDEECCDBEQfĪwRB@CEEDGG@17h͸}Q/)3?DDDHFbĺŵO?BBAA@AB@A;<;:;=?=R˺żȈ>8=<;===>?@@?@??@AAAB>=>>=<=>>>>>?AA?@ABC@Bxfeltsosrnrz~z{tjflzyvxslrn,1>A@AAAACBCCCCEFDEDDDEEGXuİѽ¸TBCDCCEED>1(/Kt~iN."'3>BDFFGHFPȻýք?=@ACBBAAA;;==?>>CxǼƸz{o39<:;><>?@@?@A@??@BBB=`܁lekrtnnrsswywnmqpopsrwpbg9,?@@??A@??@BBCB;k^els{~rkjosyyjmooqrrrwvmxxpfdI*6ACDCCBCDDDCBDDCDDCDDCEbžȵeBCCDCDDDC@:2*# $)-38;<5.*'!#+5?CDEEFFEFHHFcͿ٬T=AAA@@A@A<<=>@@@aǴ¹ʾƝI6;<;>?@????????@ABBBC?Bgaszoikqrstwusutx}{olpnsjlrhdV)2=@CCCBBCDBBCCDECDEDDCFc½¾ɳg@BDEDCDCBCB?:2)%"#(-5<@DEFGFEGGFGHIHqʲȺշb=ABAAAA@A;==?>?Coƶ˸¼¬\4:<<>>?@???@AABCBCCACB=:775579:<@BAECEEEFFDGGHHIHH}֯ʿϻk<@@A@@@@C>=Iǰ÷ѳſĹl28<;;:=>?@??A@@AAABCBBB?>>;Núöɻ³üe36::;<>>>AAAA@?@@@@CBADD=@zܒhcqrmplimmhgpurtkltwrleduG!-;AABCBBBCBBDDEEEDCFGFFYİ˼¶î͖L@BBCEDEDDFFFGEEFFEEFEEFFDDEGHGEFFGFFFFFGLӿļϾn>ABBA@AAB<>===;Vľƾ¾ǹû´o469;=>>??@?BB@@ABAABC@ABB=7b}iokhihmonnnsqgfhnmfeqw=$/=BBAA@BDDCBDEDDDFEEDECQ­ƻwFACABCDCCCDDDDFFGEEFEEFEGGFFGFEGHGGHIGIHLº˳žƵk=BAAA@ABA<<=@>?AA@?@BCB@ABAACC?6Kջi]_`dlnohadacbaesQ0'4>AAACBABCCDDDFFFDDDEEFElȺaAABCCCDDBCDFFEFFEDCDEEFGGEFDEFFFEEEFFHHHžӰ˿Ųf:?@ACAABC<<=?A?dķǣľ:6::>?@@AA@?@AAACCBABCBBBAC@77Rǝ}kccijccijp~c4#+7ABBABCBDDDCEDCDEFFDEFFFXM?BCCCBCCCFGEEEDFECDDDEFEDEFGGGGFEEGHHHDsϭĬ]:@BBCCBBB==>?@>hˬ½…;4:=>>?ABA@??ABACDCBBBCB@@BCC>56GopM.#'3>ABAABCCCEEEEEEDCCDDFDDDJunD@@ABBDDEEECDDFFFFEEFEEGGGIHEFGFHIIIIIDjɺßP=DDCCBABA==<===hо̿½l57;<<>>?AAAAA@??ABCBBABCBBBCCDC?704BZpubM=1$"*4=AAACBCB@ACCDCDDDEEEEEEFEC\T=AABCBBDDEFEEFFFGFDFGEGEGGFEEHHGFFFHJCe翷ȲͽD>CDEEDBBB<==>???>?@AAAA@BBCBCB@@ACDCCCCDDC?90,+-2563.)% "+3;?CDBBACCBACDCBCDDFECFEFGEFFKzĹzD@BDECCDEDEFFEFGFFDDEECDECDFGGGFEEFHIEXƵv?@CCCCCCDC<;=<>;]ûķT/9;=?>?@??A@A@A@AABDBACAACCCCCACCCA@?;730-,-./38=?@BDDCBBCCDCDDEDEFEDEFEDDEFFEEXVABCCEGEDDCDFDEFEDEEEFFGEFFFFFFHFHHHGGJɾb??@@ABBBABBBBABCDDBBABBBAABCBCBBDDBCBABCCDCBADEBAACBABCBCCDEDCDDEEFFDFGGgh@@AACEEDEEEEDEFEEDFFGFFGHFFGGHHGGGGHIB{ľѽ˻J>=:V͟ǽƺ£ûŒ6.9<===?@@?@A@AABBBBCBCBBBBCBCAACBCDCDBBDCACCCBBBACDDCBBCDBABDCBCDDCDDDEFFEEEGFK~z}E@@@CCDEFEEFEFFEDDGGFEEEFEEGGHHIFGGIJEUĺo:@CCCBCCCBD???>?;J䵟½ѽľÉ10:;=?@??AABB@@@CCACCCDBBABDCBBCCCBCBABCCBABCCAACCCCBCDDDDDCBCDDCCEDDDEEFEFEDGDVŴ~{M>BBBBCEFFFHGFEEEFHHGFFEDFGHHHHIGFGIJHEL;DEDDDDDDAB?===?=@ݼ´®ĸu.3:<=?A@??A@@ABBCCBBBCCCDCCBCBBCCCCECCDCDBBDCCBBBCBACDCCDDDFCBCCDFEDEEDFEEFGEDEhDZZ=CCCCDDEEEEEEEGHGFEEFGHEGHFHGFGHHHHIKFZӿžf9ACCEEEEDBAA<>>>?@:n̡ͶŻT/8;=>A@AAA@@AABBBAABCCBCCCDCCDCCDCBBCCBBDDDCABABCBACCABCCBDDCCDDCDDCEEFFEGFFEEGwȰh>ACDDDDEEEDDEHGGFEEFGGGHHHEFGFFGGGFHKKD¼D:DCDDDDCBCDD:=>>>>@A@CDDDDDEFGEDDDEHHHIIHGGFGHGFEFGHHKFK¿̼ĿS8BEFEDDEDCEED;>=<;=>?ABBAACCCCBBCDCDDEFDDDDDCCDEFDDCDFFEEIJFN¾гȴG>@ABDEDEEFGGFFEFGGFFEGFEEGGGGGGFHIIIIIE^\6?DEDDDCCDDCED====<>>:xݚ»͵=.8;=>?@@A@ACCCBCBACCBFNXjvfTE>@CBBCCBBCCDCCCDCCDCDDDDCEEFDDFGGFEFGDR»ƺdzټK?BBCDDDDEHGGFGFEFGFFEEGFEFFFEFHHIHHJIJK@qļ`48BCDDDEEDDDFFD>=@@?@@?AABBAACDCCDCBDL_w¾ŽĺnRC?ABDBABCCCCBBCDDEEEDCDCCDEDFEFGGFIEU´лJ=CBBCBCDEFEEFGHEFGGGGGHEEHHGGHIIIGHJJJKJ?rܾÿg36@BBBEFFEEEFGFD>=???@??8sᬬǽõB,8=?@@BABACDAABDBCMhĹ}XD?ACDCDCDEEDEFEDCEEDCCCCEEFFFFFGHDXͺKfϿg57BEEEEBEFEEEEFFD=>>?@???;Eʡƴȶb)1:=?@@AABABDCBBBNjU?@CBDDBCDEEDDDDCEEDECBEEDEEEEFEBWʛ{E?BBCDDEEEFFGHGGGIJGFGGHIGGIHFEHHFFIIGGHHIG==>>@@>BCAAACDEECHarG>A@AABABDDDEECDCDEFDEDFGFEFGFBSſŽƾs@?A@BCDEFFDDFFFGGHHEFGIHGFHIFFHHHGHIIFGHHIIG=CwR17DEEFEEEDDDEDCEEEC>=?==>?@@@<䦠з?'3:?>?@BDBBDDCBV˼Y?@BDCAABBCCEDECDCDDCEGEDFFFFBNκĹԭ{@?ABDDCCCEFEDDEFFEEFGFGGFFHGGFGGGHHIIIIIGGGII@;We?/;FGEEDEFGFGGGEDCCBD?>==???@>@>BǕĹŽ¼ɾJ#/;?@@ABCDDCDCC[{¥hCACBCCCCBDDDEEFDDDEEEDEFGGGDKųѡp=ABCEDCBCEFEEDFFEFFGGFGHGHGGHHHGHHFFFHHIIHFHIIC78UxhA+1?EEEDDFEEGFEFFFFDECE=><=@?@><>@;MŚ˵R"*7?AABBBCCDECEc|qC?@ADEEDDCEECDFFEFEDFHGHGHFG˵ʣc:@AAAACEFEEFFEFHHHHHGGHHHHGGIHHGGGFEEEFHIIHGHIHG@64D[{|Y7*/>??==??=??;YҠĮc''5=@AAABCDCDDEfsD?BBBDDEEFECDEECDEEEGHGGEEDv[:@CCACEFEFFFEFHIIHIIHGHGFGIIHHFHHFGGGFFGFGHHGGHGIGA6,0?O`rteP>/(/DGFEDEFFGFDFEDGFEDDCEEEDD<;==<>?>???@@?;P٥»ܿR$#1;?BAACCCEEED]ĿĻȫc>ABBC@BBCECCDDFFGFEFHGHHCSƼĻ¸B=DBBEEDEGHGEDEGHHHGFHJIHHIGFGFFFFGFGIHHIHHFFFGGIGGGIIJHE@:53313459>DGHHGGGFEDFGECDFFFFFEECDEEGGE=<>@?>??==>?@@?=Dܹȼ·A$1;@BCDCDDFFEAVˢX=CCCBDEDEDEEEFDDEFGFEGHFDœϺͪi;ABBCCCDEEFDFGFGHHHHIHGFEHIHIIIHHIHHGGFFGGGFGFHHHHIIHGHJKJIHJJJIGHIHGFGFFFFEEDGFCDEFFFGDEGGHGGHHF<<=>>>>?>????@@@<=<=>@@?>@>?BCB?@CCBBACCCCCDiʬνmABCBCEDBCCDFDDFFFGHHHGGBZݗϹoq|zƯ}>@BDEEDGGFEEFFGIHHHGIHHHGHIIIIJHGHHHGGHIHHIGIJIIGHHGFGHIGHGGJJIHGHJIHGHHHFEEEEFEEGFEDEEDDDDEGFEEDD>?=?=??=>>>A?@ABCAA<8Oȵb: $3=@AACBDDABEFFCW̥S?CAACEEDDCEEEEEFFFGGEDD?󯎛ƾ[;ABDDCDGHEFEFGGGFFGHJHHGIIHHIIIHHIJHHGFHGHIHHIHHIHHHGHHHHHFHIJJHIHIJJIFHGEEGGFFEDDEEEGFEEEEFFEEDEE>=<;=>==>@@??A@ACBBB>:9KqlK/,8@ABBCCCBCDCFGFIyyAACCBCCEEDDEFGGGEDGFEEF>Sَi<=ACCEEEDFGFFEGFEEFFHHHHHIHHGGHIGHIGGGHHHHGIIHIHGIHHIIIHHIHGHIHHIIHFHGGGGFFGGFFFGEFGGFGFEEFEEDDECGF;:;<==??=>@?@@BBBCBBAB?74@RnmT5 !+5<@A?AABDDDDDDCFCYǼ½áT>CDBBDECEFGHGHGGFGGEGGF@y㲊G7ABBCDDEEEFGHHFEGGGHIHGGHGGGFFHGGGGGHHIHJIHIHHIHJHGIJIJJIHIGIHHGHGFEGGEGHFGHFEEEGGGGFGGEDEFDEFEEFHH=??=?@A?>?@AA@AA@@@AABCC?9448@LU^eln`[_UG:, #-59>@BBBBDDDDCCCDCDDH{žwB@BBCCCDEGFFDFGHGFEFGHHEE喑d2=@BDFFEEGEFFGGEFFHGIHGGHHHIHIHHGGGGGGGIJIHGGGIHIHIHGIJIIGGIHHGIGGHHHHGGHGFGHGFEGFGFFEFEFFFEEEEHGGGF?@?>=@?=??A@@?@A@>>>@ACBCDC?7/+&%&&$!" #*3;@ACDCCBBDDDECCCCDEFASȽřO@CDDCDFFEDBDFDFGGDFGGGIAdϊ~78AADGFFFEEEGFEFFFGGGHHGGGHIHHHGFIIGGIGGHJIGHIIIHIJJJIIIGGGHHHHIJHHHIHHHHGGGFFGGFGGFDDDFEFEEDEEDGGFEF>==>>>??@?A?>?@ABAA@A@@BCCCCBA?<9630./2359;>BCCCCCEEBCCDBCCCCDDCCDEgÿɹҺȹg>BCCDDEDCFEDFFFEEEEFGHGE>寉D3?BCEEEGGFFGFFEGHHIIHFFHIGIJHGGGHIHGIIIIHHGIJIGGGIKKIIHHJIIIFFHJGGHHIHGEFFGEFFEGHFGGEDEGDDDDEEGGHEEDF==>?@?@@??>??@AABBA?@@ABBACDDCDDCDDAAAAAAAABCBABCCEDBBCDCDDDEECDEFL־ÁA>A@CDCDEDFFGEDEEHGGFFGFCD㯌J1>ABDDEEFEDFGGGGEFGHGGGFFHHHHGHHIHIIHHGIIHIHJJHHIHGFGGHHHIJIIHIIHGHGGGHHFDGHGFFFGGHGFFHHHFEFGHGHHHFGGH<>>??>>>>>@AA@@??@ACBAABDCCDDDCECCCBBCCBBABDBBCBBACCDDDEFDCDEEDS˴б̜L=CBDCDEFEDFHGEDEGFGGFHGFBF骉L.9BDEEEFEFEEEFFHFFFFEFGHHFGHGFHJJHIGIIGHIJJIHHHIJJIGFGHHGHGGGHJIIHIIIGHIHGFFHFGGGGGHGGGHFEGFDEFEFGGGGHH>???<==>=??>>?@@???@@ABBBBCAAABBCCEDCCAADEEECCEDCDBBDECDDCBCDEFFFBaƫĽƥƭY;CDDDFFFGFEFGGEEEEFFFFHGGCJഊW.8BDEFFFEFGGGHFEGGHHFGHGHHHIHGGIIJIJHGFHGGHHHHGFGGGHIJJHHHHIHGHIHHHGGHGHHGHIGGFEFHFGFEFFFEEEFDDEFFFGHFFF>>=>>>??=>>?@@@A@?AA??@A@@AAACC@BBCCCBBDEDEFEECDECCCFEDEFCCCDDEFEDpٿɽǸf9ACBBEDCDDFFFEEFFGGFFFGHHHCAּ›K)4@EEDDDEFHGFGHGGFFGGGHHHGIIIIHJJGHHHHIIIIJHIHJJIGGHJIIJIIKIIIIIHIJHHHGIFFGGIIIGFGFDEEFFGEFFEEGFFFFEEFFFF>><=<>@@?>?@AB@@???@??@A@??@ABABDCABCDCCCCBDEECDECCDDEDCDCEEECCDEHοĴȱr:@BBDEEEEFHGDDDDFFFGFFGFHHHD>kG*5>CDEDDDFGGFEEGHHHGGGFFFGEGHIHJIHGIHHJKKJIKKHHJJIHIIJIIIJIJIJKJIHIIHIIHJJHGFGHIGGFFGGFGGGGHFFFFFFFEDEGHHG<=????=>?>?=>?<==>@BA@??@A?@BCCDDCBBACDBBBCCBCECEEECBBDBBEEEDEFECKúŻЭz>?===><<=>>><>?@@@@?>?A@A@CCBAACEBBEDCCDECCCEDFEECCDEDEFEEFEEBPɾξݼÇ>ADEDDEEGFDEFGFEEEEFEEGFFFFGGE>?kP'+:BDFFFFGGGHGFGFGGFFGGHHGGIHHIHJIIIIIHIIJIIKIIJKKIJJJIHFFIIHHHIHGIKJIIJJIHIMJFGHGGFGGGGFHIGGHJHGGHGFGIHGGEFD@===>?>>>???>=@??@A??@BA?@??@ABBAABCDEDDCEDEDCDDCCEECCEGFDCCDDDECSŽȷŲÎ@>CDDEFDFGFFFEFCFEDFFEGGEFHGHHE?8EqūU(&1?EGEFHGGGGEFGHHGHGGFEFHFFGHHIJHIHHJLJIKJHHIKJJIGHIJJJIIIIIJIHGHIHIJIIIIJIKJMMIIIGFFGHGFGIIHIHHIGFGFGGGHFEGFD?>=>=<<=>?>>@??>???@@ABA@AAAABCBBCBCBDEDDCCDBCDDECCCCCDDEDBCDCDECVž¶µ?@DCEEEEEFFGFGFFFGGFGGHFGGHGHJJHE>6Eoķq@"+9BFGGFGGGGHHEDEEGFHHGFFFGHHHGFHGHGHHIIJIJJKIIIHIKIIJIIHIKLJJJJKJJIIHIJHGHIIIHIHIIGHGHGGHHHHFGHGGGFGGFFHGHFEFHH=<=>><=><=>>>?=>?@@@>>>@@A@ABBB@BAACCBDDDCCCCDDFEEDDDDCDDCEDDFFGCWƼƽƷ>?CDDDEEDDDGGGGFFGHGGHGFGGIHGIIJJFB9/1GhqS1-9BFGFEFEFFFGGGGGFFFGGGGGHHGGGGGHJIIIJIIHHIIJLLKJJJIJJIJJJLKJJJLLLJJJJJJIHHHHHIHIIHHGHGGHHGGFEEGIHIGGIGGGHGFFFGGH=<<==>=><=@>=<=??@??>?AAA?@@@?AAA@ABCBCBCEEEDDECCCDEDDEEBCEEEFGGCUþ¬{<@DECCCDDDEFGGFGFGGFGHGHGFFGGGGHIGFEB8-()0=K[jvywzvgR?. $0=DEEFHGFFGHHGGFFHIHHIHFEGIIGIGEFIIJJJJIJIIJJJIHHHGHHFGHIKKIJJIIIIKKJJJKKIHGHIHIIHHHGHHHGHIIHGFEFGHHIIHGHHGHGFHHGGG<;<===>?===>==>>??@A@@A??@A@@@@BA@BCDCABEDEEEEDCCEDDDEIIDDFFEEEGDTôǽx9?BCCCDDEFGHHGGHGEEFFGEGFFFFEFHIJIHGGFD=6.(%!! !(09?DFGFGGIIIHEFGGFGHGGGHIGHIHIIHIIGIKKHIJHIJIIKMQUZ`aba]VQLIFGJIIJIIIHIJJJJHHJKKIHIIGGGFFHHGGHFGHHGHIIGFGGEHHHGGGIGFGG=<<<>>>>>>>?>>@>>?@>?@A??@@@AAABA@@@BBAADFDBDDCCDDFEDDFGDCDDDEDEBOʽϿq6>ACDDEFGGFGFEEEEEEFEDDFFHGHGHHIGGGHFFHHHD@=874423358<@DGHHGFFEFGIHFGFFHHHHGJMIGGFIIHGGHHIIJKKIIKKJIN[k{yaNFDFHIKKJKKIIHHJKLJHHHHFGIHGGHHGGGHIHHHJJHHHFFFHHGGIHGHIH=>>==>>=<>?@>>@?@A??A@@@@@@@@@?AAAA@@ADDCDDDCCCDDEEEDCCDEDCCEGFDDHòųϼ`8@DEEDEDFEFGFFDEEGGIHFGGFFFGHGFEFEFIHHGIJHIGHHFFFGGFHHJJIHGGGHHHGHFFGGHHGHHGINKEHHHGGGFFGHIIHHJJJKTi{WCEJKJJKKKJIIIIJGHIHHIIIHHGHIIGGIIIIGHGFHIGHIGGHGGGIHIH=>>>??>><=>@A@@@AA???>>?@?AA?@BCBABBACDDCBCDEEDDDECDDDEEDCCCDEFGEDʬȹS7ADDDDECDEDEFEEGGFFHGHGHGGGHHHHFFFHHHHGHHHJHIIHGGIJIHHHGFGFGHIIHGJIHHGGGFHFGGGGGIIGHIIGFGGHHHIHIVqŹ_EFHIIHKKIJJKIHIJJJIHHIJJGGIHGFFIJIGFFFGHIJHGIIGGHGGG>>>>?>>>?A@?A?@>@@?>@?@@??AABBBCA@@BBDCCBDCCEFFFECCDEEEDDFEDEDEEEBqʵǼ̾·B9ABCDEEEFFFFEEGFFGGGFFFFFEFGHHGFGGGHGGKHEGGHGGGGGIIIHHGGFGHHHFGHIHHGHKJHHGGGGGIIGGIGGHHGIIHHJJMd־†SDHIIIKJIJKKIIIIJIIHJIIHGHHHGFFHHHHIGGHHIHHJIHHIIHI>?>==>>=>@A???>==>??A@AA@@A@AAA@ABCAABCDEEEDDDFGEDEDEDCEFFGEEEDEFB\̾źľr8>DCDDCEEEFFEEEFFGGGFHHFGFEFGGFGGGGFGFHJHFFHIHFGGGGHJIHIHGHIIIGHHIGFFGJHGHGHGGHKJIGHHFIJJJIJJKUwӹ߫gGGJJJJIIJKJJGGJJIIIIIIGFFHHIHHIIIIIHGGGGHJIIIIIGH>===?>>?>>????>@@?@A??@@AA???@BCCBBBBACDCCCCCEDDCEEEFECFEDGFEDFGGCL̔Ǹ̵Q5ADDCDEDEEDDDFGFFFGGGHGHFFHHGFGHGFFGIIGHHGGHHJHIHHHIIIGIJIGGIIIIHHGFFFGFFGHHGFFGFGFFIIIHHGIJJ[ĪyJFIJIHHHIKKIJIHHHJJHIHIHHHJJIJIIJHGHGFGJJJJHIIGH?>@?>>??>?AAB@?@@@???@BBBA?=?AACBBBBBBABCDEDDDBBEFDCDEFFFEDEEDEEEFDڦǫ>9@ACDFEEDFGEEEGGGGFEEFFGGEHHFFHGFFEHHIHHGFHIGIIIJJIIHGGHIHHHHHHHHGHHGGGHFFGHHIHFFHIIIHIIIHIK]¡ۉJFJIHGHIJKIJJIHIIJGFHJHHGHIHIKIIHGIIGIJIIJIHGGH>=???>>??@AA@@??@@@??AA@@A@@AAAABBDBABCDDCCEDCEGFFEDCDFFEDDEEEFEGG@jЙZ5>ABBCDDEEGHHGGFFFGGGIHGFFFGHFGGEEFGHIIHHGGHHIIIHKKJJHHIJJHGIIIHFGFIHHGHIGFFGILIGGHJJIGGIIIKa㷮ŽPGIJIJJJIIKLIIHHHHFHIHHGGHHJJIJJKKJIJHGGHIHHIG><<<>>=>>@A@?@??@ABB>??@??@BC@@@BCCAAAABBBBDDEFGFFEFEDCDDDDEEFDDFGEI堔87ACBBCCDEGFEDEGHGFFHGHGGFHHEFGHGFGGGGGGGGGIHIHGHJJIJJJJJJHGHHGGFHJIIIIHIGGFGIJJHFFHHIHGGHJIYϭõŽLGJJLKKJKJIHIHHHFGJJKIIIIIHGHHHIHIKJHIIIGGHGI<<===<?@?ABB@?ABA@A@A@@BBAACCBABCEEEFFGFEEFGFDDDEEFGGEEEGHH@jˍĹX/>CCCDDCEFGFEDDDGHGGEFFEFHGGFFHIGGFEEEGGHFHIIIFGHIJIIIJJJKJJHIHGGGHIHHGHHGIHGIIGHIIHGGHHIJIQ~ѽƼȽGIJIIJLLJHHIIIIGIJJIHIKIHIHIIHIIJKIHHIHHHGHH<=?@?>>>?>=>?@?@??@BA?AAAAABB@@AA@ACCDCBDEEFFEEFEDEFEGGFFFGFEEFGHHGEH񣏖ɾz36ACCEEDEFEDFGHGFFHGFFGFGGHHGHHHGFFGFFEFFHGGIGGGHIHIJIIHIJIKIIJJIGGGHHIIHHIKIHGGFGHIKHHIIIJKmѷĹɸmEIIJJJKJJIIIIIHJJIHGHJIHIJJKIJKKIJIHGHIIHHJ=?>>?@@@@AA?@AAA@?@?AA?ABA?AACCAACCABDCDCDEGGGEFEEDCDEFFGGGEEDGHHGIIB^ҏG0>BBCDDEFEEEGHHGFGHFGHHFHHGFGIGFEHGHHGGHGEEGFFHIHJJHHIIJKIJKIIJIHHJIIIJIGGIHHIIGGHHHKGHJIJI[ƱžUDJIIIIIJIJJJIIJJHIHGHHHJKJIJIJJHJKKJIHIIHJ?@?>=>?@@@@@?@A@BA@AAB?>@@?@ADDCCDDBDDCCDEEGFEDEGGFEFFEEGFEEEEEFGHGIFBþa,9ACEEEEEFFFFEEGFEFEGHGGGGGFGHIHFGIHHHGHHGFGHGHHHHHIHIKJKLJIILKHIIIHIKHGHIHIIHHJJHHIIHGGIJINz仦˷DJJJKJKKHHIJIHHIIHHIJIJJIGGIIIIIIIJIHGHHHH@AA@?>=>>=?????@@@AA@A?>>@A@ACDBABBBDCCCDDFEFEFEEFGHFGHGGFEFFDFGFGFGHEGy-2?BDEEEGGFFFFEEGGHGFHIHHHGGGGGHIHHHHHHGHIGFHHIIIHIGHJKIJKLKIIJIHHIJHHHIHGIHHHGHHIHGIJIIHJKJ_կǿľZFIJJJMKJJKJIGHIIIHJJJJJJHHJJIIIIHIHIIIIII=>????>???A?>???AABBABBA@ADCC@?@?@BCCDDDCCBCFGFEEDEFFFGGFFFFGFEEFFHIIHBYǿ7+=CDEEDFGGEEEFFFHHHFGGFFEGFEFFEFGGGGFHGHFGHFGGFGIHIJHJIGHHIJIIJIGHIIJIIJJHGGGGIIIIHIGHIIKLJM|˵ϽŻEHIHIJJHJKIHGJIJJIHIJIHJKIHHJKJJIIIIIJJII>>>>???@A@@>??A@?@BB@AAAADDBA??ABCCA@BDFDDCCDEEDEFFEEDDEDEEDDFFFGFGGGHJ@c~4(;CEEDDEEFGFFFHHGEFFGFEFFGHGHHFFHHHFGGFFGHGHHGHFGHFGIJHHHIIJJJLLKKKKKIKKJLJIHJIIJIIIIHFHHKLJ[˾UFGJJJJIIKJHHHHHIIJJKJHIJIHIIIIIIIGGHIHHI?@>?@?@@@?>??@BA?AB@BBAABBAA@ABCBBDDBBCDBBDDDEEFFGEEFGGEFGEEEEGHFEFFFGII>h}8#7ACEEEFFEGFGHFFEEDEFGGFFHHGGGIIFHIHGHHGFGIHHHGFGHGGFHIGHIKJIJKKJKMLKKKLKIJKJJKJIJIJJHHGHHJJInvFHJJIJJJJJKKJIIHIJJIIIIHIIIGGHHHGGHHGFHG@@A@ABA?@@AAB@ABBBA@BCC@AAA@BBBCCCDEFECBBDDEEEFDFDEGGGFEFHGGGEGHFFFGHHGIG?cᤄ}2!3>BCDEFFFGHHHGFDEGGGGFFGGFFFGGHLJFGGIHIHGGHGHFHFHIIHGIGGHJKJJIIIHHHIIKIIIHHJJIHHJIIJKJIIIIIHN¾JHKJJKKKKHHKKIHIJIHIIJJHGFGFHJIFFFHIGGGF>>A@A@?AAA@AAAAAA?@BA@ACCBCAABCCCDABDDCDCDDEEGFEDCFFFDDFEFGGGGEFGGGGHHGHIIBX᷉p)"5@CDEEEGGHHGGEEDDFFHHHGFEHGGHIIHHFFFHIGHHFGEGFEHFHIHGHIHHIJJIIIIJIHHJJJJIGIIIJJIIJHHJJHHIKJJIWͷXEKGHJIIIFGHHHHIJJHHIIHIHHHHIKIHGGHHHIII?@@@???AA@BCBBA@?@BBBABCBBBBBCBCBCBCCCDDECBCDEEGFDEECDEFEFFFFHHHIJIGHHHHHIHBK̙V"%7BGFFGFFFGHHHGGGGHFFGHHFFGIHGGIIIFFGHIHGGGFFGJHFHGHGHGKJIIJKIIKKLLLKJJIJJJHIIIJJKJJJIJIGHGIJIHakDHIIIKKJKIJHGHIIIGHJHGGHIJJIHGHJIHGHHGHA@@A@@@@??AA@BAA?AB@ABBCCAAABCCCCDCCCBDCCBBCDCDEGFDEDEGEDEDEFGGFEFFGGHFHJGHHDEuԷ|uu5(9BCEGFEEFHGFHGGIIHFGGEFGFGHGGGGGGGFHJIGFHHHHGFGHIJJIFHJKKJIIIHIIKLJKKKKKKKLJJHHJJKKJKIHJHIIGHHFjǼzDHKKKJJJJKJJIGHHIHHJHIIFHHIIGFGJIHFGFGG@@AA@A@A?@A@@AA@BAAACDDBAAACAABBDDCBAACDCCDDDEGFFFFFGHHEDFEFFFFFFEFHHGGJJIHIID=MɳzuyH0=CEEFGGGGGHGHGEEGHGGFFHFEHHHGIIGGGGFGIGFEGFFGGGFGHHJIIJKJIIKKHIKKJJIJIKLKKIHIIJIIIJIHHIIIIKJJIJHrɽ؂FFIIKJIIIIIIIJIIHIHIHJHHHGGHIHGIJHIHHGH?@A@?@@@@ABDB@A@AABBCCABB@CDDBCBDDBBABBDEEDDFGGGFFFHGGHFFGHIHGFFGHIHHGIJIJJIJIIA;Mtty{J'7@FGFGGFGHGFGEEGGGFGGGFFHGEFGGGHHFEFHGHHFGGGFEEEGGHGGFHJIHIHIIIJIKKKKKJIIJJIJJKKJJIGIIIIIJJKJKKKJGyۋIEIIJIIJKIGIHIIHGGHGHHHIHIGGHHHGHHHIIIH>????@AA@@ACBB@??ABAAA?ABABDCDCBEFEDCBAADFFEGGEFFEFFEFFGJIGGIHFFIIHGGFHIJHIHIJJHF?8Ddj;%4=CFGIHGGHHIHEFEEHIHHGFFHHFFFEDEFFFHHHFHIJHHHHHHFFGJJHHGGIHIHJJJKJIHHJJJIIJJIHJLLKIJKJJKKKJJJKJIJLIGĿٓLEIIJIHIJIIIIGHHGGIIIGGGGHGGHIIGGHHHIII?@>?@?@AAABCAABB@AABDCACDCCDCCCCCDDDCDDDEFEDDFFEDEGGFGGGGGFGHHGHFGGHGGGHJIJHGJJIHIF@76ARhiO5$*6?DFFGGHHHHGGGGFGGGHGHGGGHHHFGGEEFGGHIIHFEFGFFHHHGIIGHGIJHFFGIJJKJHHJPLIJJIJJKKKJJKJIIJIIJJIJKKKJIIJJH¹οӓLEHHHIIIIIJJIHGGFFFIIHHHGFGFEFGGGFHIHGIAA@AAA@@AAAAA@ABBCACCCBCCABBDCBDDCCEEDEDDEDDDEFFFHGGGHHGGFFHHHHHHGGHHGHHIGHHHHIIIIIIGA91..37<@FLID;.%&0:BFFGHHHFFGHHHGEFFGFFHGIGGGHGGFGGGGFGGHIHHHFGGEEHHHGHHGGHIHGHHIIIIIJHHKX_NJLKKKKLKIJLKJIIIHIJIKMLJIJKLLGzļΏHEHHHJJJHIGGGFGFFFGGFFGGHGIHFEIHGGFHHIH??AA???A@@>>@@?@AACDBBCAAABCCCDCBCEFEEDDCDDDFGHGGHHHGEGIHGGHHGGIIHGHFGHHIGHGIGGHHHIIIHFDA;72,)'$##"&+18>BFGGGGGGHGHIHFEGGHHHHFFGHHGGHGFGIHGGGGGHGFGGHGGIIIIHGFGHHGHGFGIJIHHJHIKKILPLIIKHIKKLKKKJJKJKJIIJHIIHHIKLKEuҾɅCDIIHJJHHGEGHGGGFGHHFFGIHHFGHHHIIHHGHGI?>>>???@@@@@@ABBABBCCABABCCDCDDBCEFFDEEDDDDEEFGGFFHIKGFGGHHHIIHHGHIJHHFHHHHGGIIGHHHHHHGIIIHEDDA?>=@ACEFHGHIIHGFHIHHGFGGFFHHHIHGFGGGIHGGGGFFEFFFFFEHHGHHIKJGGHHGGHGFGHIIIHHIIIJKLKJIJJJIIJKMLKIKJJJJJIIJKJHJIIJJJJEm|CDGHIHHHGGIJHIIHFFHGHHHJIIGGHHHIHHIGGGG@@??@@ABBAAAABCCBAACBBCCCCDDCCBBDEFEEEFFEFEEFHHHGHHIJJHFGGHGGIJHHJJIIIIHGHHIIHIHHIHHIJKJKJJJGGHGGFHHGHGHIIHIIHHIIIHFEGFDFECEHHHGIIHHGFGGFFFFEFFFFFHGFGGHGGIFFHHFHIHFHHHIIHIIJJJIJJLLKLKJJKLJKIKKKJJIHIJIIJJLKIJJJGaȹoDGGHIHHHHHHIIHIHGHHGGHHGGIIIIHHIGGHHHHG?ABBBAABCC@?ABA@@BABA@CAACCCBCCDEFECCDCCCDDFEGFGHHFGHHHIHHIGGHIJIIHGFHJIIHHIIIIJGFIIIKKJJIIJHFHHGIIHGGFHJHHHHIHFFFHLQUX[[YUPJHEDDEFFGGEGFHHHGFGHGFFGGGFFGGHGGGIHHHIIJIFGIHHGIJIILLJJKKJJKKJJIKJJIIIIIIHHIJLKJJKJLHSǫ\AGHIHHHGHGFIIGGGHHGGFFIGGGHGGHHGFEGHIGHA@BABAA??ABBA?@A@ACCCCBBBDCCDDDFDDDEDEDDDDDEFFGFHGGHGHHGGHIJJHHJIHHIHHJIHIIHHHHIIJIHHIIIKJJJJIHGHIHHGIIIHJJHGGIN[pzeRGBCDEFFEEEGGGFFFGGHHHHGIJGGHHIHGHJJIIGHHHHIIJKLKJIJKHHJKMLKIKKJIIKLJIIIKLKIGJIIKJGÖJ@GHHIJJGHIHIHHHHIGHFHGHHHFFHGGHGGGGHGGHA@????@@>>@AA>@A@BCCCCBCDCDEFDDCCEEFFEFFGGGGHHHFGGIGGGHHGHJKIJHIIJIJIHKKIIHGGHHHKKIHIJHHHFGGIIHGGIHHHHHHIJIHLZrɽlRDDFFEFEEFFGGFGIIHIHHIHHFGGGHHHIIHHHIIHJJIHJLJIIIIIIIHLLKJKKJKJJIJJIIJJJIIKLJLMD½z?CHHIGHIHIJIHHGGGHFHHHHHFHIGGHHIIIHIHGGH@A?@?@@ABA?@@A@@BCCCAACDCDBDEECCEDEEEEEEFEFEGFFGFFHHGHIIIHIIFHIIIIIJHGHHJJIIHIJIHIJHIJHIIIHIIIIHGJHHIIGGGIOb§|YGCDGGEFFFGGFHIGHIHHGFFFHGGJIGHHIIIJJJJKIIJJIJKIJJIIKLLKJKKLLJJKKHHIJIJJLJLMLH\»`>FFGHHIJJHHIIHFGHFGGGGIJGIKIGFIHHHGHGGFGAABBAAACCBABAAAACCDFDCEFEBCDDEEEEDDDDGHGGGEDFFGGHGFGFGHGHGFFHIIHIKIHHIIHIHHJJHHHGGIJJIJKIIKJJKJHHIIIHIHHM`}ϴUBCFEGFDFHGFHIGGHIHFGIJIJJIGFHHGHIJIIIJJKLKIJIHHKLLKKKLLJJLLMLJJIJIHJIIHJKKLHºGBFFGGHIIIIHGIIIHGHHHGGHJIHJIIHIIGHHFFFIHA@AAABBCBA@BBABCDDDCEECEFEDFFEFEDDCDFGGEGHGFFGGFFFEEGGGFGGEGIKIGGJHFHHHHIGIIIHHIIHIHHJJKJHIJIIHHHKJJHGJ]{ʾȹ̠kICFFEDFGGFHHFGHIIFGHHHIGHIHHIGHHHHHHIKKLKIJKJIJLLLIJIJJJKKJLKKKJIJLLJIJJJLE^һɾc=DGHGHHGHIHHHIHHHGHIHGHFGIHIIJJIIHHHHHFIIB??@@AAA@@A@ABBBBCBACDDEFGEEDEEFEEFGFEEFFFGGGHHEEFGHHGHHHIHHHIGGFGGGHIIGGIKHFIIIIKKIJIHHJJIHGGGGHIHIHOmōSDEFDDDEGGGIJJIIJIIGGHGGGIHIIHGGGIJJJIHJKKLLLJJIIJJKLLKMKJKLMKIKLKLLLLKLLLFʿƾA@GGHHGIHGHHGGHHGHFFGIIJHHHIIHHHGHJHIIHGGF@?AABA?@AA@?ABCBCCCDEEFDDDEFEGGGGFFFFFFEFEEGFFFEFGFHHGGFFIIGEFGHHHHJGGHHGHHHHIHHIKJIIJHHHJJJJIIIIHGIVzŽЧaDDFEDFHIGHIHHHJIHIIHIIIIHHHJHHIJIJKLIIJIIJLKHHIJLJJJJKJIJKKKJKKKJKKKKLMMHPůľȻY8EHGGGHJIIHHHHHGGFGGHHHHGGHIIHHGGIJJJHGHGGABBBBB@@BBBBACEDCDDDFFFEEDEEFFFFEFFFDDEGFGFGGFFGHGGGHGFFGGGGGFFGFGHHHGHIIIHIIHIHIJLJHIKLHIJKKJHIJII[ȠĿضkEEHGFFFHHIIGHIHIIIIGJKKIHHJIIJJKJKLKJKJJIJKJIJJKKJIKJKKIHHIIJIIJKKJKLMMMFgĹqjƿعw>3BFHHHIGGFGIJGHIHIIHIHFGHIGFFFFGIIHGEFGHHHGHIIIBDCCBACBBCDCBBCCBCDDEFEDDEEFFHHGHHHFFGGGHIHFFGHFGGFGGFGIJIHHGHGGFHHGHHFGHIIJIIHHJJJIJLKMKJKMLKc¬¶ͻxEEFGHHHIIHGJKKIIHHGHJIJHIJIJKKKKKLLJKMLKJIJKKJIIIHIJJKKKJKKILKJJJKKKLKLMG:LĻf94CGGGHIKHHHHIIHHHIJIHHHFFFIIGGFFGGIHGDFGHIIHJJKICDA@@@@@ACFDCBDDDDEBBDFGFDEEHHHGFFFGGFFFFIHFGHGHHIFFGHHHHGHGHIHFHIGFGGHHJHHHHIIJJJIKJJKMMKJKHS½̯԰[BGHGFFHIHFHHIJJKJGHKKIIKKLLLLLJIJKKLKJJIKKKLLKJIILKKJJLJIJJKLJKKLLLKLLKMI<GIHIHJKJIIHIHHHHHGIJHHHIHGGHHGHGGGGEFGHFHIJIHHHII@ADDCCCABBBAABDDDEDCEFGEGFFFHGFHGGGEEGIIHGGHFFFFGHHHHHGGHHGIHIIIGFGGIIGGGIHIJHGIIIJIJJLKKMLKYϴ¯˶ξcCHHGHHHHHJKKIJJJIJHIJJIIJIJLKKKLLJIJLLKILKJJJJLLKLLKKLLKLLKJIKMKKLMMLMLMNNJ>29Pt½hF0-7DJKJIIGHJIHHHIHHHHHGIIHIIIIHIHIJIGGFDEGIIHHIKKIIIIGABDCDCDDBBBCBCEDEECDFFFEDEFGHGFHGHFFHIIIIGFIHGGFGIIIGGFFIIHHHIGGGGFFHIIJKJGGHIHJJKKIKKLLKKLMy¹¼ôԎIFGFGHHIJIJKKJIHJIHJKKMLLKJMMLKMLLLKJJKIJJKKKIIKKLMKJKJKKKKKKKMMNMMMLLMLLMNMJ@503>Thv{n^L<0+1:76666:@EIKLKKJKKJKKIIIIIGFGHIIIHHFGHHIHHIIIIHGHGFIIHFGGFEGHGGIHFGICACBBAABBABBBEECCDFEFFGFGFFFFHFGGGHHIIFFGHGGHHGHHHFGIHIHHHGHHIIHFEEFHIJIIIIHGJJKIHIJJKHJKIOƼǽDZDzȹ؟MBGGFIJJJJJKIIHIKJJKJJHHGJLJJIIKMLKIIHKKIJIIKKLJJLLLLJJIIIJLKKLLMMKKKKKKMKJLLJJJJLMMLLLLLJIJJLMLKIJJJIKKHIJJJJHHIIJIIHHIHHHJJHHJJIIGHGGGFDFJIGHHFGIIIGIDDCCCCCCBBBCDDEFEEFFFFFEEHGGGGFGGFHHHJJIFFGHIIGGGHJJHHHHHJIHHHGHHHIIHHHIKIHJJHIJJHHJLJJKJG\żøĮǻɽչ]AGHGHHJHHJJLJIJKIJIHHIIIJJJHIJKLMKJJKJJKJHIKKKJKLKJJJIJLJJIJJKMMMLLMLMLLKKLMLMKJMMLLLKJKLKJJKLKKKIIIIIHHIJIJIGGHIHIJIHHIIIJJGGHIHHHGGHGFEGIHHFHIIHHHIHBCDCEEFEDCCCDCBCEEEFEFGDCEFGGGHGIIHHHGHHGFHHGFEFGHHHHHHIGHHIHGIIHJIJIGHGHHHJJJIJKKIJLKLKJJkĻʾ÷­ɸɾn=GHHHJJIHIJJJKJKKHGGHIJIIJJIKLKJLKLJKKJKKJIJJLJKKLLKJKKJKKKKKKKLNMLMLKKKKMMLLNMLKKJKKJJKLJIJIJKKKIJIJIIKKJHIHGHHHIHHJIHHHGHIGGGHHGHHHHGFFGGHGEGIIGFHHGBBDCCBCDEDBBBCCCCDDFGGHGGEGHIIIGFHIIIHIGHHIHGGGIJIFFHJIJIGHJIHHHGGHHHIJJJIIIHJIHHIKKIJLLLN}õÿǾф?DGGHHGGJJHIJJJJJHILKJIJKKLLMMJIKKLKKKKKIKKLMKMLJLLLLLJIKKKMMNLLLKMMJKKKKLMJKLMMLKKKJKILMKKKKKLJIKKKJIIIIHHJHGGIJJGGHHGHIIIHHHGHIGGHHIHGFGFGHGHHHIGFHHCCBCBBBACCCCDDEFEDDFFEEFFGFHHHGHGFFGIIHHGIHHIIGGIHHHGGGGIIHIHGHGFHHHHIIIJKIGHHIKJJIJJJKLKTл¿ʽȸŶѐDEHGHHHIJIHIJJJJHIIJJJJJIKMNNMLJJIILKIJIJIJJKKLMLKKKKJJIJJHIKNONLLLLJKLLMNLKLMMLLLMLKJJLKKJKJJLJIKKJJKKIHHJIIIIHIIHHFGHFGIIGHGGIIIHGGGFEEFGHHHIHHHHGGGCBABCEFCBBCCDDEEFGGFFFFEFGFFFFFGGHHHHHHHGHHIIHHGGFHIIGFGHHHGHGGHIIIIHHIHHIIIIJJJKLLJIJJKIWïƾҼξóĸӚHDGGHHHHIIKJHIJIIIIIIHJIIJKJJKLKJHIJJIIJLJHIHIJJJLKKKIJJIJJJJJLKKLLJKKLLMMLMNMLMKLLKKJKKJIJKJJKLKJLKLKJJJJLJHHJIGHIIGHHGGFGFFGIHIJIGGGFDDEFHGFGHIGHHGGCDBDDDCDCBACDFFEEGGFFFGGHIHHGGGGFHJIIFGGIIIGFGHIHIGHHHJIGHFGHHHIJJIHIIIIJIJKKKIHIIJJIKLLH^ƶʬǾпҫPAGHHIIHJIJKJIJIIHHIHIKKLLJHIJKIJKKIIJIIKLKKKIHIIIKJLLJJIJLKKLJJIKMKLLKKKLLKMNMKLMJKKLLLLKIJKJIKJJIJJJIIJIJJIIIJHHJJIHHIHFGHFFGIKKIIHEGGFHHGFGGGFGGHHGDEDDDBBDCCCEDEEFGGEFEEFHHGHHIGHHGGHGHGHHGGGGGGGGGHIIHGIIHGGHGIIHHHIJJJJKIIIIIKJIIJJKKKLLG`ĺ¡¿аWAHHHIIIIJIIIHIIKKJJJJKKKJLKJJJJKKJKJIJJJKLMMKKJKJJJLKKLKIIJJKJJJJLKJLLLLLJLKMLLKLLKKKKKLKJJIIJIHIIIJJJJIHGHHHIJJHIIHGFHGGHHFGGHIJJIIGGGFGHHGHIIFEHGHGCDCCDEFCBBDDFEFFFFGGFEFGGGGHGHJGGHHGHGHIFGGFHIHHHGHHIHGFGFFGGHIHHFGJKIIJKJKJHIJKJILMLLKKGe̾ƼģаWAHHHIHHIIJHIHIJKJJIJJIJIHJLKKKKKIJJJIIJJKJJLLKLLKJJKJKKKKJJJKKKLJKKKLMLLKKKLKMLLLLKJJLKKKKIHIJIHIJJIKKIIJHHHHHIJIHGGFFFHGHFFGGGIHGHIHGGFFGGGFGHGHGFEFCBBCEEECABBCFFHGHFEFEDEGFHIIHIIHJJJIJJHJIHHHGHGHIGGGJIGHFGHIJGIIJIIJIHHIKKJKJHIJIHIKJJMMGoƨøýǦƿϰWAHGGHJJHHIKIKIILKKKIIIKLLLKLKLKJJKLKKLKLLLIIJLLKLJKLKKIKKJJKMLKKJJJJJJKKJKKJJLMLMMLKLLLKIIKKKJIHJKKIKLIILLHHHIJKHGGHIIHGHHHFHIIIGFGHIGGHEFGFFGHHIGFFFCCDDDDDDBCDDDGGFEFGFFEFHHHGIIGHILLIHIJIIIHGFEGIIIHHIGGGGIIIHIIHGHJIJKJIJHHHIKJJKKIJKKKLMGiäϸƻĿ̘ɻάT?EFGIIJIIJKKJJKKKKKJJIKMNLLKLLKLMKLLJJKLLKJJJLJIJKKKKKKJIJKKMLLLKKIIHILLKJKIKLKKLKJKKJKMJJKKJIIIHHJJJJKJKLKIHIHJIIIGHGGHIHHIHGGFFHGGHFFHFFGFGHGFFGGHGCBDEEEFDCCEEDEEDFHHHHGFHIHGHHFHJIIIIGGHHIIHGGHIKJIIIHGEFIIHFGJGGGIHIJKJKKJHIJJIKKKKKKIKLIbǿ̩ĺΧP@GIIJJJKJKJJJJJJKKJJKKJJKKJJIKKLLLKJJKKJJIJLLKKKJKJJKLJJJLLKKLLKMKIJJIKLJKKKJKJIJKIJJJJKJKLJJJHGIIJJHHJJJIIIIHHIJKIGFEGIIIHHGEFFHHGGFFFFFGGEFGFGFEFHHCDCBDFGEEDDEEFFHHGGIIIGGGFIIIHHHHIKJGGGGIJIJIIGJIHHGHIHIIJIGIJIGGJJIIKJJLKKKJHIJLLJIIIJKI^Ʀ˲ƽʿɚIDIJJHHJJKJKJLKJKJKKIJKIIJIHKJKLKJLJIJKLKIIJIJIJKJIJLLKJIJMMLKJKJLKKKKKKKKKJJJJJIIJKJLLIJKJLKKIKJHJKKKIKJHGGIIIHGIKJGFHHHIGEFGHIHGHHGFEFEFFHGFHGHFEEFHEDCDGFFFFCCDEFHJGEEHGFFGIIJIHIIGIJJIIIIHHGHHGHGHHIIHGHIJJIHIIHIIJJJJJIJKKJIKKKKJKKJJJKKKHXΨͻǴƼ͊>EGHJIHHIKKKKMLKJKJIJJJJMLKKHIJKKKKLLLJILLKLJJKJJJKMMKLLKKLLKKJJJJMLLKJMKJKKJJIKKKJKKLJJJKKKJKJKKJIKJLKIHGGHJIIIJIIJHHHIHGGGGHHIHGGHHGGGFFGHGGHGFFFEHIBBDDEDEFEDFDBEFGEEFGHGIIIJIGGIIIIJHIIHKIGGGHGFFHHHHHIHHHHGIJJIILKIJJIHJJJJJKJJKKKIJJKLKKHRձȺƴûğx:EHHHJJHHHIJKJKKJJKJMLJKLLLLJKLJJLJJLMLKJKMLKKLLLLMMLLLLKLJHHILLKLLLJJJLMKIIIKKKKLKLKKKKJIJIIJKIJJIIIKKIHJIJIGIIKJHJIHGHHGHGGHHHGHFHIHGGGFGGHHFFFGFGIIBCEDCDEDCFGFFGFGFEGGFGIHGHHHHJKKKIHHIHHIJIIIHGGFIIIHIHGGGHJHIJJIIHIJIJKJJJIIIIJKIKKKKLNMKL趗ʺʸžȻa:EHJKJIJKJHIIIHHJJLKJJIJJKLLLLLKKLJJLMMMLKKLKLNMMJJJJIGJKKKIJKKKJLJJJJKKJHIKIJJLKKJJKLLJIJKLHHJIIIKHHIKIIIIIJHIIIIHJJJGHIIGFGFGGGGFGIJIHHGFHIIHFFHIIJIDEFFEGGFFGHHHGGHGHGGGGHJJJILKJJIJJJIHIHIKJGFGGHGHGHJKJHHHGHIJIHHIIJIIIKLKLLJIJIKJJKKKJLMLF̤Į̾ƿħO>GIIJIHIKKJJHHJJKMKKJIIKLMKKJKKJLMLKLLLLMLKLNMLKLJJIKJHGIJKJKLKLKLKKKKJKJJLLKJKMLKJKKJJKKKKJJIIIJJJKJHIIHGHHJJIIHHIIHJIHHGIGGHIHFGHHHHIHHGFGHIHFEHIJJHBDDEGGHGGFFFEFGIIIHGHHJJIIJJJIIHIJJIHJJJKJHFHHHIIGGHJJIIJHIIIHIIILKKJIJJLNKIHJJIJHIJKKHJLEo⳦Ķ͵Ŀ?BJIHIIHIJJJPQJJLJKJLMJIKLLJJJKKLKKLKMMKKKKLLMMMLKKKJIIKKKLLKLLLLJKKJKKJKLLKJJLLLKJJJJJJLMJIHIKJHIJIKKIIIIHHIIKJIIIIGGIKJHGHGFIIHFGGGHHHHGGHFHIGFFIIIIGDDDCEEEEDEEDDGIIHIHHHJJKKIIHHJJIIHIJIKJHIIJIJJHJJHIGGHIHKKKIGHHJJKKKKJKJKKJJJKKIJKJKLLJKLIVﵡǿʿ½`8EIIJJIIGIJKNPKJKKIJJJJJIKKKKLJJKJJKKLMLLLLMMJJMMJJKKIIJKKLLMMKKLJKLKKKJKLLJJJKLKIKKKKKIJJIIKIIIIJJHHJIHHHJJIHIJJHHHHHJJHHGGGFHHHGGFHHIIIEGHHGHHHHHHGGGDDDDCEFEEFGFGHGFGHIGHJJKJKKJIIJJIGGKLIIIHIKJIHHJLKIIIJKJJJJIHHJKJJIJJJKHIJIJLKKKMLLKKMMMLJGȞ¾ɥB>IHHIJHGGJLKKKJLKIKLIIKLKMMLLKJKKJKJLKKLKJKLKKILKJJLLJKJKKKKLLJIIKKKJJKIJJKKLLKIJJHJKKKIIJLJJJKKJKIHHHIIIIJKHGIIJJGGIIHIIHIHGGHIIGGGHHIJHHGGGGIHIIJJHFGDFFFGHIJHGFGGHFGGIJIKLIHHJKJJJIJHHIKKIJJKJKJKKJIJKHHJIIKJIGGHJIHIIIJJJJHHJIIKJJIKKIJKKKMNMGi鮣µĻ͟ül6CJHFGIHGJKKKKLJKKLKKLKMLKLMKKKKKLKKMLKKKJIJKKLLLKKMLKLLKJJLMMMMJKMJJJIIJJKLKMMJJKJIKKLKLKLKIIKLJJIIIJHIKJKIJIIJHIJIJJKHGIHIIHHHHHHHGIIHJIHGHIKJHHHHHGGFEDDGHHHHHGFFFGIIIJIIJJHHIIIJIIJIKJHJIJJIIHIJIJKIIIJJKJIKJJIIHHIJIKKJJIJJJIJJKJIIKKLLKJJJMMJNÿ̟B9FHIHHIJJKIIKJJJKKKJKKJLKLLKKJKJMMKJLMMLJKLLLKMNLMNNLMNNNMMMMNMMLLLKJIHJJKKKKMLJKKKKKLLLKJIIKIJJJJHHJJIJIJJJHJJIJIGHIHIJHHIJIGIHGHHGGHGGGGGHIIIJFFGHGHIGEFDEFEGFFFEGGHIIHGGHGGIJJJIHJKJIKJIHIKIIJIJLIHHKJJIJKJKJHIIKKHIHHJIGHKKJJJJIIKLLMMMLLKLKKMMGqƾǽħ_2BGGGIJJJJHHHIJKMLKJKKJKLLMKKJJJJLLLKLKKJJJLKIJLMKLMMMLMNMNLLLKKLKKJKKKKKJKJIIKJKKIJJJKKIJJIIKLJIIHIIIJIIIHIKHIJJJIHIIIIJHHIIIHIJIHHGGIHFHIGGIIIIFEEFGGHFDFEDDFFEEFFFFHIGHGGHIIIKKJJIJIHJLLIIJJJIHJKJJHHKKJJIIJKKIIJIJIIHJIGGGIKJKJJKJJLKJLLLLLLLKKMKLȼǻmh77GHHJJIHIIHJJJKKLMLKKKKLLLLKKKKKKLKLLLKKJKJKLKLLKIJKKLLJKJJJILLKLKKJJKLLLKKJIJKKKKKKJJKLIIKLLKLKKKIIJJKIIKJJJJJJKJKIJIHHIHHHIIIHJIGIJHHHHIKJGFHJIIGFFGHHGEDDDFGEEFGGEEGHGIIHGIJIJKKJLKHIJKLKKJHIJJIHHIIIIHIKJIIIJJIJKLLKKJIHHJJJIJIKLKKJJJJJLKJJLMMKNGg駣vi|§N.AHIHHIIHJJKJJJJJJJKLJJLLKKJIJLLLMMJJLMMLMLMKKLLLLJKJLKKMKKJHHJKLMJJKJJKKJIKLJKKJKLKKKJLLJJJKLJJKKJJIIJJJJIIIHIKIJKKIIIHGIIIIGHIIGHGIIHGFHIIKIFFGHIHGFFGGHEDEFEFFHGGGEFHGFIGGHHJKJJKJKKKLMJJJKHHJKKIHKKIHJIJJLLJHJIHIKMKKLIIIJKKKKKJKKKJIIJKJKKMMKLKMMLF֝Ǿh+GHHIIJJHKLKIHHJIJKLJJLIIKKJIJJJKLLKJILMLKJJJMLLKKKNMKKLLKKLLJKLKLLLKJJKMJIHHKJKKJJKKJIJLKKMLKLJIJLLLKJKKKKJKJIIHGHHJJKJIIHGGFFHGHHJHGIIIHHGGEHHIGGGHGFHIGHGEEGHGGHFGGEEHHGIJIIHIJHIKKJIJIIJIIIKLLJJKJLMKJKIJIJLMMLLKKKJIHIJKJIHGIIJJIIJKKJIJKKJKKKJIIJLLLMLGpﭔJ(:GIJJJJJJJJJLJHIIKJJKHIJKLLMKJLLLKLLLLKKKLKKLLMMKJKLOLLLKKLLLLKLLKMMLKLOMJIJJJKJIJKJKKIJKLJKKJJKJJJKKJJKIIJJIIJIIIHHHJKLIIHGHJHIIHGIIIHIIHIIGIGGHHHGFHHGGGGHHFEHIGGGFGGHGHIIJKIHHHJILLJJKKIJJJJKKJJIJKHIMKIIKKKIJKJKKKJKKJIHHJKJJJJHJKKIKLJKKKKKJKJJJLJJLKLNNLEn綞C$6EHHIIJLKJKKKJJJKKMLKIJJILLKLLLLMLJJKLLJIIJKKLLKLKMLNNMLMMLMKLMKKKKJJKMONLJJKLKJJHIKLJKJJKKIJLJIJIKLJIIIIJIIIHHJJIHIIIJJJKJHHJJIHIIHIHFHHIHIJIJHFHIGHGFGHHGHHHEGGFFFEEFGFHGHIHHHHHHHJJJKHJLLMJJKKKJKLLKJIHHIJJIJIJKJIIILLKKJKIHIIHJIIJKLKLKKKKJIKKJIKKLLJKMMLNMLEZš6"8EGIHJKJKJKJIJIHKLKLKJKKKJJJJKKLIJKKIIIJJIKIIIKKKJLLKLLNMLLKLJKLLLLLJKKLNOMNMLMLJJKKJKKJLLJJKKKJJKKKKJIJIIHIKIIIHIJKJHJIGGIJIHHGFGGGFGGGFGHGIIHHGGHIHGGGHGFHGGHEGGGGFFFEGGHHHHIIGHIHIJHHJIJKKKJJIHIIKKJJKLKIKLKJKJJJJJKLLMKIIHHHGHJJHJJJLKKJJJKIIKKIJKKJLKLMLKLLLKFLΫo-$8EHHIJJKLJIIHIJJJKKJIJJJIKKJIJJJKKKJKKKJKKHKKIKMKKKLLKKKLLLJKLKKKLLKLLMNMMKKKMMMMLJKLKKKKLLKJKKIJIJLLJKKIJKIIIIIHHIJJJIJIHGGHHHGGGGFGGHGHGHHIIHIIHGHGIHGHGFHHGHHFFFGGGEFHFGHGFFGIIHHHIJGGIIIKIJJKJKKIJJKKJJLKJJLLJKIJKKKJIJJJKJJJIHIKJKIIJIJJIILKKLLJJKKLLLLMKKJJKKLFEsܿM!'66D_zvP0&7BHJKJKLKLLKKKLLJKKJJKKJIIJJIJKJJJJJJIKLJLKJLLJJKKJJJJLMJJMMLMNNMMKLJKLLKKLLMLKIJLLLLKKLLKKKKJJLJJKJJJJIJJHIIIIIIJJHGHIHIHHIHHHJIIJJHHGIKIIJJIIGGGGGJJJHIJIHHHGHJIGGHHHHFGFFGGFGGFGGGHGHIIIGGIIIIIKIIJLLJIIJKLLKJIJKJHIJIIIIJIIIJKIJLJJJIJJKIIIJJIIIIJJJJLKJIILKJKKKKMKJJLLKMLKLMLLKGA7/-08AJSY^[RE9.$$/;EKKJKJKLLJKJKLLLLKIJJJJJJKLJJIJIJJJJKKIKKIJJJJJJKKJKKKJIJJKLLLLMOOKIKKKLKLKLJJLLKKLMKKKIIKKJKIKJKKKIJKJJJJKIIJIKLLJHHIJHHIIIIHHIHHHHHHIHIIJIJJJHHGHHHIIIHIIHIIHHIJIIGGHHFEFGEEGGGGEEGGFFFHIHFGGGGGGKLJJJLIJJJJLJIIJIJLKJJKKIHIIIIIJMLLIHJIJKKIHIIJJJJKJJJLMKLLKKJKKKKLMKIJMLLMNKKKKKKKMKHC<4.)'%#"!"%+09>DGKKLKKMLJKJJJJJKKJJJIJLIIHIKLKJJKKJKKIJKJJJKHHIJIIKJJJLLJHILMNLKKLMMKKKLLMLKLLJJNNLKLKJJMJJLLKJKLLJKKHILKIHIIJJJJHIKIHIGIJGGKKIGHHGHJIHIGIJIJJHHIHJHHGGGHHHHHIIIHJJIHFGFEEEFEGFFEFFFHGHGFGGGIIGHHHGHHIKJLJLKIIIIJHHJKJIKKIJJKKJJJJKJKKKKKIJIJJIJKKJIHIKLKKKKKIJKKLJKJJJJKNLLLKKLKLLLMLKLNNOPOKHFEDCBDCEHIJJKJKKJJLKKKKKKJLKHIJKLKLLKKJJIJJIIHGIJKKJJKIKKLKIIKJIKLKKLKJLLMNNKKMMLMNMKJKLKJIKLNLLLKLKKKKLMKKKKLLKJJJIIJKJJIJIIHGHJJIHHIJIHIJIIIIHIIJJHHHIKJJHHIJIGGHGIHGGIIGHIIIIGGFEGFEFGGGFHFFFGGGFGGGHIGGIHGGHIHIHHJJJJIGIJKJIIJJKKJIKKJJJIJKLJJJIJKIIIJJHIIIIIIJJKJJJJJJJJJKKKIKMLKLKKMLJKLMLMMMMMNMNLNNMNMONLLKKKLLJLMKLLKKJKKKLKKJIJJKKKJJIKJKJJJIIJHHIJJKKLJKKJKKJKKIJJLJJKLLMLMLNMLLLLMMMLLKMLJKKLKLLLKMLLKKKKKLKKKLKKKKKKKKJIIIIIIHHIIHIIJJIHHIIIKJJJIHJHJIIJJIHIIJHGGIHIHHHIHGGHIHIHHGGHFEFFGHGEGFGGGHGHGFGHHHIIHIHHGJIIIIHIIHJLLLKKIKLJIJJKKIKJJIIIIJIJIIIIHHJHIIGIIJIHJIIIIJIIKKKLIJONKKIKMLLLLLKKKMLKKKKKLNNMMNMLLLLMMLKMMKJKJKJLJJJKJIKKJKJIKJIIHJJIHIIJKKKJJJJJKKKKIIKMLKKJKJJKLLLLLMNLMLKKMLJLLKLLLLLKJKMLJKKLMLKMMKKLMKLLKIJLKKIIHIKKIJJHGGHHJJJJIIJJIHIHGHJJJIHHHGGIIIHHGGHGGGHHHJIIGHHHHHIIFDEFFFFEFEEFGGFFFFGGGGFGGJJHHHHIJIIJJJJKKKKKJKKIKKKJKKKIKJHJKKJKJJKHHHHIJJIJKIIHIIJKIJJKJJIIJJMLJKLMMKKLMMLKJKJIKKLMMLLLJLMLLLKKLLKKLIJKLJKLKLLLKJKJJKKMONIIKIHHJKJIJKKLLKKKJKJKKKKLKJJJJKKLKLLLLLLMMLLMLMMLLLJIKKIKLMMKLKLKLLLMKLMKNMLKJJJKJIIHHGJJJJIHHIIIJGHIHHIIHHJJIIIIIHHIHGGHJJIHHFFGHHJIHIIHGHHIIGHHGDEGFFFFFEFGFDDEDDFGHGHHFHIHHGGGGILLKJJJIJJJIJKHKKJJKKJIJJIIIJIIIHIIIIHHIHHIKJKLKJIJIIJKKKIJKKLKJKMNKJJIKLKJKKJILLLMMNLMKLKKLKJJJJJJKJJJLKLKIIKLJJJIJIKKKLJKKJIIKLIJJJLLKJIIJJJIJJKJIJJIJKKLJJLNMMMLLMKKJMNLLMKJIIJMLKLLLLLKKJJKKLKKLKIKKJKLJJKJJHHIKKIJIHIIHGGGHIKIIHHKJJIIHFGIHHIHIGIIIHGHHIJJJJHHHGHGHGFHFEFEGGGFGFGFGGFGGGGHGHIIIFFGHHGHHJHJKKKJIIJIKKJIKLKKJIIKJJIHIIKIIJKJKLJHHHJJHIJIIIIIJIGIJKKKLJIKLLKLKKJHKKKMKIMLKKLMJLMLMMMLLIKLLLKKLLKKJLMLKKKIIKHHIIJIJKMKJKJKKIHKKKKKJKIIILKKLJLKIKJJKKKLLKLMLLLKKMMMLMLMMLLLLIJKLKKLKLMMKIJLMMLLKKJLLIJKJJIJIJJIIIIJJIJIIJHHHGHJKJIIIJJIGFGHHGHJHHIJHHIIIHGIJJIHIHGEFHGGFHGEDEFEDEGGFGFFGGFHGHHGGFGGHHHJIHHIJIJLLIJJJJKKJKJJJKKJKJIJJIJJJIJIKJJJJKJIIJJIIJJJLKIJJKJKJIIKKLMLJJIJLLLKJKMMMLKKKJIKMLLIJKLKKKMKKJJJIJJKJJJHJKIIJJJKKLLJJKKJJJIJIJIKKIJKLKKKLKKJJKJKJIJJKMLLLMLKJLMMLLLJLLKJKJJJLLKLMMKKLKGGMMKMMMLJJLKLKIJKKJJJIIHIKJHHHIKIGHHJJIIJJJIHHGHHHHHIJJIIIGHHHGIJGHHGGGEEFHIGGHG \ No newline at end of file diff --git a/corr/corrShared.cuh b/corr/corrShared.cuh new file mode 100644 index 0000000..ac7a951 --- /dev/null +++ b/corr/corrShared.cuh @@ -0,0 +1,150 @@ +/* + * + * corrShared.cuh + * + * Header file for implementation of normalized correlation + * that reads the image from shared memory and the template + * from constant memory. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +extern __shared__ unsigned char LocalBlock[]; + + +__global__ void +corrShared_kernel( + float *pCorr, size_t CorrPitch, + int wTile, + int wTemplate, int hTemplate, + float xOffset, float yOffset, + float cPixels, float fDenomExp, int SharedPitch, + float xUL, float yUL, int w, int h ) +{ + int uTile = blockIdx.x*wTile; + int vTile = blockIdx.y*blockDim.y; + int v = vTile + threadIdx.y; + + float *pOut = (float *) (((char *) pCorr)+v*CorrPitch); + + for ( int row = threadIdx.y; + row < blockDim.y+hTemplate; + row += blockDim.y ) { + int SharedIdx = row * SharedPitch; + for ( int col = threadIdx.x; + col < wTile+wTemplate; + col += blockDim.x ) { + + LocalBlock[SharedIdx+col] = + tex2D( texImage, + (float) (uTile+col+xUL+xOffset), + (float) (vTile+row+yUL+yOffset) ); + + } + } + + __syncthreads(); + + for ( int col = threadIdx.x; + col < wTile; + col += blockDim.x ) { + + int SumI = 0; + int SumISq = 0; + int SumIT = 0; + int idx = 0; + int SharedIdx = threadIdx.y * SharedPitch + col; + for ( int j = 0; j < hTemplate; j++ ) { + for ( int i = 0; i < wTemplate; i++) { + unsigned char I = LocalBlock[SharedIdx+i]; + unsigned char T = g_Tpix[idx++]; + SumI += I; + SumISq += I*I; + SumIT += I*T; + } + SharedIdx += SharedPitch; + } + if ( uTile+col < w && v < h ) { + pOut[uTile+col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, fDenomExp ); + } + } + __syncthreads(); +} + + +void +corrShared( + float *dCorr, int CorrPitch, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + int device; + cudaDeviceProp props; + cudaError_t status; + + CUDART_CHECK( cudaGetDevice( &device ) ); + CUDART_CHECK( cudaGetDeviceProperties( &props, device ) ); + if ( sharedMem > props.sharedMemPerBlock ) { + dim3 tcThreads(32, 16, 1); + dim3 tcBlocks; + tcBlocks.x = INTCEIL(w,threads.x); + tcBlocks.y = INTCEIL(h,threads.y); + tcBlocks.z = 1; + return corrTexConstant( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + cPixels, + fDenomExp, + sharedPitch, + xOffset, yOffset, + xTemplate, yTemplate, + xUL, yUL, w, h, + tcThreads, tcBlocks, + sharedMem ); + } + corrShared_kernel<<>>( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); +Error: + return; +} diff --git a/corr/corrShared4.cuh b/corr/corrShared4.cuh new file mode 100644 index 0000000..22f668b --- /dev/null +++ b/corr/corrShared4.cuh @@ -0,0 +1,186 @@ +/* + * + * corrShared.cuh + * + * Header file for implementation of normalized correlation + * that reads the image from shared memory and the template + * from constant memory. This implementation includes the + * SM-aware optimizations of corrSharedSMSums.cuh and also + * unrolls the innermost loop 4x. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +extern __shared__ unsigned char LocalBlock[]; + + +template +__global__ void +corrShared4_kernel( + float *pCorr, size_t CorrPitch, + int wTile, + int wTemplate, int hTemplate, + float xOffset, float yOffset, + float cPixels, float fDenomExp, int SharedPitch, + float xUL, float yUL, int w, int h ) +{ + int uTile = blockIdx.x*wTile; + int vTile = blockIdx.y*blockDim.y; + int v = vTile + threadIdx.y; + + float *pOut = (float *) (((char *) pCorr)+v*CorrPitch); + + for ( int row = threadIdx.y; + row < blockDim.y+hTemplate; + row += blockDim.y ) { + int SharedIdx = row * SharedPitch; + for ( int col = threadIdx.x; + col < wTile+wTemplate; + col += blockDim.x ) { + + LocalBlock[SharedIdx+col] = + tex2D( texImage, + (float) (uTile+col+xUL+xOffset), + (float) (vTile+row+yUL+yOffset) ); + + } + } + + __syncthreads(); + + for ( int col = threadIdx.x; + col < wTile; + col += blockDim.x ) { + + int SumI = 0; + int SumISq = 0; + int SumIT = 0; + int idx = 0; + int SharedIdx = threadIdx.y * SharedPitch + col; + for ( int j = 0; j < hTemplate; j++ ) { + for ( int i = 0; i < wTemplate/4; i++) { + corrSharedAccumulate( SumI, SumISq, SumIT, LocalBlock[SharedIdx+i*4+0], g_Tpix[idx++] ); + corrSharedAccumulate( SumI, SumISq, SumIT, LocalBlock[SharedIdx+i*4+1], g_Tpix[idx++] ); + corrSharedAccumulate( SumI, SumISq, SumIT, LocalBlock[SharedIdx+i*4+2], g_Tpix[idx++] ); + corrSharedAccumulate( SumI, SumISq, SumIT, LocalBlock[SharedIdx+i*4+3], g_Tpix[idx++] ); + } + SharedIdx += SharedPitch; + } + if ( uTile+col < w && v < h ) { + pOut[uTile+col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, fDenomExp ); + } + } + __syncthreads(); +} + + +void +corrShared4( + float *dCorr, int CorrPitch, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + int device; + cudaDeviceProp props; + cudaError_t status; + + CUDART_CHECK( cudaGetDevice( &device ) ); + CUDART_CHECK( cudaGetDeviceProperties( &props, device ) ); + if ( sharedMem > props.sharedMemPerBlock ) { + dim3 tcThreads(32, 16, 1); + dim3 tcBlocks; + tcBlocks.x = INTCEIL(w,threads.x); + tcBlocks.y = INTCEIL(h,threads.y); + tcBlocks.z = 1; + return corrTexConstant( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + cPixels, + fDenomExp, + sharedPitch, + xOffset, yOffset, + xTemplate, yTemplate, + xUL, yUL, w, h, + tcThreads, tcBlocks, + sharedMem ); + } + if ( wTemplate % 4 ) { + if ( props.major == 1 ) { + corrSharedSM_kernel<<>>( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } + else { + corrSharedSM_kernel<<>>( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } + } + if ( props.major == 1 ) { + corrShared4_kernel<<>>( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } + else { + corrShared4_kernel<<>>( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } +Error: + return; +} diff --git a/corr/corrShared4Sums.cuh b/corr/corrShared4Sums.cuh new file mode 100644 index 0000000..c47a6f5 --- /dev/null +++ b/corr/corrShared4Sums.cuh @@ -0,0 +1,217 @@ +/* + * + * corrShared4Sums.cuh + * + * Header file for implementation of normalized correlation + * that reads the image from shared memory and the template + * from constant memory, and reports the sums (I, IT, ISq) + * as well as the output coefficients. This implementation + * includes the SM-aware optimizations of corrSharedSMSums.cuh + * and also unrolls the innermost loop 4x. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +extern __shared__ unsigned char LocalBlock[]; + +template +__device__ void +corrSharedAccumulate( + int& SumI, int& SumISq, int& SumIT, + unsigned char I, unsigned char T ) +{ + SumI += I; + if ( bSM1 ) { + SumISq += __umul24(I,I); + SumIT += __umul24(I,T); + } + else { + SumISq += I*I; + SumIT += I*T; + } +} + +template +__global__ void +corrShared4Sums_kernel( + float *pCorr, size_t CorrPitch, + int *pI, int *pISq, int *pIT, + int wTile, + int wTemplate, int hTemplate, + float xOffset, float yOffset, + float cPixels, float fDenomExp, int SharedPitch, + float xUL, float yUL, int w, int h ) +{ + int uTile = blockIdx.x*wTile; + int vTile = blockIdx.y*blockDim.y; + int v = vTile + threadIdx.y; + + float *pOut = (float *) (((char *) pCorr)+v*CorrPitch); + pI = (int *) (((char *) pI)+v*CorrPitch); + pISq = (int *) (((char *) pISq)+v*CorrPitch); + pIT = (int *) (((char *) pIT)+v*CorrPitch); + + for ( int row = threadIdx.y; + row < blockDim.y+hTemplate; + row += blockDim.y ) { + int SharedIdx = row * SharedPitch; + + for ( int col = threadIdx.x; + col < wTile+wTemplate; + col += blockDim.x ) { + + LocalBlock[SharedIdx+col] = + tex2D( texImage, + (float) (uTile+col+xUL+xOffset), + (float) (vTile+row+yUL+yOffset) ); + + } + } + + __syncthreads(); + + for ( int col = threadIdx.x; + col < wTile; + col += blockDim.x ) { + + int SumI = 0; + int SumISq = 0; + int SumIT = 0; + int idx = 0; + int SharedIdx = threadIdx.y * SharedPitch + col; + for ( int j = 0; j < hTemplate; j++ ) { + for ( int i = 0; i < wTemplate/4; i++) { + corrSharedAccumulate( SumI, SumISq, SumIT, LocalBlock[SharedIdx+i*4+0], g_Tpix[idx++] ); + corrSharedAccumulate( SumI, SumISq, SumIT, LocalBlock[SharedIdx+i*4+1], g_Tpix[idx++] ); + corrSharedAccumulate( SumI, SumISq, SumIT, LocalBlock[SharedIdx+i*4+2], g_Tpix[idx++] ); + corrSharedAccumulate( SumI, SumISq, SumIT, LocalBlock[SharedIdx+i*4+3], g_Tpix[idx++] ); + } + SharedIdx += SharedPitch; + } + if ( uTile+col < w && v < h ) { + pI[uTile+col] = SumI; + pISq[uTile+col] = SumISq; + pIT[uTile+col] = SumIT; + pOut[uTile+col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, g_fDenomExp ); + } + } + __syncthreads(); +} + + +void +corrShared4Sums( + float *dCorr, int CorrPitch, + int *dSumI, int *dSumISq, int *dSumIT, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + int device; + cudaDeviceProp props; + cudaError_t status; + + CUDART_CHECK( cudaGetDevice( &device ) ); + CUDART_CHECK( cudaGetDeviceProperties( &props, device ) ); + if ( sharedMem > props.sharedMemPerBlock ) { + dim3 tcThreads(32, 16, 1); + dim3 tcBlocks; + tcBlocks.x = INTCEIL(w,threads.x); + tcBlocks.y = INTCEIL(h,threads.y); + tcBlocks.z = 1; + return corrTexConstantSums( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + cPixels, + fDenomExp, + sharedPitch, + xOffset, yOffset, + xTemplate, yTemplate, + xUL, yUL, w, h, + tcThreads, tcBlocks, + sharedMem ); + } + if ( wTemplate%4 ) { + if ( props.major == 1 ) { + corrSharedSMSums_kernel<<>>( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } + else { + corrSharedSMSums_kernel<<>>( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } + } + if ( props.major == 1 ) { + corrShared4Sums_kernel<<>>( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } + else { + corrShared4Sums_kernel<<>>( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } +Error: + return; +} diff --git a/corr/corrSharedSM.cuh b/corr/corrSharedSM.cuh new file mode 100644 index 0000000..0c5711e --- /dev/null +++ b/corr/corrSharedSM.cuh @@ -0,0 +1,171 @@ +/* + * + * corrSharedSM.cuh + * + * Header file for implementation of normalized correlation + * that reads the image from shared memory and the template + * from constant memory. This implementation is SM-aware; + * it launches different kernels depending on whether it's + * running on SM 1.x hardware. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +extern __shared__ unsigned char LocalBlock[]; + + +template +__global__ void +corrSharedSM_kernel( + float *pCorr, size_t CorrPitch, + int wTile, + int wTemplate, int hTemplate, + float xOffset, float yOffset, + float cPixels, float fDenomExp, int SharedPitch, + float xUL, float yUL, int w, int h ) +{ + int uTile = blockIdx.x*wTile; + int vTile = blockIdx.y*blockDim.y; + int v = vTile + threadIdx.y; + + float *pOut = (float *) (((char *) pCorr)+v*CorrPitch); + + for ( int row = threadIdx.y; + row < blockDim.y+hTemplate; + row += blockDim.y ) { + int SharedIdx = row * SharedPitch; + for ( int col = threadIdx.x; + col < wTile+wTemplate; + col += blockDim.x ) { + + LocalBlock[SharedIdx+col] = + tex2D( texImage, + (float) (uTile+col+xUL+xOffset), + (float) (vTile+row+yUL+yOffset) ); + + } + } + + __syncthreads(); + + for ( int col = threadIdx.x; + col < wTile; + col += blockDim.x ) { + + int SumI = 0; + int SumISq = 0; + int SumIT = 0; + int idx = 0; + int SharedIdx = threadIdx.y * SharedPitch + col; + for ( int j = 0; j < hTemplate; j++ ) { + for ( int i = 0; i < wTemplate; i++) { + unsigned char I = LocalBlock[SharedIdx+i]; + unsigned char T = g_Tpix[idx++]; + SumI += I; + if ( bSM1 ) { + SumISq += __umul24(I, I); + SumIT += __umul24(I, T); + } + else { + SumISq += I*I; + SumIT += I*T; + } + } + SharedIdx += SharedPitch; + } + if ( uTile+col < w && v < h ) { + pOut[uTile+col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, fDenomExp ); + } + } + __syncthreads(); +} + + +void +corrSharedSM( + float *dCorr, int CorrPitch, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + int device; + cudaDeviceProp props; + cudaError_t status; + + CUDART_CHECK( cudaGetDevice( &device ) ); + CUDART_CHECK( cudaGetDeviceProperties( &props, device ) ); + if ( sharedMem > props.sharedMemPerBlock ) { + dim3 tcThreads(32, 16, 1); + dim3 tcBlocks; + tcBlocks.x = INTCEIL(w,threads.x); + tcBlocks.y = INTCEIL(h,threads.y); + tcBlocks.z = 1; + return corrTexConstant( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + cPixels, + fDenomExp, + sharedPitch, + xOffset, yOffset, + xTemplate, yTemplate, + xUL, yUL, w, h, + tcThreads, tcBlocks, + sharedMem ); + } + if ( props.major == 1 ) { + corrSharedSM_kernel<<>>( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } + else { + corrSharedSM_kernel<<>>( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } +Error: + return; +} diff --git a/corr/corrSharedSMSums.cuh b/corr/corrSharedSMSums.cuh new file mode 100644 index 0000000..23cdde4 --- /dev/null +++ b/corr/corrSharedSMSums.cuh @@ -0,0 +1,183 @@ +/* + * + * corrSharedSums.cuh + * + * Header file for implementation of normalized correlation + * that reads the image from shared memory and the template + * from constant memory, and reports the sums (I, IT, ISq) + * as well as the output coefficients. This implementation + * is SM-aware; it launches different kernels depending on + * whether it's running on SM 1.x hardware. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +extern __shared__ unsigned char LocalBlock[]; + +template +__global__ void +corrSharedSMSums_kernel( + float *pCorr, size_t CorrPitch, + int *pI, int *pISq, int *pIT, + int wTile, + int wTemplate, int hTemplate, + float xOffset, float yOffset, + float cPixels, float fDenomExp, int SharedPitch, + float xUL, float yUL, int w, int h ) +{ + int uTile = blockIdx.x*wTile; + int vTile = blockIdx.y*blockDim.y; + int v = vTile + threadIdx.y; + + float *pOut = (float *) (((char *) pCorr)+v*CorrPitch); + pI = (int *) (((char *) pI)+v*CorrPitch); + pISq = (int *) (((char *) pISq)+v*CorrPitch); + pIT = (int *) (((char *) pIT)+v*CorrPitch); + + for ( int row = threadIdx.y; + row < blockDim.y+hTemplate; + row += blockDim.y ) { + int SharedIdx = row * SharedPitch; + + for ( int col = threadIdx.x; + col < wTile+wTemplate; + col += blockDim.x ) { + + LocalBlock[SharedIdx+col] = + tex2D( texImage, + (float) (uTile+col+xUL+xOffset), + (float) (vTile+row+yUL+yOffset) ); + + } + } + + __syncthreads(); + + for ( int col = threadIdx.x; + col < wTile; + col += blockDim.x ) { + + int SumI = 0; + int SumISq = 0; + int SumIT = 0; + int idx = 0; + int SharedIdx = threadIdx.y * SharedPitch + col; + for ( int j = 0; j < hTemplate; j++ ) { + for ( int i = 0; i < wTemplate; i++) { + unsigned char I = LocalBlock[SharedIdx+i]; + unsigned char T = g_Tpix[idx++]; + SumI += I; + if ( bSM1 ) { + SumISq += __umul24( I, I ); + SumIT += __umul24( I, T ); + } + else { + SumISq += I*I; + SumIT += I*T; + } + } + SharedIdx += SharedPitch; + } + if ( uTile+col < w && v < h ) { + pI[uTile+col] = SumI; + pISq[uTile+col] = SumISq; + pIT[uTile+col] = SumIT; + pOut[uTile+col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, g_fDenomExp ); + } + } + __syncthreads(); +} + + +void +corrSharedSMSums( + float *dCorr, int CorrPitch, + int *dSumI, int *dSumISq, int *dSumIT, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + int device; + cudaDeviceProp props; + cudaError_t status; + + CUDART_CHECK( cudaGetDevice( &device ) ); + CUDART_CHECK( cudaGetDeviceProperties( &props, device ) ); + if ( sharedMem > props.sharedMemPerBlock ) { + dim3 tcThreads(32, 16, 1); + dim3 tcBlocks; + tcBlocks.x = INTCEIL(w,threads.x); + tcBlocks.y = INTCEIL(h,threads.y); + tcBlocks.z = 1; + return corrTexTexSums( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + cPixels, + fDenomExp, + sharedPitch, + xOffset, yOffset, + xTemplate, yTemplate, + xUL, yUL, w, h, + tcThreads, tcBlocks, + sharedMem ); + } + if ( props.major == 1 ) { + corrSharedSMSums_kernel<<>>( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } + else { + corrSharedSMSums_kernel<<>>( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); + } +Error: + return; +} diff --git a/corr/corrSharedSums.cuh b/corr/corrSharedSums.cuh new file mode 100644 index 0000000..371ede6 --- /dev/null +++ b/corr/corrSharedSums.cuh @@ -0,0 +1,161 @@ +/* + * + * corrSharedSums.cuh + * + * Header file for implementation of normalized correlation + * that reads the image from shared memory and the template + * from constant memory, and reports the sums (I, IT, ISq) + * as well as the output coefficients. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +extern __shared__ unsigned char LocalBlock[]; + +__global__ void +corrSharedSums_kernel( + float *pCorr, size_t CorrPitch, + int *pI, int *pISq, int *pIT, + int wTile, + int wTemplate, int hTemplate, + float xOffset, float yOffset, + float cPixels, float fDenomExp, int SharedPitch, + float xUL, float yUL, int w, int h ) +{ + int uTile = blockIdx.x*wTile; + int vTile = blockIdx.y*blockDim.y; + int v = vTile + threadIdx.y; + + float *pOut = (float *) (((char *) pCorr)+v*CorrPitch); + pI = (int *) (((char *) pI)+v*CorrPitch); + pISq = (int *) (((char *) pISq)+v*CorrPitch); + pIT = (int *) (((char *) pIT)+v*CorrPitch); + + for ( int row = threadIdx.y; + row < blockDim.y+hTemplate; + row += blockDim.y ) { + int SharedIdx = row * SharedPitch; + + for ( int col = threadIdx.x; + col < wTile+wTemplate; + col += blockDim.x ) { + + LocalBlock[SharedIdx+col] = + tex2D( texImage, + (float) (uTile+col+xUL+xOffset), + (float) (vTile+row+yUL+yOffset) ); + + } + } + + __syncthreads(); + + for ( int col = threadIdx.x; + col < wTile; + col += blockDim.x ) { + + int SumI = 0; + int SumISq = 0; + int SumIT = 0; + int idx = 0; + int SharedIdx = threadIdx.y * SharedPitch + col; + for ( int j = 0; j < hTemplate; j++ ) { + for ( int i = 0; i < wTemplate; i++) { + unsigned char I = LocalBlock[SharedIdx+i]; + unsigned char T = g_Tpix[idx++]; + SumI += I; + SumISq += I*I; + SumIT += I*T; + } + SharedIdx += SharedPitch; + } + if ( uTile+col < w && v < h ) { + pI[uTile+col] = SumI; + pISq[uTile+col] = SumISq; + pIT[uTile+col] = SumIT; + pOut[uTile+col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, g_fDenomExp ); + } + } + __syncthreads(); +} + + +void +corrSharedSums( + float *dCorr, int CorrPitch, + int *dSumI, int *dSumISq, int *dSumIT, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + int device; + cudaDeviceProp props; + cudaError_t status; + + CUDART_CHECK( cudaGetDevice( &device ) ); + CUDART_CHECK( cudaGetDeviceProperties( &props, device ) ); + if ( sharedMem > props.sharedMemPerBlock ) { + dim3 tcThreads(32, 16, 1); + dim3 tcBlocks; + tcBlocks.x = INTCEIL(w,threads.x); + tcBlocks.y = INTCEIL(h,threads.y); + tcBlocks.z = 1; + return corrTexConstantSums( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + cPixels, + fDenomExp, + sharedPitch, + xOffset, yOffset, + xTemplate, yTemplate, + xUL, yUL, w, h, + tcThreads, tcBlocks, + sharedMem ); + } + corrSharedSums_kernel<<>>( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + (float) xOffset, (float) yOffset, + cPixels, fDenomExp, + sharedPitch, + (float) xUL, (float) yUL, w, h ); +Error: + return; +} diff --git a/corr/corrTexConstant.cuh b/corr/corrTexConstant.cuh new file mode 100644 index 0000000..41ec3cf --- /dev/null +++ b/corr/corrTexConstant.cuh @@ -0,0 +1,94 @@ +/* + * + * corrTexConstant.cuh + * + * Header file for 2D implementation of normalized correlation + * that reads the template from constant memory. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +__global__ void +corrTexConstant_kernel( + float *pCorr, size_t CorrPitch, + float cPixels, float fDenomExp, + float xUL, float yUL, int w, int h, + int xOffset, int yOffset, + int wTemplate, int hTemplate ) +{ + size_t row = blockIdx.y*blockDim.y + threadIdx.y; + size_t col = blockIdx.x*blockDim.x + threadIdx.x; + + // adjust pointers to row + pCorr = (float *) ((char *) pCorr+row*CorrPitch); + + // No __syncthreads in this kernel, so we can early-out + // without worrying about the effects of divergence. + if ( col >= w || row >= h ) + return; + + int SumI = 0; + int SumISq = 0; + int SumIT = 0; + int inx = 0; + + for ( int j = 0; j < hTemplate; j++ ) { + for ( int i = 0; i < wTemplate; i++ ) { + unsigned char I = tex2D( texImage, + (float) col+xUL+xOffset+i, + (float) row+yUL+yOffset+j ); + unsigned char T = g_Tpix[inx++]; + SumI += I; + SumISq += I*I; + SumIT += I*T; + } + } + pCorr[col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, fDenomExp ); +} + +void +corrTexConstant( + float *dCorr, int CorrPitch, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + corrTexConstant_kernel<<>>( + dCorr, CorrPitch, + cPixels, fDenomExp, + (float) xUL, (float) yUL, w, h, + xOffset, yOffset, wTemplate, hTemplate ); +} diff --git a/corr/corrTexConstantSums.cuh b/corr/corrTexConstantSums.cuh new file mode 100644 index 0000000..dc61ca1 --- /dev/null +++ b/corr/corrTexConstantSums.cuh @@ -0,0 +1,105 @@ +/* + * + * corrTexConstant.cuh + * + * Header file for implementation of normalized correlation + * that reads the template from constant memory, and reports + * the sums (I, IT, ISq) as well as the output coefficients. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +__global__ void +corrTexConstantSums_kernel( + float *pCorr, size_t CorrPitch, + int *pI, int *pISq, int *pIT, + float cPixels, float fDenomExp, + float xUL, float yUL, int w, int h, + int xOffset, int yOffset, + int wTemplate, int hTemplate ) +{ + size_t row = blockIdx.y*blockDim.y + threadIdx.y; + size_t col = blockIdx.x*blockDim.x + threadIdx.x; + + // adjust pointers to row + pCorr = (float *) ((char *) pCorr+row*CorrPitch); + pI = (int *) ((char *) pI +row*CorrPitch); + pISq = (int *) ((char *) pISq+row*CorrPitch); + pIT = (int *) ((char *) pIT +row*CorrPitch); + + // No __syncthreads in this kernel, so we can early-out + // without worrying about the effects of divergence. + if ( col >= w || row >= h ) + return; + + int SumI = 0; + int SumISq = 0; + int SumIT = 0; + + int inx = 0; + + for ( int j = 0; j < hTemplate; j++ ) { + for ( int i = 0; i < wTemplate; i++ ) { + unsigned char I = tex2D( texImage, + (float) col+xUL+xOffset+i, + (float) row+yUL+yOffset+j ); + unsigned char T = g_Tpix[inx++]; + SumI += I; + SumISq += I*I; + SumIT += I*T; + } + } + pCorr[col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, fDenomExp ); + pI[col] = SumI; + pISq[col] = SumISq; + pIT[col] = SumIT; +} + +void +corrTexConstantSums( + float *dCorr, int CorrPitch, + int *dSumI, int *dSumISq, int *dSumIT, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + corrTexConstantSums_kernel<<>>( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + cPixels, fDenomExp, + (float) xUL, (float) yUL, w, h, + xOffset, yOffset, wTemplate, hTemplate ); +} diff --git a/corr/corrTexTex.cuh b/corr/corrTexTex.cuh new file mode 100644 index 0000000..9fc9a1d --- /dev/null +++ b/corr/corrTexTex.cuh @@ -0,0 +1,101 @@ +/* + * + * corrTexTex.cuh + * + * Header file for implementation of normalized correlation + * that reads both image and template from texture. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +__global__ void +corrTexTex_kernel( + float *pCorr, size_t CorrPitch, + float cPixels, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int wTemplate, int hTemplate, + float xUL, float yUL, int w, int h ) +{ + size_t row = blockIdx.y*blockDim.y + threadIdx.y; + size_t col = blockIdx.x*blockDim.x + threadIdx.x; + + // adjust pCorr to point to row + pCorr = (float *) ((char *) pCorr+row*CorrPitch); + + // No __syncthreads in this kernel, so we can early-out + // without worrying about the effects of divergence. + if ( col >= w || row >= h ) + return; + + int SumI = 0; + int SumT = 0; + int SumISq = 0; + int SumTSq = 0; + int SumIT = 0; + for ( int y = 0; y < hTemplate; y++ ) { + for ( int x = 0; x < wTemplate; x++ ) { + unsigned char I = tex2D( texImage, + (float) col+xUL+xOffset+x, (float) row+yUL+yOffset+y ); + unsigned char T = tex2D( texTemplate, + (float) xTemplate+x, (float) yTemplate+y); + SumI += I; + SumT += T; + SumISq += I*I; + SumTSq += T*T; + SumIT += I*T; + } + float fDenomExp = (float) + ( (double) cPixels*SumTSq - (double) SumT*SumT); + pCorr[col] = CorrelationValue( SumI, SumISq, SumIT, SumT, cPixels, fDenomExp ); + } +} + +void +corrTexTex( + float *dCorr, int CorrPitch, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + corrTexTex_kernel<<>>( + dCorr, CorrPitch, + cPixels, + xOffset, yOffset, + xTemplate+xOffset, yTemplate+yOffset, + wTemplate, hTemplate, + (float) xUL, (float) yUL, w, h ); +} diff --git a/corr/corrTexTexSums.cuh b/corr/corrTexTexSums.cuh new file mode 100644 index 0000000..a2939c5 --- /dev/null +++ b/corr/corrTexTexSums.cuh @@ -0,0 +1,115 @@ +/* + * + * corrTexTexSums.cuh + * + * Header file for implementation of normalized correlation + * that reads the template from texture, and reports the + * sums (I, IT, ISq) in addition to the coefficients. + * + * Copyright (c) 2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +__global__ void +corrTexTexSums_kernel( + float *pCorr, size_t CorrPitch, + int *pI, int *pISq, int *pIT, + float cPixels, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int wTemplate, int hTemplate, + float xUL, float yUL, int w, int h ) +{ + size_t row = blockIdx.y*blockDim.y + threadIdx.y; + size_t col = blockIdx.x*blockDim.x + threadIdx.x; + + // adjust pointers to row + pCorr = (float *) ((char *) pCorr+row*CorrPitch); + pI = (int *) ((char *) pI +row*CorrPitch); + pISq = (int *) ((char *) pISq+row*CorrPitch); + pIT = (int *) ((char *) pIT +row*CorrPitch); + + // No __syncthreads in this kernel, so we can early-out + // without worrying about the effects of divergence. + if ( col >= w || row >= h ) + return; + + int SumI = 0; + int SumT = 0; + int SumISq = 0; + int SumTSq = 0; + int SumIT = 0; + for ( int y = 0; y < hTemplate; y++ ) { + for ( int x = 0; x < wTemplate; x++ ) { + unsigned char I = + tex2D( texImage, + (float) col+xUL+xOffset+x, + (float) row+yUL+yOffset+y ); + unsigned char T = + tex2D( texTemplate, + (float) xTemplate+x, + (float) yTemplate+y); + SumI += I; + SumT += T; + SumISq += I*I; + SumTSq += T*T; + SumIT += I*T; + } + float fDenomExp = (float) + ((double) cPixels*SumTSq - (double) SumT*SumT); + pI[col] = SumI; + pISq[col] = SumISq; + pIT[col] = SumIT; + pCorr[col] = CorrelationValue( SumI, SumISq, SumIT, SumT, cPixels, fDenomExp ); + } +} + +void +corrTexTexSums( + float *dCorr, int CorrPitch, + int *pI, int *pISq, int *pIT, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ) +{ + corrTexTexSums_kernel<<>>( + dCorr, CorrPitch, + pI, pISq, pIT, + cPixels, + xOffset, yOffset, + xTemplate+xOffset, yTemplate+yOffset, + wTemplate, hTemplate, + (float) xUL, (float) yUL, w, h ); +} diff --git a/corr/normalizedCrossCorrelation.cu b/corr/normalizedCrossCorrelation.cu new file mode 100644 index 0000000..5398893 --- /dev/null +++ b/corr/normalizedCrossCorrelation.cu @@ -0,0 +1,641 @@ +/* + * + * normalizedCrossCorrelation.cu + * + * Microbenchmark for normalized cross correlation, a template- + * matching algorithm for computer vision. + * + * Build with: nvcc -I ../chLib normalizedCrossCorrelation.cu pgm.cu + * + * Make sure to include pgm.cu for the image file I/O support. + * + * To avoid warnings about double precision support, specify the + * target gpu-architecture, e.g.: + * nvcc --gpu-architecture sm_13 -I ../chLib normalizedCrossCorrelation.cu pgm.cu + * + * Requires: No minimum SM requirement. + * + * Copyright (c) 2011-2012, Archaea Software, LLC. + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include "pgm.h" + +texture texImage; +texture texTemplate; + +const int maxTemplatePixels = 3072; +__constant__ int g_xOffset[maxTemplatePixels]; +__constant__ int g_yOffset[maxTemplatePixels]; +__constant__ unsigned char g_Tpix[maxTemplatePixels]; +__constant__ float g_cPixels, g_SumT, g_fDenomExp; +unsigned int g_cpuSumT, g_cpuSumTSq; + +const float fThreshold = 1e-3f; + +#define INTCEIL(a,b) ( ((a)+(b)-1) / (b) ) + +__device__ __host__ inline float +CorrelationValue( float SumI, float SumISq, float SumIT, float SumT, float cPixels, float fDenomExp ) +{ + float Numerator = cPixels*SumIT - SumI*SumT; + float Denominator = rsqrtf( (cPixels*SumISq - SumI*SumI)*fDenomExp ); + return Numerator * Denominator; +} + +#include "corrTexTexSums.cuh" +#include "corrTexTex.cuh" + +#include "corrTexConstantSums.cuh" +#include "corrTexConstant.cuh" + +extern __shared__ unsigned char LocalBlock[]; + +#include "corrSharedSMSums.cuh" +#include "corrSharedSM.cuh" + +#include "corrSharedSums.cuh" +#include "corrShared.cuh" + +#include "corrShared4Sums.cuh" +#include "corrShared4.cuh" + +int poffsetx[maxTemplatePixels]; +int poffsety[maxTemplatePixels]; + +cudaError_t +CopyToTemplate( + unsigned char *img, size_t imgPitch, + int xTemplate, int yTemplate, + int wTemplate, int hTemplate, + int OffsetX, int OffsetY +) +{ + cudaError_t status; + unsigned char pixels[maxTemplatePixels]; + + int inx = 0; + int SumT = 0; + int SumTSq = 0; + int cPixels = wTemplate*hTemplate; + size_t sizeOffsets = cPixels*sizeof(int); + float fSumT, fDenomExp, fcPixels; + + CUDART_CHECK( + cudaMemcpy2D( + pixels, wTemplate, + img+yTemplate*imgPitch+xTemplate, imgPitch, + wTemplate, hTemplate, + cudaMemcpyDeviceToHost ) ); + + CUDART_CHECK( cudaMemcpyToSymbol( g_Tpix, pixels, cPixels ) ); + + for ( int i = OffsetY; i < OffsetY+hTemplate; i++ ) { + for ( int j = OffsetX; j < OffsetX+wTemplate; j++) { + SumT += pixels[inx]; + SumTSq += pixels[inx]*pixels[inx]; + poffsetx[inx] = j; + poffsety[inx] = i; + inx += 1; + } + } + g_cpuSumT = SumT; + g_cpuSumTSq = SumTSq; + + CUDART_CHECK( cudaMemcpyToSymbol(g_xOffset, poffsetx, sizeOffsets) ); + CUDART_CHECK( cudaMemcpyToSymbol(g_yOffset, poffsety, sizeOffsets) ); + + fSumT = (float) SumT; + CUDART_CHECK( cudaMemcpyToSymbol(g_SumT, &fSumT, sizeof(float)) ); + + fDenomExp = (float) ( (double)cPixels*SumTSq - (double) SumT*SumT); + CUDART_CHECK( cudaMemcpyToSymbol(g_fDenomExp, &fDenomExp, sizeof(float)) ); + + fcPixels = (float) cPixels; + CUDART_CHECK( cudaMemcpyToSymbol(g_cPixels, &fcPixels, sizeof(float)) ); +Error: + return status; +} + +int +bCompareCorrValues( const float *pBase0, + const float *pBase1, + int w, int h ) +{ + for ( int j = 0; j < h; j++ ) { + + float *pf0 = (float *) ((char *) pBase0+j*w*sizeof(float)); + float *pf1 = (float *) ((char *) pBase1+j*w*sizeof(float)); + + for ( int i = 0; i < w; i++ ) { + if ( fabsf(pf0[i]-pf1[i]) > fThreshold ) { + printf( "Mismatch pf0[%d] = %.5f, pf1[%d] = %.5f\n", i, pf0[i], i, pf1[i] ); + fflush( stdout ); + //CH_ASSERT(0); + return 1; + } + } + } + return 0; +} + +int +bCompareSums( const int *pBaseI0, const int *pBaseISq0, const int *pBaseIT0, + const int *pBaseI1, const int *pBaseISq1, const int *pBaseIT1, + int w, int h ) +{ + for ( int j = 0; j < h; j++ ) { + + const int *pi0 = (const int *) ((char *) pBaseI0+j*w*sizeof(int)); + const int *pi1 = (const int *) ((char *) pBaseI1+j*w*sizeof(int)); + + const int *pisq0 = (const int *) ((char *) pBaseISq0+j*w*sizeof(int)); + const int *pisq1 = (const int *) ((char *) pBaseISq1+j*w*sizeof(int)); + + const int *pit0 = (const int *) ((char *) pBaseIT0+j*w*sizeof(int)); + const int *pit1 = (const int *) ((char *) pBaseIT1+j*w*sizeof(int)); + for ( int i = 0; i < w; i++ ) { + if ( pi0[i] != pi1[i] || + pisq0[i] != pisq1[i] || + pit0[i] != pit1[i] ) { + printf( "Mismatch pi[%d] = %d, reference = %d\n", i, pi0[i], pi1[i] ); + printf( "Mismatch pisq[%d] = %d, reference = %d\n", i, pisq0[i], pisq1[i] ); + printf( "Mismatch pit[%d] = %d, reference = %d\n", i, pit0[i], pit1[i] ); + fflush( stdout ); + //CH_ASSERT(0); + return 1; + } + } + } + return 0; +} + +unsigned char +ReadPixel( unsigned char *base, int pitch, int w, int h, int x, int y ) +{ + if ( x < 0 ) x = 0; + if ( x >= w ) x = w-1; + if ( y < 0 ) y = 0; + if ( y >= h ) y = h-1; + return base[y*pitch+x]; +} + +void +corrCPU( float *pCorr, + int *_pI, int *_pISq, int *_pIT, + size_t CorrPitch, + int cPixels, + int xTemplate, int yTemplate, + int w, int h, + unsigned char *img, int imgPitch, + unsigned char *tmp, int tmpPitch ) +{ + for ( int row = 0; row < h; row += 1 ) { + float *pOut = (float *) (((char *) pCorr)+row*CorrPitch); + int *pI = (int *) (((char *) _pI)+row*CorrPitch); + int *pISq = (int *) (((char *) _pISq)+row*CorrPitch); + int *pIT = (int *) (((char *) _pIT)+row*CorrPitch); + for ( int col = 0; col < w; col += 1 ) { + int SumI = 0; + int SumT = 0; + int SumISq = 0; + int SumTSq = 0; + int SumIT = 0; + for ( int j = 0; j < cPixels; j++ ) { + unsigned char I = ReadPixel( img, imgPitch, w, h, col+poffsetx[j], row+poffsety[j] ); + unsigned char T = ReadPixel( tmp, tmpPitch, w, h, xTemplate+poffsetx[j], yTemplate+poffsety[j] ); + SumI += I; + SumT += T; + SumISq += I*I; + SumTSq += T*T; + SumIT += I*T; + } + float fDenomExp = (float) ((double) cPixels*SumTSq - (double) SumT*SumT); + pI[col] = SumI; + pISq[col] = SumISq; + pIT[col] = SumIT; + pOut[col] = CorrelationValue( (float) SumI, (float) SumISq, (float) SumIT, (float) SumT, (float) cPixels, fDenomExp ); + } + } +} + +bool +TestCorrelation( + double *pixelsPerSecond, // passbacks to report performance + double *templatePixelsPerSecond, // + int xOffset, int yOffset, // offset into image + int w, int h, // width and height of output + const float *hrefCorr, // host reference data + const int *hrefSumI, + const int *hrefSumISq, + const int *hrefSumIT, + int xTemplate, int yTemplate, // reference point in template image + int wTemplate, int hTemplate, + int wTile, // width of image tile + int sharedPitch, int sharedMem, + dim3 threads, dim3 blocks, + void (*pfnCorrelationSums)( + float *dCorr, int CorrPitch, + int *dSumI, int *dSumISq, int *dSumIT, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ), + void (*pfnCorrelation)( + float *dCorr, int CorrPitch, + int wTile, + int wTemplate, int hTemplate, + float cPixels, + float fDenomExp, + int sharedPitch, + int xOffset, int yOffset, + int xTemplate, int yTemplate, + int xUL, int yUL, int w, int h, + dim3 threads, dim3 blocks, + int sharedMem ), + bool bPrintNeighborhood = false, + int cIterations = 1, + const char *outputFilename = NULL +) +{ + cudaError_t status; + bool ret = false; + size_t CorrPitch; + + float cPixels = (float) wTemplate*hTemplate; + float fDenomExp = (float) ((double) cPixels*g_cpuSumTSq - (double) g_cpuSumT*g_cpuSumT); + + float *hCorr = NULL, *dCorr = NULL; + int *hSumI = NULL, *dSumI = NULL; + int *hSumISq = NULL, *dSumISq = NULL; + int *hSumIT = NULL, *dSumIT = NULL; + + cudaEvent_t start = 0, stop = 0; + + hCorr = (float *) malloc( w*sizeof(float)*h ); + hSumI = (int *) malloc( w*sizeof(int)*h ); + hSumISq = (int *) malloc( w*sizeof(int)*h ); + hSumIT = (int *) malloc( w*sizeof(int)*h ); + if ( NULL == hCorr || NULL == hSumI || NULL == hSumISq || NULL == hSumIT ) + goto Error; + + CUDART_CHECK( cudaMallocPitch( (void **) &dCorr, &CorrPitch, w*sizeof(float), h ) ); + CUDART_CHECK( cudaMallocPitch( (void **) &dSumI, &CorrPitch, w*sizeof(int), h ) ); + CUDART_CHECK( cudaMallocPitch( (void **) &dSumISq, &CorrPitch, w*sizeof(int), h ) ); + CUDART_CHECK( cudaMallocPitch( (void **) &dSumIT, &CorrPitch, w*sizeof(int), h ) ); + + CUDART_CHECK( cudaMemset( dCorr, 0, CorrPitch*h ) ); + CUDART_CHECK( cudaMemset( dSumI, 0, CorrPitch*h ) ); + CUDART_CHECK( cudaMemset( dSumISq, 0, CorrPitch*h ) ); + CUDART_CHECK( cudaMemset( dSumIT, 0, CorrPitch*h ) ); + + CUDART_CHECK( cudaEventCreate( &start, 0 ) ); + CUDART_CHECK( cudaEventCreate( &stop, 0 ) ); + + pfnCorrelationSums( + dCorr, CorrPitch, + dSumI, dSumISq, dSumIT, + wTile, + wTemplate, hTemplate, + cPixels, fDenomExp, + sharedPitch, + xOffset, yOffset, + xTemplate, yTemplate, + 0, 0, w, h, + threads, blocks, sharedMem ); + + CUDART_CHECK( cudaMemcpy2D( hSumI, w*sizeof(int), dSumI, CorrPitch, w*sizeof(int), h, cudaMemcpyDeviceToHost ) ); + CUDART_CHECK( cudaMemcpy2D( hSumISq, w*sizeof(int), dSumISq, CorrPitch, w*sizeof(int), h, cudaMemcpyDeviceToHost ) ); + CUDART_CHECK( cudaMemcpy2D( hSumIT, w*sizeof(int), dSumIT, CorrPitch, w*sizeof(int), h, cudaMemcpyDeviceToHost ) ); + + if ( bCompareSums( hSumI, hSumISq, hSumIT, + hrefSumI, hrefSumISq, hrefSumIT, + w, h ) ) { + //CH_ASSERT(0); + printf( "Sums miscompare\n" ); + goto Error; + } + + CUDART_CHECK( cudaMemcpy2D( hCorr, w*sizeof(float), dCorr, CorrPitch, w*sizeof(float), h, cudaMemcpyDeviceToHost ) ); + + if ( bCompareCorrValues( hrefCorr, hCorr, w, h ) ) { + //CH_ASSERT(0); + printf( "Correlation coefficients generated by sums kernel mismatch\n" ); + return 1; + } + + CUDART_CHECK( cudaMemset2D( dCorr, CorrPitch, 0, w*sizeof(float), h ) ); + CUDART_CHECK( cudaDeviceSynchronize() ); + CUDART_CHECK( cudaEventRecord( start, 0 ) ); + + for ( int i = 0; i < cIterations; i++ ) { + pfnCorrelation( + dCorr, CorrPitch, + wTile, + wTemplate, hTemplate, + cPixels, fDenomExp, + sharedPitch, + xOffset, yOffset, + xTemplate, yTemplate, + 0, 0, w, h, + threads, blocks, sharedMem ); + } + + CUDART_CHECK( cudaEventRecord( stop, 0 ) ); + CUDART_CHECK( cudaMemcpy2D( hCorr, w*sizeof(float), dCorr, CorrPitch, w*sizeof(float), h, cudaMemcpyDeviceToHost ) ); + + if ( bCompareCorrValues( hrefCorr, hCorr, w, h ) ) { + CH_ASSERT(0); + printf( "Correlation coefficients generated by coefficient-only kernel mismatch\n" ); + return 1; + } + + { + float ms; + CUDART_CHECK( cudaEventElapsedTime( &ms, start, stop ) ); + *pixelsPerSecond = (double) w*h*cIterations*1000.0 / ms; + *templatePixelsPerSecond = *pixelsPerSecond*wTemplate*hTemplate; + } + + if ( bPrintNeighborhood ) { + printf( "\nNeighborhood around template:\n" ); + for ( int VertOffset = -4; VertOffset <= 4; VertOffset++ ) { + const float *py = hrefCorr+w*(VertOffset+yTemplate); + for ( int HorzOffset = -4; HorzOffset <= 4; HorzOffset++ ) { + printf( "%6.2f", py[xTemplate+HorzOffset] ); + } + printf("\n"); + } + } + + if ( outputFilename ) { + unsigned char *correlationValues = (unsigned char *) malloc( w*h ); + if ( ! correlationValues ) { + status = cudaErrorMemoryAllocation; + goto Error; + } + for ( int row = 0; row < h; row++ ) { + for ( int col = 0; col < w; col++ ) { + int index = row*w+col; + float value = hCorr[index] < 0.0f ? 0.0f : logf( 1.0f+hCorr[index] )/logf(2.0f); + if ( value < 0.5f ) value = 0.0f; + value = 2.0f * (value - 0.5f); + correlationValues[index] = (unsigned char) (255.0f*value+0.5f); + } + } + if ( 0 != pgmSave( outputFilename, correlationValues, w, h ) ) { + status = cudaErrorUnknown; + goto Error; + } + free( correlationValues ); + } + + ret = true; + +Error: + cudaEventDestroy( start ); + cudaEventDestroy( stop ); + free( hCorr ); + free( hSumI ); + free( hSumISq ); + free( hSumIT ); + if ( dCorr ) cudaFree( dCorr ); + if ( dSumI ) cudaFree( dSumI ); + if ( dSumI ) cudaFree( dSumISq ); + if ( dSumI ) cudaFree( dSumIT ); + return ret; +} + +int +main(int argc, char *argv[]) +{ + int ret = 1; + cudaError_t status; + + unsigned char *hidata = NULL; + unsigned char *didata = NULL; + float *hoCorrCPU = NULL; + + int *hoCorrCPUI = NULL; + int *hoCorrCPUISq = NULL; + int *hoCorrCPUIT = NULL; + unsigned int HostPitch, DevicePitch; + int w, h; + + int wTemplate = 52; + int hTemplate = 52; + int xOffset, yOffset; + + int xTemplate = 210; + int yTemplate = 148; + + int wTile; + dim3 threads; + dim3 blocks; + + int sharedPitch; + int sharedMem; + char *inputFilename = "coins.pgm"; + char *outputFilename = NULL; + + cudaArray *pArrayImage = NULL; + cudaArray *pArrayTemplate = NULL; + cudaChannelFormatDesc desc = cudaCreateChannelDesc(); + + if ( chCommandLineGetBool( "help", argc, argv ) ) { + printf( "Usage:\n" ); + printf( " --input : specify input filename (must be PGM)\n" ); + printf( " --output : Write PGM of correlation values (0..255) to .\n" ); + printf( " --padWidth : pad input image width to specified value\n" ); + printf( " --padHeight : pad input image height to specified value\n" ); + printf( " --xTemplate : X coordinate of upper left corner of template\n" ); + printf( " --yTemplate : Y coordinate of upper left corner of template\n" ); + printf( " --wTemplate : Width of template\n" ); + printf( " --hTemplate : Height of template\n" ); + printf( "\nDefault values are coins.pgm, no output file or padding, and template of the dime in the\n" ); + printf("lower right corner of coins.pgm: xTemplate=210, yTemplate=148, wTemplate=hTemplate=52\n" ); + + return 0; + } + + CUDART_CHECK( cudaSetDeviceFlags( cudaDeviceMapHost ) ); + CUDART_CHECK( cudaDeviceSetCacheConfig( cudaFuncCachePreferShared ) ); + + if ( chCommandLineGet( &inputFilename, "input", argc, argv ) ) { + printf( "Reading from image file %s\n", inputFilename ); + } + chCommandLineGet( &outputFilename, "output", argc, argv ); + { + int padWidth = 0; + int padHeight = 0; + if ( chCommandLineGet( &padWidth, "padWidth", argc, argv ) ) { + if ( ! chCommandLineGet( &padHeight, "padHeight", argc, argv ) ) { + printf( "Must specify both --padWidth and --padHeight\n" ); + goto Error; + } + } + else { + if ( chCommandLineGet( &padHeight, "padHeight", argc, argv ) ) { + printf( "Must specify both --padWidth and --padHeight\n" ); + goto Error; + } + } + if ( pgmLoad(inputFilename, &hidata, &HostPitch, &didata, &DevicePitch, &w, &h, padWidth, padHeight) ) + goto Error; + } + chCommandLineGet( &xTemplate, "xTemplate", argc, argv ); + chCommandLineGet( &yTemplate, "yTemplate", argc, argv ); + chCommandLineGet( &wTemplate, "wTemplate", argc, argv ); + chCommandLineGet( &hTemplate, "hTemplate", argc, argv ); + + xOffset = -wTemplate/2; + yOffset = -wTemplate/2; + + hoCorrCPU = (float *) malloc(w*h*sizeof(float)); if ( ! hoCorrCPU ) return 1; + hoCorrCPUI = (int *) malloc(w*h*sizeof(int)); if ( ! hoCorrCPUI ) return 1; + hoCorrCPUISq = (int *) malloc(w*h*sizeof(int)); if ( ! hoCorrCPUISq ) return 1; + hoCorrCPUIT = (int *) malloc(w*h*sizeof(int)); if ( ! hoCorrCPUIT ) return 1; + if ( NULL == hoCorrCPU || + NULL == hoCorrCPUI || + NULL == hoCorrCPUISq || + NULL == hoCorrCPUIT ) + goto Error; + + CUDART_CHECK( cudaMallocArray( &pArrayImage, &desc, w, h ) ); + CUDART_CHECK( cudaMallocArray( &pArrayTemplate, &desc, w, h ) ); + CUDART_CHECK( cudaMemcpyToArray( pArrayImage, 0, 0, hidata, w*h, cudaMemcpyHostToDevice ) ); + + CUDART_CHECK( cudaMemcpy2DArrayToArray( pArrayTemplate, 0, 0, pArrayImage, 0, 0, w, h, cudaMemcpyDeviceToDevice ) ); + + CUDART_CHECK( cudaBindTextureToArray( texImage, pArrayImage ) ); + CUDART_CHECK( cudaBindTextureToArray( texTemplate, pArrayTemplate ) ); + + CopyToTemplate( didata, DevicePitch, + xTemplate, yTemplate, + wTemplate, hTemplate, + xOffset, yOffset ); + + corrCPU( hoCorrCPU, hoCorrCPUI, hoCorrCPUISq, hoCorrCPUIT, + w*sizeof(float), wTemplate*hTemplate, xTemplate-xOffset, yTemplate-yOffset, w, h, + hidata, HostPitch, hidata, HostPitch ); + + // height of thread block must be >= hTemplate + wTile = 32; + threads = dim3(32,8); + blocks = dim3(w/wTile+(0!=w%wTile),h/threads.y+(0!=h%threads.y)); + + sharedPitch = ~63&(wTile+wTemplate+63); + sharedMem = sharedPitch*(threads.y+hTemplate); + +#define TEST_VECTOR( baseName, bPrintNeighborhood, cIterations, outfile ) \ + { \ + double pixelsPerSecond; \ + double templatePixelsPerSecond; \ + if ( ! TestCorrelation( &pixelsPerSecond, \ + &templatePixelsPerSecond, \ + xOffset, yOffset, \ + w, h, \ + hoCorrCPU, \ + hoCorrCPUI, \ + hoCorrCPUISq, \ + hoCorrCPUIT, \ + xTemplate-xOffset, yTemplate-yOffset, \ + wTemplate, hTemplate, \ + wTile, sharedPitch, sharedMem, \ + threads, blocks, \ + baseName##Sums, \ + baseName, \ + bPrintNeighborhood, cIterations, outfile ) ) { \ + printf( "Error\n" ); \ + } \ + printf( "%s: %.2f Mpix/s\t%.2fGtpix/s\n", \ + #baseName, pixelsPerSecond/1e6, templatePixelsPerSecond/1e9 ); \ + } + + TEST_VECTOR( corrShared, false, 100, NULL ); + + // height of thread block must be >= hTemplate + wTile = 32; + threads = dim3(32,8); + blocks = dim3(w/wTile+(0!=w%wTile),h/threads.y+(0!=h%threads.y)); + + sharedPitch = ~63&(((wTile+wTemplate)+63)); + sharedMem = sharedPitch*(threads.y+hTemplate); + + TEST_VECTOR( corrSharedSM, false, 100, NULL ); + + TEST_VECTOR( corrShared4, false, 100, NULL ); + + // set up blocking parameters for 2D tex-constant formulation + threads.x = 32; threads.y = 16; threads.z = 1; + blocks.x = INTCEIL(w,threads.x); blocks.y = INTCEIL(h,threads.y); blocks.z = 1; + TEST_VECTOR( corrTexConstant, false, 100, NULL ); + + if ( outputFilename ) { + printf( "Writing graymap of correlation values to %s\n", outputFilename ); + } + + // set up blocking parameters for 2D tex-tex formulation + threads.x = 16; threads.y = 8; threads.z = 1; + blocks.x = INTCEIL(w,threads.x); blocks.y = INTCEIL(h,threads.y); blocks.z = 1; + TEST_VECTOR( corrTexTex, false, 100, outputFilename ); + + ret = 0; +Error: + free( hoCorrCPU ); + free( hoCorrCPUI ); + free( hoCorrCPUISq ); + free( hoCorrCPUIT ); + + free( hidata ); + + cudaFree(didata); + + cudaFreeArray(pArrayImage); + cudaFreeArray(pArrayTemplate); + + return ret; + +} diff --git a/corr/pgm.cu b/corr/pgm.cu new file mode 100644 index 0000000..80fccec --- /dev/null +++ b/corr/pgm.cu @@ -0,0 +1,133 @@ +/* + * + * pgm.cu + * + * Functions to load and store PGM (portable gray map) files. + * + * Copyright (c) 2011-2012, Archaea Software, LLC. + * All rights reserved. + + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * + * 1. Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * 2. Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in + * the documentation and/or other materials provided with the + * distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT + * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS + * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE + * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, + * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, + * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; + * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER + * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT + * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN + * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + * + */ + +#include +#include +#include +#include +#include "pgm.h" + +int +pgmLoad( + const char *filename, + unsigned char **pHostData, unsigned int *pHostPitch, + unsigned char **pDeviceData, unsigned int *pDevicePitch, + int *pWidth, int *pHeight, int padWidth, int padHeight) +{ + int ret = 1; + const int hsize = 0x40; + int w, h; + FILE *fp = NULL; + int maxval; + char header[hsize]; + unsigned char *idata = NULL; + unsigned char *ddata = NULL; + size_t dPitch; + + fp = fopen( filename, "rb" ); + if ( fp == NULL) { + fprintf( stderr, "Failed to open %s.\n", filename ); + goto Error; + } + + if (NULL == fgets(header, hsize, fp)) { + fprintf(stderr, "Invalid PGM file.\n"); + goto Error; + } + + if ( strncmp(header, "P5", 2) ) { + fprintf(stderr, "File is not a PGM image.\n"); + goto Error; + } + if ( 1 != fscanf( fp, "%d", &w ) ) + goto Error; + if ( 1 != fscanf( fp, "%d", &h ) ) + goto Error; + if ( 1 != fscanf( fp, "%d", &maxval ) ) + goto Error; + if ( padWidth == 0 && padHeight == 0 ) { + padWidth = w; + padHeight = h; + } + idata = (unsigned char *) malloc( padWidth*padHeight ); + if ( ! idata ) + goto Error; +/* if (((size_t)w*h) != fread(idata, sizeof(unsigned char), w*h, fp)) + goto Error;*/ + for ( int row = 0; row < h; row++ ) { + if ( (size_t) w != fread( idata+row*padWidth, 1, w, fp ) ) + goto Error; + } + if ( cudaSuccess != cudaMallocPitch( (void **) &ddata, &dPitch, padWidth, padHeight ) ) + goto Error; + *pWidth = padWidth; + *pHeight = padHeight; + *pHostPitch = padWidth; + *pHostData = idata; + *pDeviceData = ddata; + *pDevicePitch = (unsigned int) dPitch; + cudaMemcpy2D( ddata, dPitch, idata, padWidth, padWidth, padHeight, cudaMemcpyHostToDevice ); + fclose(fp); + return 0; +Error: + free( idata ); + cudaFree( ddata ); + if ( fp ) { + fclose( fp ); + } + return ret; +} + +int +pgmSave(const char* filename, unsigned char *data, int w, int h) +{ + int ret = 1; + FILE *fp = fopen( filename, "wb" ); + if ( NULL == fp ) { + fprintf( stderr, "Failed to open %s\n", filename ); + goto Error; + } + + fprintf( fp, "P5\n%d\n%d\n%d\n", w, h, 0xff ); + if ( w*h != fwrite(data, sizeof(unsigned char), w*h, fp) ) { + fprintf( stderr, "Write failed\n" ); + goto Error; + } + + fclose(fp); + ret = 0; +Error: + return ret; +} diff --git a/corr/pgm.h b/corr/pgm.h new file mode 100644 index 0000000..6198974 --- /dev/null +++ b/corr/pgm.h @@ -0,0 +1,10 @@ +#ifndef _IMAGE_H_ +#define _IMAGE_H_ + +int pgmLoad( const char *filename, + unsigned char **pHostData, unsigned int *pSysPitch, + unsigned char **pDeviceData, unsigned int *pDevPitch, + int *pWidth, int *pHeight, int padWidth=0, int padHeight=0 ); +int pgmSave( const char *filename, unsigned char *data, int w, int h); + +#endif diff --git a/corr/values.bmp b/corr/values.bmp new file mode 100644 index 0000000..01d841c Binary files /dev/null and b/corr/values.bmp differ