kernel_gpu_opencl.cl 73 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540541542543544545546547548549550551552553554555556557558559560561562563564565566567568569570571572573574575576577578579580581582583584585586587588589590591592593594595596597598599600601602603604605606607608609610611612613614615616617618619620621622623624625626627628629630631632633634635636637638639640641642643644645646647648649650651652653654655656657658659660661662663664665666667668669670671672673674675676677678679680681682683684685686687688689690691692693694695696697698699700701702703704705706707708709710711712713714715716717718719720721722723724725726727728729730731732733734735736737738739740741742743744745746747748749750751752753754755756757758759760761762763764765766767768769770771772773774775776777778779780781782783784785786787788789790791792793794795796797798799800801802803804805806807808809810811812813814815816817818819820821822823824825826827828829830831832833834835836837838839840841842843844845846847848849850851852853854855856857858859860861862863864865866867868869870871872873874875876877878879880881882883884885886887888889890891892893894895896897898899900901902903904905906907908909910911912913914915916917918919920921922923924925926927928929930931932933934935936937938939940941942943944945946947948949950951952953954955956957958959960961962963964965966967968969970971972973974975976977978979980981982983984985986987988989990991992993994995996997998999100010011002100310041005100610071008100910101011101210131014101510161017101810191020102110221023102410251026102710281029103010311032103310341035103610371038103910401041104210431044104510461047104810491050105110521053105410551056105710581059106010611062106310641065106610671068106910701071107210731074107510761077107810791080108110821083108410851086108710881089109010911092109310941095109610971098109911001101110211031104110511061107110811091110111111121113111411151116111711181119112011211122112311241125112611271128112911301131113211331134113511361137113811391140114111421143114411451146114711481149115011511152115311541155115611571158115911601161116211631164116511661167116811691170117111721173117411751176117711781179118011811182118311841185118611871188118911901191119211931194119511961197119811991200120112021203120412051206120712081209121012111212121312141215121612171218121912201221122212231224122512261227122812291230123112321233123412351236123712381239124012411242124312441245124612471248124912501251125212531254125512561257125812591260126112621263126412651266126712681269127012711272127312741275127612771278127912801281128212831284128512861287128812891290129112921293129412951296129712981299130013011302130313041305130613071308130913101311131213131314131513161317131813191320132113221323132413251326132713281329133013311332133313341335133613371338133913401341134213431344134513461347134813491350135113521353135413551356135713581359136013611362136313641365136613671368136913701371137213731374137513761377137813791380138113821383138413851386138713881389139013911392139313941395139613971398139914001401140214031404140514061407140814091410141114121413141414151416141714181419142014211422142314241425142614271428142914301431143214331434143514361437143814391440144114421443144414451446144714481449145014511452145314541455145614571458145914601461146214631464146514661467146814691470147114721473147414751476147714781479148014811482148314841485148614871488148914901491149214931494149514961497149814991500150115021503150415051506150715081509151015111512151315141515151615171518151915201521152215231524152515261527152815291530153115321533153415351536153715381539154015411542154315441545154615471548154915501551155215531554155515561557155815591560156115621563156415651566156715681569157015711572157315741575157615771578157915801581158215831584158515861587158815891590159115921593159415951596159715981599160016011602160316041605160616071608160916101611161216131614161516161617161816191620162116221623162416251626162716281629163016311632163316341635163616371638163916401641164216431644164516461647164816491650165116521653165416551656165716581659166016611662166316641665166616671668166916701671167216731674167516761677167816791680168116821683168416851686168716881689169016911692169316941695169616971698169917001701170217031704170517061707170817091710171117121713171417151716171717181719172017211722172317241725172617271728172917301731173217331734173517361737173817391740174117421743174417451746174717481749175017511752175317541755175617571758175917601761176217631764176517661767176817691770177117721773177417751776177717781779178017811782178317841785178617871788178917901791179217931794179517961797179817991800180118021803180418051806180718081809181018111812181318141815181618171818181918201821182218231824182518261827182818291830183118321833183418351836183718381839184018411842184318441845184618471848184918501851185218531854185518561857185818591860186118621863186418651866186718681869187018711872187318741875187618771878187918801881188218831884188518861887188818891890189118921893189418951896189718981899190019011902190319041905190619071908190919101911191219131914191519161917191819191920192119221923192419251926192719281929193019311932193319341935193619371938193919401941194219431944194519461947194819491950195119521953195419551956195719581959196019611962196319641965196619671968196919701971197219731974197519761977197819791980198119821983198419851986198719881989199019911992199319941995199619971998199920002001200220032004200520062007200820092010201120122013201420152016201720182019202020212022202320242025202620272028202920302031203220332034203520362037203820392040204120422043204420452046204720482049205020512052205320542055205620572058205920602061206220632064206520662067206820692070207120722073207420752076207720782079208020812082208320842085208620872088208920902091209220932094209520962097209820992100210121022103210421052106210721082109211021112112211321142115211621172118211921202121212221232124212521262127212821292130213121322133213421352136213721382139214021412142214321442145214621472148214921502151215221532154215521562157215821592160216121622163216421652166216721682169217021712172217321742175217621772178217921802181218221832184218521862187218821892190219121922193219421952196219721982199220022012202220322042205220622072208220922102211221222132214221522162217221822192220222122222223222422252226222722282229223022312232223322342235
  1. //========================================================================================================================================================================================================200
  2. // DEFINE / INCLUDE
  3. //========================================================================================================================================================================================================200
  4. //======================================================================================================================================================150
  5. // MAIN FUNCTION HEADER
  6. //======================================================================================================================================================150
  7. #include "./main.h"
  8. //======================================================================================================================================================150
  9. // End
  10. //======================================================================================================================================================150
  11. //========================================================================================================================================================================================================200
  12. // KERNEL
  13. //========================================================================================================================================================================================================200
  14. __kernel void
  15. kernel_gpu_opencl( // structures
  16. params_common d_common, // 0
  17. // common_change
  18. __global fp* d_frame, // 1 INPUT
  19. int d_frame_no, // 2 INPUT
  20. // common
  21. __global int* d_endoRow, // 3 INPUT
  22. __global int* d_endoCol, // 4 INPUT
  23. __global int* d_tEndoRowLoc, // 5 OUTPUT common.endoPoints * common.no_frames
  24. __global int* d_tEndoColLoc, // 6 OUTPUT common.endoPoints * common.no_frames
  25. __global int* d_epiRow, // 7 INPUT
  26. __global int* d_epiCol, // 8 INPUT
  27. __global int* d_tEpiRowLoc, // 9 OUTPUT common.epiPoints * common.no_frames
  28. __global int* d_tEpiColLoc, // 10 OUTPUT common.epiPoints * common.no_frames
  29. // common_unique
  30. __global fp* d_endoT, // 11 OUTPUT common.in_elem * common.endoPoints
  31. __global fp* d_epiT, // 12 OUTPUT common.in_elem * common.epiPoints
  32. __global fp* d_in2_all, // 13 OUTPUT common.in2_elem * common.allPoints
  33. __global fp* d_conv_all, // 14 OUTPUT common.conv_elem * common.allPoints
  34. __global fp* d_in2_pad_cumv_all, // 15 OUTPUT common.in2_pad_cumv_elem * common.allPoints
  35. __global fp* d_in2_pad_cumv_sel_all, // 16 OUTPUT common.in2_pad_cumv_sel_elem * common.allPoints
  36. __global fp* d_in2_sub_cumh_all, // 17 OUTPUT common.in2_sub_cumh_elem * common.allPoints
  37. __global fp* d_in2_sub_cumh_sel_all, // 18 OUTPUT common.in2_sub_cumh_sel_elem * common.allPoints
  38. __global fp* d_in2_sub2_all, // 19 OUTPUT common.in2_sub2_elem * common.allPoints
  39. __global fp* d_in2_sqr_all, // 20 OUTPUT common.in2_elem * common.allPoints
  40. __global fp* d_in2_sqr_sub2_all, // 21 OUTPUT common.in2_sub2_elem * common.allPoints
  41. __global fp* d_in_sqr_all, // 22 OUTPUT common.in_elem * common.allPoints
  42. __global fp* d_tMask_all, // 23 OUTPUT common.tMask_elem * common.allPoints
  43. __global fp* d_mask_conv_all, // 24 OUTPUT common.mask_conv_elem * common.allPoints
  44. // // local
  45. // __local fp* d_in_mod_temp, // 25 OUTPUT common.in_elem
  46. // __local fp* in_partial_sum, // 26 OUTPUT common.in_cols
  47. // __local fp* in_sqr_partial_sum, // 27 OUTPUT common.in_sqr_rows
  48. // __local fp* par_max_val, // 28 OUTPUT common.mask_conv_rows
  49. // __local int* par_max_coo) // 29 OUTPUT common.mask_conv_rows
  50. // local
  51. __global fp* d_in_mod_temp_all, // 25 OUTPUT common.in_elem * common.allPoints
  52. __global fp* in_partial_sum_all, // 26 OUTPUT common.in_cols * common.allPoints
  53. __global fp* in_sqr_partial_sum_all, // 27 OUTPUT common.in_sqr_rows * common.allPoints
  54. __global fp* par_max_val_all, // 28 OUTPUT common.mask_conv_rows * common.allPoints
  55. __global int* par_max_coo_all, // 29 OUTPUT common.mask_conv_rows * common.allPoints
  56. __global fp* in_final_sum_all, // 30 OUTPUT common.allPoints
  57. __global fp* in_sqr_final_sum_all, // 31 OUTPUT common.allPoints
  58. __global fp* denomT_all, // 32 OUTPUT common.allPoints
  59. __global fp* checksum) // 33 OUTPUT 100
  60. {
  61. //======================================================================================================================================================150
  62. // COMMON VARIABLES
  63. //======================================================================================================================================================150
  64. // __global fp* d_in;
  65. int rot_row;
  66. int rot_col;
  67. int in2_rowlow;
  68. int in2_collow;
  69. int ic;
  70. int jc;
  71. int jp1;
  72. int ja1, ja2;
  73. int ip1;
  74. int ia1, ia2;
  75. int ja, jb;
  76. int ia, ib;
  77. fp s;
  78. int i;
  79. int j;
  80. int row;
  81. int col;
  82. int ori_row;
  83. int ori_col;
  84. int position;
  85. fp sum;
  86. int pos_ori;
  87. fp temp;
  88. fp temp2;
  89. int location;
  90. int cent;
  91. int tMask_row;
  92. int tMask_col;
  93. fp largest_value_current = 0;
  94. fp largest_value = 0;
  95. int largest_coordinate_current = 0;
  96. int largest_coordinate = 0;
  97. fp fin_max_val = 0;
  98. int fin_max_coo = 0;
  99. int largest_row;
  100. int largest_col;
  101. int offset_row;
  102. int offset_col;
  103. fp mean;
  104. fp mean_sqr;
  105. fp variance;
  106. fp deviation;
  107. int pointer;
  108. int ori_pointer;
  109. int loc_pointer;
  110. // __local fp in_final_sum;
  111. // __local fp in_sqr_final_sum;
  112. // __local fp denomT;
  113. //======================================================================================================================================================150
  114. // BLOCK/THREAD IDs
  115. //======================================================================================================================================================150
  116. int bx = get_group_id(0); // get current horizontal block index (0-n)
  117. int tx = get_local_id(0); // get current horizontal thread index (0-n)
  118. int ei_new;
  119. //======================================================================================================================================================150
  120. // UNIQUE STRUCTURE RECONSTRUCTED HERE
  121. //======================================================================================================================================================150
  122. // common
  123. __global fp* d_common_change_d_frame = &d_frame[0];
  124. // offsets for either endo or epi points (separate arrays for endo and epi points)
  125. int d_unique_point_no;
  126. __global int* d_unique_d_Row;
  127. __global int* d_unique_d_Col;
  128. __global int* d_unique_d_tRowLoc;
  129. __global int* d_unique_d_tColLoc;
  130. __global fp* d_in;
  131. if(bx < d_common.endoPoints){
  132. d_unique_point_no = bx; // endo point number 0-???
  133. d_unique_d_Row = d_endoRow; // initial endo row coordinates
  134. d_unique_d_Col = d_endoCol; // initial endo col coordinates
  135. d_unique_d_tRowLoc = d_tEndoRowLoc; // all endo row coordinates
  136. d_unique_d_tColLoc = d_tEndoColLoc; // all endo col coordinates
  137. d_in = &d_endoT[d_unique_point_no * d_common.in_elem]; // endo templates
  138. }
  139. else{
  140. d_unique_point_no = bx-d_common.endoPoints; // epi point number 0-???
  141. d_unique_d_Row = d_epiRow; // initial epi row coordinates
  142. d_unique_d_Col = d_epiCol; // initial epi col coordinates
  143. d_unique_d_tRowLoc = d_tEpiRowLoc; // all epi row coordinates
  144. d_unique_d_tColLoc = d_tEpiColLoc; // all epi col coordinates
  145. d_in = &d_epiT[d_unique_point_no * d_common.in_elem]; // epi templates
  146. }
  147. // offsets for all points (one array for all points)
  148. __global fp* d_unique_d_in2 = &d_in2_all[bx*d_common.in2_elem];
  149. __global fp* d_unique_d_conv = &d_conv_all[bx*d_common.conv_elem];
  150. __global fp* d_unique_d_in2_pad_cumv = &d_in2_pad_cumv_all[bx*d_common.in2_pad_cumv_elem];
  151. __global fp* d_unique_d_in2_pad_cumv_sel = &d_in2_pad_cumv_sel_all[bx*d_common.in2_pad_cumv_sel_elem];
  152. __global fp* d_unique_d_in2_sub_cumh = &d_in2_sub_cumh_all[bx*d_common.in2_sub_cumh_elem];
  153. __global fp* d_unique_d_in2_sub_cumh_sel = &d_in2_sub_cumh_sel_all[bx*d_common.in2_sub_cumh_sel_elem];
  154. __global fp* d_unique_d_in2_sub2 = &d_in2_sub2_all[bx*d_common.in2_sub2_elem];
  155. __global fp* d_unique_d_in2_sqr = &d_in2_sqr_all[bx*d_common.in2_sqr_elem];
  156. __global fp* d_unique_d_in2_sqr_sub2 = &d_in2_sqr_sub2_all[bx*d_common.in2_sqr_sub2_elem];
  157. __global fp* d_unique_d_in_sqr = &d_in_sqr_all[bx*d_common.in_sqr_elem];
  158. __global fp* d_unique_d_tMask = &d_tMask_all[bx*d_common.tMask_elem];
  159. __global fp* d_unique_d_mask_conv = &d_mask_conv_all[bx*d_common.mask_conv_elem];
  160. // used to be local
  161. __global fp* d_in_mod_temp = &d_in_mod_temp_all[bx*d_common.in_elem];
  162. __global fp* in_partial_sum = &in_partial_sum_all[bx*d_common.in_cols];
  163. __global fp* in_sqr_partial_sum = &in_sqr_partial_sum_all[bx*d_common.in_sqr_rows];
  164. __global fp* par_max_val = &par_max_val_all[bx*d_common.mask_conv_rows];
  165. __global int* par_max_coo = &par_max_coo_all[bx*d_common.mask_conv_rows];
  166. __global fp* in_final_sum = &in_final_sum_all[bx];
  167. __global fp* in_sqr_final_sum = &in_sqr_final_sum_all[bx];
  168. __global fp* denomT = &denomT_all[bx];
  169. //======================================================================================================================================================150
  170. // END
  171. //======================================================================================================================================================150
  172. //======================================================================================================================================================150
  173. // Initialize checksum
  174. //======================================================================================================================================================150
  175. #ifdef TEST_CHECKSUM
  176. if(bx==0 && tx==0){
  177. for(i=0; i<CHECK; i++){
  178. checksum[i] = 0;
  179. }
  180. }
  181. #endif
  182. //======================================================================================================================================================150
  183. // INITIAL COORDINATE AND TEMPLATE UPDATE
  184. //======================================================================================================================================================150
  185. // generate templates based on the first frame only
  186. if(d_frame_no == 0){
  187. //====================================================================================================100
  188. // Initialize cross-frame variables
  189. //====================================================================================================100
  190. #ifdef INIT
  191. // only the first thread initializes
  192. if(tx==0){
  193. // this block and for all frames
  194. for(i=0; i<d_common.no_frames; i++){
  195. d_unique_d_tRowLoc[d_unique_point_no*d_common.no_frames+i] = 0;
  196. d_unique_d_tColLoc[d_unique_point_no*d_common.no_frames+i] = 0;
  197. }
  198. // this block
  199. for(i=0; i<d_common.in_elem; i++){
  200. d_in[i] = 0;
  201. }
  202. }
  203. #endif
  204. //====================================================================================================100
  205. // SYNCHRONIZE THREADS
  206. //====================================================================================================100
  207. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  208. //====================================================================================================100
  209. // UPDATE ROW LOC AND COL LOC
  210. //====================================================================================================100
  211. // uptade temporary endo/epi row/col coordinates (in each block corresponding to point, narrow work to one thread)
  212. ei_new = tx;
  213. if(ei_new == 0){
  214. // update temporary row/col coordinates
  215. pointer = d_unique_point_no*d_common.no_frames+d_frame_no;
  216. d_unique_d_tRowLoc[pointer] = d_unique_d_Row[d_unique_point_no];
  217. d_unique_d_tColLoc[pointer] = d_unique_d_Col[d_unique_point_no];
  218. }
  219. //====================================================================================================100
  220. // CREATE TEMPLATES
  221. //====================================================================================================100
  222. // work
  223. ei_new = tx;
  224. while(ei_new < d_common.in_elem){
  225. // figure out row/col location in new matrix
  226. row = (ei_new+1) % d_common.in_rows - 1; // (0-n) row
  227. col = (ei_new+1) / d_common.in_rows + 1 - 1; // (0-n) column
  228. if((ei_new+1) % d_common.in_rows == 0){
  229. row = d_common.in_rows - 1;
  230. col = col-1;
  231. }
  232. // figure out row/col location in corresponding new template area in image and give to every thread (get top left corner and progress down and right)
  233. ori_row = d_unique_d_Row[d_unique_point_no] - 25 + row - 1;
  234. ori_col = d_unique_d_Col[d_unique_point_no] - 25 + col - 1;
  235. ori_pointer = ori_col*d_common.frame_rows+ori_row;
  236. // update template
  237. d_in[col*d_common.in_rows+row] = d_common_change_d_frame[ori_pointer];
  238. // go for second round
  239. ei_new = ei_new + NUMBER_THREADS;
  240. }
  241. //====================================================================================================100
  242. // SYNCHRONIZE THREADS
  243. //====================================================================================================100
  244. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  245. //====================================================================================================100
  246. // CHECKSUM
  247. //====================================================================================================100
  248. #ifdef TEST_CHECKSUM
  249. if(bx==0 && tx==0){
  250. for(i=0; i<d_common.in_elem; i++){
  251. checksum[0] = checksum[0]+d_in[i];
  252. }
  253. }
  254. //====================================================================================================100
  255. // SYNCHRONIZE THREADS
  256. //====================================================================================================100
  257. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  258. #endif
  259. //====================================================================================================100
  260. // End
  261. //====================================================================================================100
  262. }
  263. //======================================================================================================================================================150
  264. // PROCESS POINTS
  265. //======================================================================================================================================================150
  266. // process points in all frames except for the first one
  267. if(d_frame_no != 0){
  268. //====================================================================================================100
  269. // Initialize frame-specific variables
  270. //====================================================================================================100
  271. #ifdef INIT
  272. // only the first thread initializes
  273. if(tx==0){
  274. // this block
  275. for(i=0; i<d_common.in2_elem; i++){
  276. d_unique_d_in2[i] = 0;
  277. }
  278. for(i=0; i<d_common.conv_elem; i++){
  279. d_unique_d_conv[i] = 0;
  280. }
  281. for(i=0; i<d_common.in2_pad_cumv_elem; i++){
  282. d_unique_d_in2_pad_cumv[i] = 0;
  283. }
  284. for(i=0; i<d_common.in2_pad_cumv_sel_elem; i++){
  285. d_unique_d_in2_pad_cumv_sel[i] = 0;
  286. }
  287. for(i=0; i<d_common.in2_sub_cumh_elem; i++){
  288. d_unique_d_in2_sub_cumh[i] = 0;
  289. }
  290. for(i=0; i<d_common.in2_sub_cumh_sel_elem; i++){
  291. d_unique_d_in2_sub_cumh_sel[i] = 0;
  292. }
  293. for(i=0; i<d_common.in2_sub2_elem; i++){
  294. d_unique_d_in2_sub2[i] = 0;
  295. }
  296. for(i=0; i<d_common.in2_sqr_elem; i++){
  297. d_unique_d_in2_sqr[i] = 0;
  298. }
  299. for(i=0; i<d_common.in2_sqr_sub2_elem; i++){
  300. d_unique_d_in2_sqr_sub2[i] = 0;
  301. }
  302. for(i=0; i<d_common.in_sqr_elem; i++){
  303. d_unique_d_in_sqr[i] = 0;
  304. }
  305. for(i=0; i<d_common.tMask_elem; i++){
  306. d_unique_d_tMask[i] = 0;
  307. }
  308. for(i=0; i<d_common.mask_conv_elem; i++){
  309. d_unique_d_mask_conv[i] = 0;
  310. }
  311. for(i=0; i<d_common.in_elem; i++){
  312. d_in_mod_temp[i] = 0;
  313. }
  314. for(i=0; i<d_common.in_cols; i++){
  315. in_partial_sum[i] = 0;
  316. }
  317. for(i=0; i<d_common.in_sqr_rows; i++){
  318. in_sqr_partial_sum[i] = 0;
  319. }
  320. for(i=0; i<d_common.mask_conv_rows; i++){
  321. par_max_val[i] = 0;
  322. }
  323. for(i=0; i<d_common.mask_conv_rows; i++){
  324. par_max_coo[i] = 0;
  325. }
  326. in_final_sum[0] = 0;
  327. in_sqr_final_sum[0] = 0;
  328. denomT[0] = 0;
  329. }
  330. #endif
  331. //====================================================================================================100
  332. // SYNCHRONIZE THREADS
  333. //====================================================================================================100
  334. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  335. //====================================================================================================100
  336. // SELECTION
  337. //====================================================================================================100
  338. in2_rowlow = d_unique_d_Row[d_unique_point_no] - d_common.sSize; // (1 to n+1)
  339. in2_collow = d_unique_d_Col[d_unique_point_no] - d_common.sSize;
  340. // work
  341. ei_new = tx;
  342. while(ei_new < d_common.in2_elem){
  343. // figure out row/col location in new matrix
  344. row = (ei_new+1) % d_common.in2_rows - 1; // (0-n) row
  345. col = (ei_new+1) / d_common.in2_rows + 1 - 1; // (0-n) column
  346. if((ei_new+1) % d_common.in2_rows == 0){
  347. row = d_common.in2_rows - 1;
  348. col = col-1;
  349. }
  350. // figure out corresponding location in old matrix and copy values to new matrix
  351. ori_row = row + in2_rowlow - 1;
  352. ori_col = col + in2_collow - 1;
  353. d_unique_d_in2[ei_new] = d_common_change_d_frame[ori_col*d_common.frame_rows+ori_row];
  354. // go for second round
  355. ei_new = ei_new + NUMBER_THREADS;
  356. }
  357. //====================================================================================================100
  358. // SYNCHRONIZE THREADS
  359. //====================================================================================================100
  360. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  361. //====================================================================================================100
  362. // CHECKSUM
  363. //====================================================================================================100
  364. #ifdef TEST_CHECKSUM
  365. if(bx==0 && tx==0){
  366. for(i=0; i<d_common.in2_elem; i++){
  367. checksum[1] = checksum[1]+d_unique_d_in2[i];
  368. }
  369. }
  370. //====================================================================================================100
  371. // SYNCHRONIZE THREADS
  372. //====================================================================================================100
  373. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  374. #endif
  375. //====================================================================================================100
  376. // CONVOLUTION
  377. //====================================================================================================100
  378. //==================================================50
  379. // ROTATION
  380. //==================================================50
  381. // work
  382. ei_new = tx;
  383. while(ei_new < d_common.in_elem){
  384. // while(ei_new < 1){
  385. // figure out row/col location in padded array
  386. row = (ei_new+1) % d_common.in_rows - 1; // (0-n) row
  387. col = (ei_new+1) / d_common.in_rows + 1 - 1; // (0-n) column
  388. if((ei_new+1) % d_common.in_rows == 0){
  389. row = d_common.in_rows - 1;
  390. col = col-1;
  391. }
  392. // execution
  393. rot_row = (d_common.in_rows-1) - row;
  394. rot_col = (d_common.in_rows-1) - col;
  395. d_in_mod_temp[ei_new] = d_in[rot_col*d_common.in_rows+rot_row];
  396. // d_in_mod_temp[ei_new] = d_in[0];
  397. // d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_in_pointer];
  398. // d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_point_no * d_common.in_elem];
  399. // d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_point_no * 2601];
  400. // if((d_unique_point_no * d_common.in_elem) > (2601*51
  401. // printf("frame_no IS %d\n", d_common_change[0].frame_no);
  402. // d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_point_no];
  403. // d_in_mod_temp[ei_new] = 1;
  404. // kot = d_in[rot_col*d_common.in_rows+rot_row];
  405. // d_in_mod_temp[ei_new] = kot;
  406. // d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_in_pointer+rot_col*d_common.in_rows+rot_row];
  407. // d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_point_no * d_common.in_elem+rot_col*d_common.in_rows+rot_row];
  408. //d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_point_no];
  409. // d_unique_d_T[d_unique_in_pointer+rot_col*d_common.in_rows+rot_row] = 1;
  410. // d_unique_d_T[d_unique_in_pointer] = 1;
  411. // d_endoT[d_unique_in_pointer] = 1;
  412. // d_in_mod_temp[ei_new] = 1;
  413. // go for second round
  414. ei_new = ei_new + NUMBER_THREADS;
  415. }
  416. //==================================================50
  417. // SYNCHRONIZE THREADS
  418. //==================================================50
  419. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  420. //==================================================50
  421. // CHECKSUM
  422. //==================================================50
  423. #ifdef TEST_CHECKSUM
  424. if(bx==0 && tx==0){
  425. for(i=0; i<d_common.in_elem; i++){
  426. checksum[2] = checksum[2]+d_in_mod_temp[i];
  427. }
  428. }
  429. //==================================================50
  430. // SYNCHRONIZE THREADS
  431. //==================================================50
  432. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  433. #endif
  434. //==================================================50
  435. // ACTUAL CONVOLUTION
  436. //==================================================50
  437. // work
  438. ei_new = tx;
  439. while(ei_new < d_common.conv_elem){
  440. // figure out row/col location in array
  441. ic = (ei_new+1) % d_common.conv_rows; // (1-n)
  442. jc = (ei_new+1) / d_common.conv_rows + 1; // (1-n)
  443. if((ei_new+1) % d_common.conv_rows == 0){
  444. ic = d_common.conv_rows;
  445. jc = jc-1;
  446. }
  447. //
  448. j = jc + d_common.joffset;
  449. jp1 = j + 1;
  450. if(d_common.in2_cols < jp1){
  451. ja1 = jp1 - d_common.in2_cols;
  452. }
  453. else{
  454. ja1 = 1;
  455. }
  456. if(d_common.in_cols < j){
  457. ja2 = d_common.in_cols;
  458. }
  459. else{
  460. ja2 = j;
  461. }
  462. i = ic + d_common.ioffset;
  463. ip1 = i + 1;
  464. if(d_common.in2_rows < ip1){
  465. ia1 = ip1 - d_common.in2_rows;
  466. }
  467. else{
  468. ia1 = 1;
  469. }
  470. if(d_common.in_rows < i){
  471. ia2 = d_common.in_rows;
  472. }
  473. else{
  474. ia2 = i;
  475. }
  476. s = 0;
  477. for(ja=ja1; ja<=ja2; ja++){
  478. jb = jp1 - ja;
  479. for(ia=ia1; ia<=ia2; ia++){
  480. ib = ip1 - ia;
  481. s = s + d_in_mod_temp[d_common.in_rows*(ja-1)+ia-1] * d_unique_d_in2[d_common.in2_rows*(jb-1)+ib-1];
  482. }
  483. }
  484. //d_unique_d_conv[d_common.conv_rows*(jc-1)+ic-1] = s;
  485. d_unique_d_conv[ei_new] = s;
  486. // go for second round
  487. ei_new = ei_new + NUMBER_THREADS;
  488. }
  489. //==================================================50
  490. // SYNCHRONIZE THREADS
  491. //==================================================50
  492. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  493. //==================================================50
  494. // CHECKSUM
  495. //==================================================50
  496. #ifdef TEST_CHECKSUM
  497. if(bx==0 && tx==0){
  498. for(i=0; i<d_common.conv_elem; i++){
  499. checksum[3] = checksum[3]+d_unique_d_conv[i];
  500. }
  501. }
  502. //==================================================50
  503. // SYNCHRONIZE THREADS
  504. //==================================================50
  505. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  506. #endif
  507. //==================================================50
  508. // End
  509. //==================================================50
  510. //====================================================================================================100
  511. // CUMULATIVE SUM (LOCAL)
  512. //====================================================================================================100
  513. //==================================================50
  514. // PADD ARRAY
  515. //==================================================50
  516. // work
  517. ei_new = tx;
  518. while(ei_new < d_common.in2_pad_cumv_elem){
  519. // figure out row/col location in padded array
  520. row = (ei_new+1) % d_common.in2_pad_cumv_rows - 1; // (0-n) row
  521. col = (ei_new+1) / d_common.in2_pad_cumv_rows + 1 - 1; // (0-n) column
  522. if((ei_new+1) % d_common.in2_pad_cumv_rows == 0){
  523. row = d_common.in2_pad_cumv_rows - 1;
  524. col = col-1;
  525. }
  526. // execution
  527. if( row > (d_common.in2_pad_add_rows-1) && // do if has numbers in original array
  528. row < (d_common.in2_pad_add_rows+d_common.in2_rows) &&
  529. col > (d_common.in2_pad_add_cols-1) &&
  530. col < (d_common.in2_pad_add_cols+d_common.in2_cols)){
  531. ori_row = row - d_common.in2_pad_add_rows;
  532. ori_col = col - d_common.in2_pad_add_cols;
  533. d_unique_d_in2_pad_cumv[ei_new] = d_unique_d_in2[ori_col*d_common.in2_rows+ori_row];
  534. }
  535. else{ // do if otherwise
  536. d_unique_d_in2_pad_cumv[ei_new] = 0;
  537. }
  538. // go for second round
  539. ei_new = ei_new + NUMBER_THREADS;
  540. }
  541. //==================================================50
  542. // SYNCHRONIZE THREADS
  543. //==================================================50
  544. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  545. //==================================================50
  546. // CHECKSUM
  547. //==================================================50
  548. #ifdef TEST_CHECKSUM
  549. if(bx==0 && tx==0){
  550. for(i=0; i<d_common.in2_pad_cumv_elem; i++){
  551. checksum[4] = checksum[4]+d_unique_d_in2_pad_cumv[i];
  552. }
  553. }
  554. //==================================================50
  555. // SYNCHRONIZE THREADS
  556. //==================================================50
  557. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  558. #endif
  559. //==================================================50
  560. // VERTICAL CUMULATIVE SUM
  561. //==================================================50
  562. //work
  563. ei_new = tx;
  564. while(ei_new < d_common.in2_pad_cumv_cols){
  565. // figure out column position
  566. pos_ori = ei_new*d_common.in2_pad_cumv_rows;
  567. // variables
  568. sum = 0;
  569. // loop through all rows
  570. for(position = pos_ori; position < pos_ori+d_common.in2_pad_cumv_rows; position = position + 1){
  571. d_unique_d_in2_pad_cumv[position] = d_unique_d_in2_pad_cumv[position] + sum;
  572. sum = d_unique_d_in2_pad_cumv[position];
  573. }
  574. // go for second round
  575. ei_new = ei_new + NUMBER_THREADS;
  576. }
  577. //==================================================50
  578. // SYNCHRONIZE THREADS
  579. //==================================================50
  580. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  581. //==================================================50
  582. // CHECKSUM
  583. //==================================================50
  584. #ifdef TEST_CHECKSUM
  585. if(bx==0 && tx==0){
  586. for(i=0; i<d_common.in2_pad_cumv_cols; i++){
  587. checksum[5] = checksum[5]+d_unique_d_in2_pad_cumv[i];
  588. }
  589. }
  590. //==================================================50
  591. // SYNCHRONIZE THREADS
  592. //==================================================50
  593. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  594. #endif
  595. //==================================================50
  596. // SELECTION
  597. //==================================================50
  598. // work
  599. ei_new = tx;
  600. while(ei_new < d_common.in2_pad_cumv_sel_elem){
  601. // figure out row/col location in new matrix
  602. row = (ei_new+1) % d_common.in2_pad_cumv_sel_rows - 1; // (0-n) row
  603. col = (ei_new+1) / d_common.in2_pad_cumv_sel_rows + 1 - 1; // (0-n) column
  604. if((ei_new+1) % d_common.in2_pad_cumv_sel_rows == 0){
  605. row = d_common.in2_pad_cumv_sel_rows - 1;
  606. col = col-1;
  607. }
  608. // figure out corresponding location in old matrix and copy values to new matrix
  609. ori_row = row + d_common.in2_pad_cumv_sel_rowlow - 1;
  610. ori_col = col + d_common.in2_pad_cumv_sel_collow - 1;
  611. d_unique_d_in2_pad_cumv_sel[ei_new] = d_unique_d_in2_pad_cumv[ori_col*d_common.in2_pad_cumv_rows+ori_row];
  612. // go for second round
  613. ei_new = ei_new + NUMBER_THREADS;
  614. }
  615. //==================================================50
  616. // SYNCHRONIZE THREADS
  617. //==================================================50
  618. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  619. //==================================================50
  620. // CHECKSUM
  621. //==================================================50
  622. #ifdef TEST_CHECKSUM
  623. if(bx==0 && tx==0){
  624. for(i=0; i<d_common.in2_pad_cumv_sel_elem; i++){
  625. checksum[6] = checksum[6]+d_unique_d_in2_pad_cumv_sel[i];
  626. }
  627. }
  628. //==================================================50
  629. // SYNCHRONIZE THREADS
  630. //==================================================50
  631. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  632. #endif
  633. //==================================================50
  634. // SELECTION 2
  635. //==================================================50
  636. // work
  637. ei_new = tx;
  638. while(ei_new < d_common.in2_sub_cumh_elem){
  639. // figure out row/col location in new matrix
  640. row = (ei_new+1) % d_common.in2_sub_cumh_rows - 1; // (0-n) row
  641. col = (ei_new+1) / d_common.in2_sub_cumh_rows + 1 - 1; // (0-n) column
  642. if((ei_new+1) % d_common.in2_sub_cumh_rows == 0){
  643. row = d_common.in2_sub_cumh_rows - 1;
  644. col = col-1;
  645. }
  646. // figure out corresponding location in old matrix and copy values to new matrix
  647. ori_row = row + d_common.in2_pad_cumv_sel2_rowlow - 1;
  648. ori_col = col + d_common.in2_pad_cumv_sel2_collow - 1;
  649. d_unique_d_in2_sub_cumh[ei_new] = d_unique_d_in2_pad_cumv[ori_col*d_common.in2_pad_cumv_rows+ori_row];
  650. // go for second round
  651. ei_new = ei_new + NUMBER_THREADS;
  652. }
  653. //==================================================50
  654. // SYNCHRONIZE THREADS
  655. //==================================================50
  656. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  657. //==================================================50
  658. // CHECKSUM
  659. //==================================================50
  660. #ifdef TEST_CHECKSUM
  661. if(bx==0 && tx==0){
  662. for(i=0; i<d_common.in2_sub_cumh_elem; i++){
  663. checksum[7] = checksum[7]+d_unique_d_in2_sub_cumh[i];
  664. }
  665. }
  666. #endif
  667. //==================================================50
  668. // SYNCHRONIZE THREADS
  669. //==================================================50
  670. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  671. //==================================================50
  672. // SUBTRACTION
  673. //==================================================50
  674. // work
  675. ei_new = tx;
  676. while(ei_new < d_common.in2_sub_cumh_elem){
  677. // subtract
  678. d_unique_d_in2_sub_cumh[ei_new] = d_unique_d_in2_pad_cumv_sel[ei_new] - d_unique_d_in2_sub_cumh[ei_new];
  679. // go for second round
  680. ei_new = ei_new + NUMBER_THREADS;
  681. }
  682. //==================================================50
  683. // SYNCHRONIZE THREADS
  684. //==================================================50
  685. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  686. //==================================================50
  687. // CHECKSUM
  688. //==================================================50
  689. #ifdef TEST_CHECKSUM
  690. if(bx==0 && tx==0){
  691. for(i=0; i<d_common.in2_sub_cumh_elem; i++){
  692. checksum[8] = checksum[8]+d_unique_d_in2_sub_cumh[i];
  693. }
  694. }
  695. //==================================================50
  696. // SYNCHRONIZE THREADS
  697. //==================================================50
  698. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  699. #endif
  700. //==================================================50
  701. // HORIZONTAL CUMULATIVE SUM
  702. //==================================================50
  703. // work
  704. ei_new = tx;
  705. while(ei_new < d_common.in2_sub_cumh_rows){
  706. // figure out row position
  707. pos_ori = ei_new;
  708. // variables
  709. sum = 0;
  710. // loop through all rows
  711. for(position = pos_ori; position < pos_ori+d_common.in2_sub_cumh_elem; position = position + d_common.in2_sub_cumh_rows){
  712. d_unique_d_in2_sub_cumh[position] = d_unique_d_in2_sub_cumh[position] + sum;
  713. sum = d_unique_d_in2_sub_cumh[position];
  714. }
  715. // go for second round
  716. ei_new = ei_new + NUMBER_THREADS;
  717. }
  718. //==================================================50
  719. // SYNCHRONIZE THREADS
  720. //==================================================50
  721. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  722. //==================================================50
  723. // CHECKSUM
  724. //==================================================50
  725. #ifdef TEST_CHECKSUM
  726. if(bx==0 && tx==0){
  727. for(i=0; i<d_common.in2_sub_cumh_elem; i++){
  728. checksum[9] = checksum[9]+d_unique_d_in2_sub_cumh[i];
  729. }
  730. }
  731. //==================================================50
  732. // SYNCHRONIZE THREADS
  733. //==================================================50
  734. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  735. #endif
  736. //==================================================50
  737. // SELECTION
  738. //==================================================50
  739. // work
  740. ei_new = tx;
  741. while(ei_new < d_common.in2_sub_cumh_sel_elem){
  742. // figure out row/col location in new matrix
  743. row = (ei_new+1) % d_common.in2_sub_cumh_sel_rows - 1; // (0-n) row
  744. col = (ei_new+1) / d_common.in2_sub_cumh_sel_rows + 1 - 1; // (0-n) column
  745. if((ei_new+1) % d_common.in2_sub_cumh_sel_rows == 0){
  746. row = d_common.in2_sub_cumh_sel_rows - 1;
  747. col = col - 1;
  748. }
  749. // figure out corresponding location in old matrix and copy values to new matrix
  750. ori_row = row + d_common.in2_sub_cumh_sel_rowlow - 1;
  751. ori_col = col + d_common.in2_sub_cumh_sel_collow - 1;
  752. d_unique_d_in2_sub_cumh_sel[ei_new] = d_unique_d_in2_sub_cumh[ori_col*d_common.in2_sub_cumh_rows+ori_row];
  753. // go for second round
  754. ei_new = ei_new + NUMBER_THREADS;
  755. }
  756. //==================================================50
  757. // SYNCHRONIZE THREADS
  758. //==================================================50
  759. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  760. //==================================================50
  761. // CHECKSUM
  762. //==================================================50
  763. #ifdef TEST_CHECKSUM
  764. if(bx==0 && tx==0){
  765. for(i=0; i<d_common.in2_sub_cumh_sel_elem; i++){
  766. checksum[10] = checksum[10]+d_unique_d_in2_sub_cumh_sel[i];
  767. }
  768. }
  769. //==================================================50
  770. // SYNCHRONIZE THREADS
  771. //==================================================50
  772. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  773. #endif
  774. //==================================================50
  775. // SELECTION 2
  776. //==================================================50
  777. // work
  778. ei_new = tx;
  779. while(ei_new < d_common.in2_sub2_elem){
  780. // figure out row/col location in new matrix
  781. row = (ei_new+1) % d_common.in2_sub2_rows - 1; // (0-n) row
  782. col = (ei_new+1) / d_common.in2_sub2_rows + 1 - 1; // (0-n) column
  783. if((ei_new+1) % d_common.in2_sub2_rows == 0){
  784. row = d_common.in2_sub2_rows - 1;
  785. col = col-1;
  786. }
  787. // figure out corresponding location in old matrix and copy values to new matrix
  788. ori_row = row + d_common.in2_sub_cumh_sel2_rowlow - 1;
  789. ori_col = col + d_common.in2_sub_cumh_sel2_collow - 1;
  790. d_unique_d_in2_sub2[ei_new] = d_unique_d_in2_sub_cumh[ori_col*d_common.in2_sub_cumh_rows+ori_row];
  791. // go for second round
  792. ei_new = ei_new + NUMBER_THREADS;
  793. }
  794. //==================================================50
  795. // SYNCHRONIZE THREADS
  796. //==================================================50
  797. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  798. //==================================================50
  799. // CHECKSUM
  800. //==================================================50
  801. #ifdef TEST_CHECKSUM
  802. if(bx==0 && tx==0){
  803. for(i=0; i<d_common.in2_sub2_elem; i++){
  804. checksum[11] = checksum[11]+d_unique_d_in2_sub2[i];
  805. }
  806. }
  807. //==================================================50
  808. // SYNCHRONIZE THREADS
  809. //==================================================50
  810. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  811. #endif
  812. //==================================================50
  813. // SUBTRACTION
  814. //==================================================50
  815. // work
  816. ei_new = tx;
  817. while(ei_new < d_common.in2_sub2_elem){
  818. // subtract
  819. d_unique_d_in2_sub2[ei_new] = d_unique_d_in2_sub_cumh_sel[ei_new] - d_unique_d_in2_sub2[ei_new];
  820. // go for second round
  821. ei_new = ei_new + NUMBER_THREADS;
  822. }
  823. //==================================================50
  824. // SYNCHRONIZE THREADS
  825. //==================================================50
  826. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  827. //==================================================50
  828. // CHECKSUM
  829. //==================================================50
  830. #ifdef TEST_CHECKSUM
  831. if(bx==0 && tx==0){
  832. for(i=0; i<d_common.in2_sub2_elem; i++){
  833. checksum[12] = checksum[12]+d_unique_d_in2_sub2[i];
  834. }
  835. }
  836. //==================================================50
  837. // SYNCHRONIZE THREADS
  838. //==================================================50
  839. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  840. #endif
  841. //==================================================50
  842. // End
  843. //==================================================50
  844. //====================================================================================================100
  845. // CUMULATIVE SUM 2
  846. //====================================================================================================100
  847. //==================================================50
  848. // MULTIPLICATION
  849. //==================================================50
  850. // work
  851. ei_new = tx;
  852. while(ei_new < d_common.in2_sqr_elem){
  853. temp = d_unique_d_in2[ei_new];
  854. d_unique_d_in2_sqr[ei_new] = temp * temp;
  855. // go for second round
  856. ei_new = ei_new + NUMBER_THREADS;
  857. }
  858. //==================================================50
  859. // SYNCHRONIZE THREADS
  860. //==================================================50
  861. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  862. //==================================================50
  863. // CHECKSUM
  864. //==================================================50
  865. #ifdef TEST_CHECKSUM
  866. if(bx==0 && tx==0){
  867. for(i=0; i<d_common.in2_sqr_elem; i++){
  868. checksum[13] = checksum[13]+d_unique_d_in2_sqr[i];
  869. }
  870. }
  871. //==================================================50
  872. // SYNCHRONIZE THREADS
  873. //==================================================50
  874. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  875. #endif
  876. //==================================================50
  877. // PAD ARRAY, VERTICAL CUMULATIVE SUM
  878. //==================================================50
  879. //==================================================50
  880. // PAD ARRAY
  881. //==================================================50
  882. // work
  883. ei_new = tx;
  884. while(ei_new < d_common.in2_pad_cumv_elem){
  885. // figure out row/col location in padded array
  886. row = (ei_new+1) % d_common.in2_pad_cumv_rows - 1; // (0-n) row
  887. col = (ei_new+1) / d_common.in2_pad_cumv_rows + 1 - 1; // (0-n) column
  888. if((ei_new+1) % d_common.in2_pad_cumv_rows == 0){
  889. row = d_common.in2_pad_cumv_rows - 1;
  890. col = col-1;
  891. }
  892. // execution
  893. if( row > (d_common.in2_pad_add_rows-1) && // do if has numbers in original array
  894. row < (d_common.in2_pad_add_rows+d_common.in2_sqr_rows) &&
  895. col > (d_common.in2_pad_add_cols-1) &&
  896. col < (d_common.in2_pad_add_cols+d_common.in2_sqr_cols)){
  897. ori_row = row - d_common.in2_pad_add_rows;
  898. ori_col = col - d_common.in2_pad_add_cols;
  899. d_unique_d_in2_pad_cumv[ei_new] = d_unique_d_in2_sqr[ori_col*d_common.in2_sqr_rows+ori_row];
  900. }
  901. else{ // do if otherwise
  902. d_unique_d_in2_pad_cumv[ei_new] = 0;
  903. }
  904. // go for second round
  905. ei_new = ei_new + NUMBER_THREADS;
  906. }
  907. //==================================================50
  908. // SYNCHRONIZE THREADS
  909. //==================================================50
  910. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  911. //==================================================50
  912. // CHECKSUM
  913. //==================================================50
  914. #ifdef TEST_CHECKSUM
  915. if(bx==0 && tx==0){
  916. for(i=0; i<d_common.in2_pad_cumv_elem; i++){
  917. checksum[14] = checksum[14]+d_unique_d_in2_pad_cumv[i];
  918. }
  919. }
  920. //==================================================50
  921. // SYNCHRONIZE THREADS
  922. //==================================================50
  923. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  924. #endif
  925. //==================================================50
  926. // VERTICAL CUMULATIVE SUM
  927. //==================================================50
  928. //work
  929. ei_new = tx;
  930. while(ei_new < d_common.in2_pad_cumv_cols){
  931. // figure out column position
  932. pos_ori = ei_new*d_common.in2_pad_cumv_rows;
  933. // variables
  934. sum = 0;
  935. // loop through all rows
  936. for(position = pos_ori; position < pos_ori+d_common.in2_pad_cumv_rows; position = position + 1){
  937. d_unique_d_in2_pad_cumv[position] = d_unique_d_in2_pad_cumv[position] + sum;
  938. sum = d_unique_d_in2_pad_cumv[position];
  939. }
  940. // go for second round
  941. ei_new = ei_new + NUMBER_THREADS;
  942. }
  943. //==================================================50
  944. // SYNCHRONIZE THREADS
  945. //==================================================50
  946. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  947. //==================================================50
  948. // CHECKSUM
  949. //==================================================50
  950. #ifdef TEST_CHECKSUM
  951. if(bx==0 && tx==0){
  952. for(i=0; i<d_common.in2_pad_cumv_elem; i++){
  953. checksum[15] = checksum[15]+d_unique_d_in2_pad_cumv[i];
  954. }
  955. }
  956. //==================================================50
  957. // SYNCHRONIZE THREADS
  958. //==================================================50
  959. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  960. #endif
  961. //==================================================50
  962. // SELECTION
  963. //==================================================50
  964. // work
  965. ei_new = tx;
  966. while(ei_new < d_common.in2_pad_cumv_sel_elem){
  967. // figure out row/col location in new matrix
  968. row = (ei_new+1) % d_common.in2_pad_cumv_sel_rows - 1; // (0-n) row
  969. col = (ei_new+1) / d_common.in2_pad_cumv_sel_rows + 1 - 1; // (0-n) column
  970. if((ei_new+1) % d_common.in2_pad_cumv_sel_rows == 0){
  971. row = d_common.in2_pad_cumv_sel_rows - 1;
  972. col = col-1;
  973. }
  974. // figure out corresponding location in old matrix and copy values to new matrix
  975. ori_row = row + d_common.in2_pad_cumv_sel_rowlow - 1;
  976. ori_col = col + d_common.in2_pad_cumv_sel_collow - 1;
  977. d_unique_d_in2_pad_cumv_sel[ei_new] = d_unique_d_in2_pad_cumv[ori_col*d_common.in2_pad_cumv_rows+ori_row];
  978. // go for second round
  979. ei_new = ei_new + NUMBER_THREADS;
  980. }
  981. //==================================================50
  982. // SYNCHRONIZE THREADS
  983. //==================================================50
  984. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  985. //==================================================50
  986. // CHECKSUM
  987. //==================================================50
  988. #ifdef TEST_CHECKSUM
  989. if(bx==0 && tx==0){
  990. for(i=0; i<d_common.in2_pad_cumv_sel_elem; i++){
  991. checksum[16] = checksum[16]+d_unique_d_in2_pad_cumv_sel[i];
  992. }
  993. }
  994. //==================================================50
  995. // SYNCHRONIZE THREADS
  996. //==================================================50
  997. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  998. #endif
  999. //==================================================50
  1000. // SELECTION 2
  1001. //==================================================50
  1002. // work
  1003. ei_new = tx;
  1004. while(ei_new < d_common.in2_sub_cumh_elem){
  1005. // figure out row/col location in new matrix
  1006. row = (ei_new+1) % d_common.in2_sub_cumh_rows - 1; // (0-n) row
  1007. col = (ei_new+1) / d_common.in2_sub_cumh_rows + 1 - 1; // (0-n) column
  1008. if((ei_new+1) % d_common.in2_sub_cumh_rows == 0){
  1009. row = d_common.in2_sub_cumh_rows - 1;
  1010. col = col-1;
  1011. }
  1012. // figure out corresponding location in old matrix and copy values to new matrix
  1013. ori_row = row + d_common.in2_pad_cumv_sel2_rowlow - 1;
  1014. ori_col = col + d_common.in2_pad_cumv_sel2_collow - 1;
  1015. d_unique_d_in2_sub_cumh[ei_new] = d_unique_d_in2_pad_cumv[ori_col*d_common.in2_pad_cumv_rows+ori_row];
  1016. // go for second round
  1017. ei_new = ei_new + NUMBER_THREADS;
  1018. }
  1019. //==================================================50
  1020. // SYNCHRONIZE THREADS
  1021. //==================================================50
  1022. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1023. //==================================================50
  1024. // CHECKSUM
  1025. //==================================================50
  1026. #ifdef TEST_CHECKSUM
  1027. if(bx==0 && tx==0){
  1028. for(i=0; i<d_common.in2_sub_cumh_elem; i++){
  1029. checksum[17] = checksum[17]+d_unique_d_in2_sub_cumh[i];
  1030. }
  1031. }
  1032. //==================================================50
  1033. // SYNCHRONIZE THREADS
  1034. //==================================================50
  1035. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1036. #endif
  1037. //==================================================50
  1038. // SUBTRACTION
  1039. //==================================================50
  1040. // work
  1041. ei_new = tx;
  1042. while(ei_new < d_common.in2_sub_cumh_elem){
  1043. // subtract
  1044. d_unique_d_in2_sub_cumh[ei_new] = d_unique_d_in2_pad_cumv_sel[ei_new] - d_unique_d_in2_sub_cumh[ei_new];
  1045. // go for second round
  1046. ei_new = ei_new + NUMBER_THREADS;
  1047. }
  1048. //==================================================50
  1049. // SYNCHRONIZE THREADS
  1050. //==================================================50
  1051. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1052. //==================================================50
  1053. // CHECKSUM
  1054. //==================================================50
  1055. #ifdef TEST_CHECKSUM
  1056. if(bx==0 && tx==0){
  1057. for(i=0; i<d_common.in2_sub_cumh_elem; i++){
  1058. checksum[18] = checksum[18]+d_unique_d_in2_sub_cumh[i];
  1059. }
  1060. }
  1061. //==================================================50
  1062. // SYNCHRONIZE THREADS
  1063. //==================================================50
  1064. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1065. #endif
  1066. //==================================================50
  1067. // HORIZONTAL CUMULATIVE SUM
  1068. //==================================================50
  1069. // work
  1070. ei_new = tx;
  1071. while(ei_new < d_common.in2_sub_cumh_rows){
  1072. // figure out row position
  1073. pos_ori = ei_new;
  1074. // variables
  1075. sum = 0;
  1076. // loop through all rows
  1077. for(position = pos_ori; position < pos_ori+d_common.in2_sub_cumh_elem; position = position + d_common.in2_sub_cumh_rows){
  1078. d_unique_d_in2_sub_cumh[position] = d_unique_d_in2_sub_cumh[position] + sum;
  1079. sum = d_unique_d_in2_sub_cumh[position];
  1080. }
  1081. // go for second round
  1082. ei_new = ei_new + NUMBER_THREADS;
  1083. }
  1084. //==================================================50
  1085. // SYNCHRONIZE THREADS
  1086. //==================================================50
  1087. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1088. //==================================================50
  1089. // CHECKSUM
  1090. //==================================================50
  1091. #ifdef TEST_CHECKSUM
  1092. if(bx==0 && tx==0){
  1093. for(i=0; i<d_common.in2_sub_cumh_rows; i++){
  1094. checksum[19] = checksum[19]+d_unique_d_in2_sub_cumh[i];
  1095. }
  1096. }
  1097. //==================================================50
  1098. // SYNCHRONIZE THREADS
  1099. //==================================================50
  1100. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1101. #endif
  1102. //==================================================50
  1103. // SELECTION
  1104. //==================================================50
  1105. // work
  1106. ei_new = tx;
  1107. while(ei_new < d_common.in2_sub_cumh_sel_elem){
  1108. // figure out row/col location in new matrix
  1109. row = (ei_new+1) % d_common.in2_sub_cumh_sel_rows - 1; // (0-n) row
  1110. col = (ei_new+1) / d_common.in2_sub_cumh_sel_rows + 1 - 1; // (0-n) column
  1111. if((ei_new+1) % d_common.in2_sub_cumh_sel_rows == 0){
  1112. row = d_common.in2_sub_cumh_sel_rows - 1;
  1113. col = col - 1;
  1114. }
  1115. // figure out corresponding location in old matrix and copy values to new matrix
  1116. ori_row = row + d_common.in2_sub_cumh_sel_rowlow - 1;
  1117. ori_col = col + d_common.in2_sub_cumh_sel_collow - 1;
  1118. d_unique_d_in2_sub_cumh_sel[ei_new] = d_unique_d_in2_sub_cumh[ori_col*d_common.in2_sub_cumh_rows+ori_row];
  1119. // go for second round
  1120. ei_new = ei_new + NUMBER_THREADS;
  1121. }
  1122. //==================================================50
  1123. // SYNCHRONIZE THREADS
  1124. //==================================================50
  1125. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1126. //==================================================50
  1127. // CHECKSUM
  1128. //==================================================50
  1129. #ifdef TEST_CHECKSUM
  1130. if(bx==0 && tx==0){
  1131. for(i=0; i<d_common.in2_sub_cumh_sel_elem; i++){
  1132. checksum[20] = checksum[20]+d_unique_d_in2_sub_cumh_sel[i];
  1133. }
  1134. }
  1135. //==================================================50
  1136. // SYNCHRONIZE THREADS
  1137. //==================================================50
  1138. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1139. #endif
  1140. //==================================================50
  1141. // SELECTION 2
  1142. //==================================================50
  1143. // work
  1144. ei_new = tx;
  1145. while(ei_new < d_common.in2_sub2_elem){
  1146. // figure out row/col location in new matrix
  1147. row = (ei_new+1) % d_common.in2_sub2_rows - 1; // (0-n) row
  1148. col = (ei_new+1) / d_common.in2_sub2_rows + 1 - 1; // (0-n) column
  1149. if((ei_new+1) % d_common.in2_sub2_rows == 0){
  1150. row = d_common.in2_sub2_rows - 1;
  1151. col = col-1;
  1152. }
  1153. // figure out corresponding location in old matrix and copy values to new matrix
  1154. ori_row = row + d_common.in2_sub_cumh_sel2_rowlow - 1;
  1155. ori_col = col + d_common.in2_sub_cumh_sel2_collow - 1;
  1156. d_unique_d_in2_sqr_sub2[ei_new] = d_unique_d_in2_sub_cumh[ori_col*d_common.in2_sub_cumh_rows+ori_row];
  1157. // go for second round
  1158. ei_new = ei_new + NUMBER_THREADS;
  1159. }
  1160. //==================================================50
  1161. // SYNCHRONIZE THREADS
  1162. //==================================================50
  1163. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1164. //==================================================50
  1165. // CHECKSUM
  1166. //==================================================50
  1167. #ifdef TEST_CHECKSUM
  1168. if(bx==0 && tx==0){
  1169. for(i=0; i<d_common.in2_sub2_elem; i++){
  1170. checksum[21] = checksum[21]+d_unique_d_in2_sqr_sub2[i];
  1171. }
  1172. }
  1173. //==================================================50
  1174. // SYNCHRONIZE THREADS
  1175. //==================================================50
  1176. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1177. #endif
  1178. //==================================================50
  1179. // SUBTRACTION
  1180. //==================================================50
  1181. // work
  1182. ei_new = tx;
  1183. while(ei_new < d_common.in2_sub2_elem){
  1184. // subtract
  1185. d_unique_d_in2_sqr_sub2[ei_new] = d_unique_d_in2_sub_cumh_sel[ei_new] - d_unique_d_in2_sqr_sub2[ei_new];
  1186. // go for second round
  1187. ei_new = ei_new + NUMBER_THREADS;
  1188. }
  1189. //==================================================50
  1190. // SYNCHRONIZE THREADS
  1191. //==================================================50
  1192. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1193. //==================================================50
  1194. // CHECKSUM
  1195. //==================================================50
  1196. #ifdef TEST_CHECKSUM
  1197. if(bx==0 && tx==0){
  1198. for(i=0; i<d_common.in2_sub2_elem; i++){
  1199. checksum[22] = checksum[22]+d_unique_d_in2_sqr_sub2[i];
  1200. }
  1201. }
  1202. //==================================================50
  1203. // SYNCHRONIZE THREADS
  1204. //==================================================50
  1205. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1206. #endif
  1207. //==================================================50
  1208. // End
  1209. //==================================================50
  1210. //====================================================================================================100
  1211. // FINAL
  1212. //====================================================================================================100
  1213. //==================================================50
  1214. // DENOMINATOR A SAVE RESULT IN CUMULATIVE SUM A2
  1215. //==================================================50
  1216. // work
  1217. ei_new = tx;
  1218. while(ei_new < d_common.in2_sub2_elem){
  1219. temp = d_unique_d_in2_sub2[ei_new];
  1220. temp2 = d_unique_d_in2_sqr_sub2[ei_new] - (temp * temp / d_common.in_elem);
  1221. if(temp2 < 0){
  1222. temp2 = 0;
  1223. }
  1224. d_unique_d_in2_sqr_sub2[ei_new] = sqrt(temp2);
  1225. // go for second round
  1226. ei_new = ei_new + NUMBER_THREADS;
  1227. }
  1228. //==================================================50
  1229. // SYNCHRONIZE THREADS
  1230. //==================================================50
  1231. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1232. //==================================================50
  1233. // CHECKSUM
  1234. //==================================================50
  1235. #ifdef TEST_CHECKSUM
  1236. if(bx==0 && tx==0){
  1237. for(i=0; i<d_common.in2_sub2_elem; i++){
  1238. checksum[23] = checksum[23]+d_unique_d_in2_sqr_sub2[i];
  1239. }
  1240. }
  1241. //==================================================50
  1242. // SYNCHRONIZE THREADS
  1243. //==================================================50
  1244. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1245. #endif
  1246. //==================================================50
  1247. // MULTIPLICATION
  1248. //==================================================50
  1249. // work
  1250. ei_new = tx;
  1251. while(ei_new < d_common.in_sqr_elem){
  1252. temp = d_in[ei_new];
  1253. d_unique_d_in_sqr[ei_new] = temp * temp;
  1254. // go for second round
  1255. ei_new = ei_new + NUMBER_THREADS;
  1256. }
  1257. //==================================================50
  1258. // SYNCHRONIZE THREADS
  1259. //==================================================50
  1260. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1261. //==================================================50
  1262. // CHECKSUM
  1263. //==================================================50
  1264. #ifdef TEST_CHECKSUM
  1265. if(bx==0 && tx==0){
  1266. for(i=0; i<d_common.in_sqr_elem; i++){
  1267. checksum[24] = checksum[24]+d_unique_d_in_sqr[i];
  1268. }
  1269. }
  1270. //==================================================50
  1271. // SYNCHRONIZE THREADS
  1272. //==================================================50
  1273. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1274. #endif
  1275. //==================================================50
  1276. // IN SUM
  1277. //==================================================50
  1278. // work
  1279. ei_new = tx;
  1280. while(ei_new < d_common.in_cols){
  1281. sum = 0;
  1282. for(i = 0; i < d_common.in_rows; i++){
  1283. sum = sum + d_in[ei_new*d_common.in_rows+i];
  1284. }
  1285. in_partial_sum[ei_new] = sum;
  1286. // go for second round
  1287. ei_new = ei_new + NUMBER_THREADS;
  1288. }
  1289. //==================================================50
  1290. // SYNCHRONIZE THREADS
  1291. //==================================================50
  1292. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1293. //==================================================50
  1294. // CHECKSUM
  1295. //==================================================50
  1296. #ifdef TEST_CHECKSUM
  1297. if(bx==0 && tx==0){
  1298. for(i=0; i<d_common.in_cols; i++){
  1299. checksum[25] = checksum[25]+in_partial_sum[i];
  1300. }
  1301. }
  1302. //==================================================50
  1303. // SYNCHRONIZE THREADS
  1304. //==================================================50
  1305. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1306. #endif
  1307. //==================================================50
  1308. // IN_SQR SUM
  1309. //==================================================50
  1310. ei_new = tx;
  1311. while(ei_new < d_common.in_sqr_rows){
  1312. sum = 0;
  1313. for(i = 0; i < d_common.in_sqr_cols; i++){
  1314. sum = sum + d_unique_d_in_sqr[ei_new+d_common.in_sqr_rows*i];
  1315. }
  1316. in_sqr_partial_sum[ei_new] = sum;
  1317. // go for second round
  1318. ei_new = ei_new + NUMBER_THREADS;
  1319. }
  1320. //==================================================50
  1321. // SYNCHRONIZE THREADS
  1322. //==================================================50
  1323. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1324. //==================================================50
  1325. // CHECKSUM
  1326. //==================================================50
  1327. #ifdef TEST_CHECKSUM
  1328. if(bx==0 && tx==0){
  1329. for(i=0; i<d_common.in_sqr_rows; i++){
  1330. checksum[26] = checksum[26]+in_sqr_partial_sum[i];
  1331. }
  1332. }
  1333. //==================================================50
  1334. // SYNCHRONIZE THREADS
  1335. //==================================================50
  1336. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1337. #endif
  1338. //==================================================50
  1339. // FINAL SUMMATION
  1340. //==================================================50
  1341. if(tx == 0){
  1342. in_final_sum[0] = 0;
  1343. for(i = 0; i<d_common.in_cols; i++){
  1344. // in_final_sum = in_final_sum + in_partial_sum[i];
  1345. in_final_sum[0] = in_final_sum[0] + in_partial_sum[i];
  1346. }
  1347. }else if(tx == 1){
  1348. in_sqr_final_sum[0] = 0;
  1349. for(i = 0; i<d_common.in_sqr_cols; i++){
  1350. // in_sqr_final_sum = in_sqr_final_sum + in_sqr_partial_sum[i];
  1351. in_sqr_final_sum[0] = in_sqr_final_sum[0] + in_sqr_partial_sum[i];
  1352. }
  1353. }
  1354. //==================================================50
  1355. // SYNCHRONIZE THREADS
  1356. //==================================================50
  1357. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1358. //==================================================50
  1359. // CHECKSUM
  1360. //==================================================50
  1361. #ifdef TEST_CHECKSUM
  1362. if(bx==0 && tx==0){
  1363. checksum[27] = checksum[27]+in_final_sum[0]+in_sqr_final_sum[0];
  1364. }
  1365. //==================================================50
  1366. // SYNCHRONIZE THREADS
  1367. //==================================================50
  1368. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1369. #endif
  1370. //==================================================50
  1371. // DENOMINATOR T
  1372. //==================================================50
  1373. if(tx == 0){
  1374. // mean = in_final_sum / d_common.in_elem; // gets mean (average) value of element in ROI
  1375. mean = in_final_sum[0] / d_common.in_elem; // gets mean (average) value of element in ROI
  1376. mean_sqr = mean * mean;
  1377. // variance = (in_sqr_final_sum / d_common.in_elem) - mean_sqr; // gets variance of ROI
  1378. variance = (in_sqr_final_sum[0] / d_common.in_elem) - mean_sqr; // gets variance of ROI
  1379. deviation = sqrt(variance); // gets standard deviation of ROI
  1380. // denomT = sqrt((float)(d_common.in_elem-1))*deviation;
  1381. denomT[0] = sqrt((float)(d_common.in_elem-1))*deviation;
  1382. }
  1383. //==================================================50
  1384. // SYNCHRONIZE THREADS
  1385. //==================================================50
  1386. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1387. //==================================================50
  1388. // CHECKSUM
  1389. //==================================================50
  1390. #ifdef TEST_CHECKSUM
  1391. if(bx==0 && tx==0){
  1392. checksum[28] = checksum[28]+denomT[i];
  1393. }
  1394. //==================================================50
  1395. // SYNCHRONIZE THREADS
  1396. //==================================================50
  1397. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1398. #endif
  1399. //==================================================50
  1400. // DENOMINATOR SAVE RESULT IN CUMULATIVE SUM A2
  1401. //==================================================50
  1402. // work
  1403. ei_new = tx;
  1404. while(ei_new < d_common.in2_sub2_elem){
  1405. // d_unique_d_in2_sqr_sub2[ei_new] = d_unique_d_in2_sqr_sub2[ei_new] * denomT;
  1406. d_unique_d_in2_sqr_sub2[ei_new] = d_unique_d_in2_sqr_sub2[ei_new] * denomT[0];
  1407. // go for second round
  1408. ei_new = ei_new + NUMBER_THREADS;
  1409. }
  1410. //==================================================50
  1411. // SYNCHRONIZE THREADS
  1412. //==================================================50
  1413. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1414. //==================================================50
  1415. // CHECKSUM
  1416. //==================================================50
  1417. #ifdef TEST_CHECKSUM
  1418. if(bx==0 && tx==0){
  1419. for(i=0; i<d_common.in2_sub2_elem; i++){
  1420. checksum[29] = checksum[29]+d_unique_d_in2_sqr_sub2[i];
  1421. }
  1422. }
  1423. //==================================================50
  1424. // SYNCHRONIZE THREADS
  1425. //==================================================50
  1426. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1427. #endif
  1428. //==================================================50
  1429. // NUMERATOR SAVE RESULT IN CONVOLUTION
  1430. //==================================================50
  1431. // work
  1432. ei_new = tx;
  1433. while(ei_new < d_common.conv_elem){
  1434. // d_unique_d_conv[ei_new] = d_unique_d_conv[ei_new] - d_unique_d_in2_sub2[ei_new] * in_final_sum / d_common.in_elem;
  1435. d_unique_d_conv[ei_new] = d_unique_d_conv[ei_new] - d_unique_d_in2_sub2[ei_new] * in_final_sum[0] / d_common.in_elem;
  1436. // go for second round
  1437. ei_new = ei_new + NUMBER_THREADS;
  1438. }
  1439. //==================================================50
  1440. // SYNCHRONIZE THREADS
  1441. //==================================================50
  1442. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1443. //==================================================50
  1444. // CHECKSUM
  1445. //==================================================50
  1446. #ifdef TEST_CHECKSUM
  1447. if(bx==0 && tx==0){
  1448. for(i=0; i<d_common.conv_elem; i++){
  1449. checksum[30] = checksum[30]+d_unique_d_conv[i];
  1450. }
  1451. }
  1452. //==================================================50
  1453. // SYNCHRONIZE THREADS
  1454. //==================================================50
  1455. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1456. #endif
  1457. //==================================================50
  1458. // CORRELATION SAVE RESULT IN CUMULATIVE SUM A2
  1459. //==================================================50
  1460. // work
  1461. ei_new = tx;
  1462. while(ei_new < d_common.in2_sub2_elem){
  1463. d_unique_d_in2_sqr_sub2[ei_new] = d_unique_d_conv[ei_new] / d_unique_d_in2_sqr_sub2[ei_new];
  1464. // go for second round
  1465. ei_new = ei_new + NUMBER_THREADS;
  1466. }
  1467. //==================================================50
  1468. // SYNCHRONIZE THREADS
  1469. //==================================================50
  1470. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1471. //==================================================50
  1472. // CHECKSUM
  1473. //==================================================50
  1474. #ifdef TEST_CHECKSUM
  1475. if(bx==0 && tx==0){
  1476. for(i=0; i<d_common.in2_sub2_elem; i++){
  1477. checksum[31] = checksum[31]+d_unique_d_in2_sqr_sub2[i];
  1478. }
  1479. }
  1480. //==================================================50
  1481. // SYNCHRONIZE THREADS
  1482. //==================================================50
  1483. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1484. #endif
  1485. //==================================================50
  1486. // End
  1487. //==================================================50
  1488. //====================================================================================================100
  1489. // TEMPLATE MASK CREATE
  1490. //====================================================================================================100
  1491. cent = d_common.sSize + d_common.tSize + 1;
  1492. if(d_frame_no == 0){
  1493. tMask_row = cent + d_unique_d_Row[d_unique_point_no] - d_unique_d_Row[d_unique_point_no] - 1;
  1494. tMask_col = cent + d_unique_d_Col[d_unique_point_no] - d_unique_d_Col[d_unique_point_no] - 1;
  1495. }
  1496. else{
  1497. pointer = d_unique_point_no*d_common.no_frames+d_frame_no-1;
  1498. tMask_row = cent + d_unique_d_tRowLoc[pointer] - d_unique_d_Row[d_unique_point_no] - 1;
  1499. tMask_col = cent + d_unique_d_tColLoc[pointer] - d_unique_d_Col[d_unique_point_no] - 1;
  1500. }
  1501. //work
  1502. ei_new = tx;
  1503. while(ei_new < d_common.tMask_elem){
  1504. location = tMask_col*d_common.tMask_rows + tMask_row;
  1505. if(ei_new==location){
  1506. d_unique_d_tMask[ei_new] = 1;
  1507. }
  1508. else{
  1509. d_unique_d_tMask[ei_new] = 0;
  1510. }
  1511. //go for second round
  1512. ei_new = ei_new + NUMBER_THREADS;
  1513. }
  1514. //==================================================50
  1515. // SYNCHRONIZE THREADS
  1516. //==================================================50
  1517. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1518. //==================================================50
  1519. // CHECKSUM
  1520. //==================================================50
  1521. #ifdef TEST_CHECKSUM
  1522. if(bx==0 && tx==0){
  1523. for(i=0; i<d_common.tMask_elem; i++){
  1524. checksum[32] = checksum[32]+d_unique_d_tMask[i];
  1525. }
  1526. }
  1527. //==================================================50
  1528. // SYNCHRONIZE THREADS
  1529. //==================================================50
  1530. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1531. #endif
  1532. //==================================================50
  1533. // End
  1534. //==================================================50
  1535. //====================================================================================================100
  1536. // MASK CONVOLUTION
  1537. //====================================================================================================100
  1538. // work
  1539. ei_new = tx;
  1540. while(ei_new < d_common.mask_conv_elem){
  1541. // figure out row/col location in array
  1542. ic = (ei_new+1) % d_common.mask_conv_rows; // (1-n)
  1543. jc = (ei_new+1) / d_common.mask_conv_rows + 1; // (1-n)
  1544. if((ei_new+1) % d_common.mask_conv_rows == 0){
  1545. ic = d_common.mask_conv_rows;
  1546. jc = jc-1;
  1547. }
  1548. //
  1549. j = jc + d_common.mask_conv_joffset;
  1550. jp1 = j + 1;
  1551. if(d_common.mask_cols < jp1){
  1552. ja1 = jp1 - d_common.mask_cols;
  1553. }
  1554. else{
  1555. ja1 = 1;
  1556. }
  1557. if(d_common.tMask_cols < j){
  1558. ja2 = d_common.tMask_cols;
  1559. }
  1560. else{
  1561. ja2 = j;
  1562. }
  1563. i = ic + d_common.mask_conv_ioffset;
  1564. ip1 = i + 1;
  1565. if(d_common.mask_rows < ip1){
  1566. ia1 = ip1 - d_common.mask_rows;
  1567. }
  1568. else{
  1569. ia1 = 1;
  1570. }
  1571. if(d_common.tMask_rows < i){
  1572. ia2 = d_common.tMask_rows;
  1573. }
  1574. else{
  1575. ia2 = i;
  1576. }
  1577. s = 0;
  1578. for(ja=ja1; ja<=ja2; ja++){
  1579. jb = jp1 - ja;
  1580. for(ia=ia1; ia<=ia2; ia++){
  1581. ib = ip1 - ia;
  1582. s = s + d_unique_d_tMask[d_common.tMask_rows*(ja-1)+ia-1] * 1;
  1583. }
  1584. }
  1585. // //d_unique_d_mask_conv[d_common.mask_conv_rows*(jc-1)+ic-1] = s;
  1586. d_unique_d_mask_conv[ei_new] = d_unique_d_in2_sqr_sub2[ei_new] * s;
  1587. // go for second round
  1588. ei_new = ei_new + NUMBER_THREADS;
  1589. }
  1590. //==================================================50
  1591. // SYNCHRONIZE THREADS
  1592. //==================================================50
  1593. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1594. //==================================================50
  1595. // CHECKSUM
  1596. //==================================================50
  1597. #ifdef TEST_CHECKSUM
  1598. if(bx==0 && tx==0){
  1599. for(i=0; i<d_common.mask_conv_elem; i++){
  1600. checksum[33] = checksum[33]+d_unique_d_mask_conv[i];
  1601. }
  1602. }
  1603. //==================================================50
  1604. // SYNCHRONIZE THREADS
  1605. //==================================================50
  1606. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1607. #endif
  1608. //==================================================50
  1609. // End
  1610. //==================================================50
  1611. //====================================================================================================100
  1612. // MAXIMUM VALUE
  1613. //====================================================================================================100
  1614. //==================================================50
  1615. // INITIAL SEARCH
  1616. //==================================================50
  1617. ei_new = tx;
  1618. while(ei_new < d_common.mask_conv_rows){
  1619. for(i=0; i<d_common.mask_conv_cols; i++){
  1620. largest_coordinate_current = ei_new*d_common.mask_conv_rows+i;
  1621. largest_value_current = fabs(d_unique_d_mask_conv[largest_coordinate_current]);
  1622. if(largest_value_current > largest_value){
  1623. largest_coordinate = largest_coordinate_current;
  1624. largest_value = largest_value_current;
  1625. }
  1626. }
  1627. par_max_coo[ei_new] = largest_coordinate;
  1628. par_max_val[ei_new] = largest_value;
  1629. // go for second round
  1630. ei_new = ei_new + NUMBER_THREADS;
  1631. }
  1632. //==================================================50
  1633. // SYNCHRONIZE THREADS
  1634. //==================================================50
  1635. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1636. //==================================================50
  1637. // CHECKSUM
  1638. //==================================================50
  1639. #ifdef TEST_CHECKSUM
  1640. if(bx==0 && tx==0){
  1641. for(i=0; i<d_common.mask_conv_rows; i++){
  1642. checksum[34] = checksum[34]+par_max_coo[i]+par_max_val[i];
  1643. }
  1644. }
  1645. //==================================================50
  1646. // SYNCHRONIZE THREADS
  1647. //==================================================50
  1648. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1649. #endif
  1650. //==================================================50
  1651. // FINAL SEARCH
  1652. //==================================================50
  1653. if(tx == 0){
  1654. for(i = 0; i < d_common.mask_conv_rows; i++){
  1655. if(par_max_val[i] > fin_max_val){
  1656. fin_max_val = par_max_val[i];
  1657. fin_max_coo = par_max_coo[i];
  1658. }
  1659. }
  1660. // convert coordinate to row/col form
  1661. largest_row = (fin_max_coo+1) % d_common.mask_conv_rows - 1; // (0-n) row
  1662. largest_col = (fin_max_coo+1) / d_common.mask_conv_rows; // (0-n) column
  1663. if((fin_max_coo+1) % d_common.mask_conv_rows == 0){
  1664. largest_row = d_common.mask_conv_rows - 1;
  1665. largest_col = largest_col - 1;
  1666. }
  1667. // calculate offset
  1668. largest_row = largest_row + 1; // compensate to match MATLAB format (1-n)
  1669. largest_col = largest_col + 1; // compensate to match MATLAB format (1-n)
  1670. offset_row = largest_row - d_common.in_rows - (d_common.sSize - d_common.tSize);
  1671. offset_col = largest_col - d_common.in_cols - (d_common.sSize - d_common.tSize);
  1672. pointer = d_unique_point_no*d_common.no_frames+d_frame_no;
  1673. d_unique_d_tRowLoc[pointer] = d_unique_d_Row[d_unique_point_no] + offset_row;
  1674. d_unique_d_tColLoc[pointer] = d_unique_d_Col[d_unique_point_no] + offset_col;
  1675. }
  1676. //==================================================50
  1677. // SYNCHRONIZE THREADS
  1678. //==================================================50
  1679. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1680. //==================================================50
  1681. // CHECKSUM
  1682. //==================================================50
  1683. #ifdef TEST_CHECKSUM
  1684. if(bx==0 && tx==0){
  1685. checksum[35] = checksum[35]+d_unique_d_tRowLoc[pointer]+d_unique_d_tColLoc[pointer];
  1686. }
  1687. //==================================================50
  1688. // SYNCHRONIZE THREADS
  1689. //==================================================50
  1690. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1691. #endif
  1692. //==================================================50
  1693. // End
  1694. //==================================================50
  1695. //====================================================================================================100
  1696. // End
  1697. //====================================================================================================100
  1698. }
  1699. //======================================================================================================================================================150
  1700. // PERIODIC COORDINATE AND TEMPLATE UPDATE
  1701. //======================================================================================================================================================150
  1702. if(d_frame_no != 0 && (d_frame_no)%10 == 0){
  1703. //====================================================================================================100
  1704. // initialize cross-frame variables
  1705. //====================================================================================================100
  1706. #ifdef INIT
  1707. // only the first thread initializes
  1708. if(tx==0){
  1709. // this block
  1710. for(i=0; i<d_common.in_elem; i++){
  1711. d_in[i] = 0;
  1712. }
  1713. }
  1714. #endif
  1715. //====================================================================================================100
  1716. // SYNCHRONIZE THREADS
  1717. //====================================================================================================100
  1718. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1719. //====================================================================================================100
  1720. // if the last frame in the bath, update template
  1721. //====================================================================================================100
  1722. // update coordinate
  1723. loc_pointer = d_unique_point_no*d_common.no_frames+d_frame_no;
  1724. d_unique_d_Row[d_unique_point_no] = d_unique_d_tRowLoc[loc_pointer];
  1725. d_unique_d_Col[d_unique_point_no] = d_unique_d_tColLoc[loc_pointer];
  1726. // work
  1727. ei_new = tx;
  1728. while(ei_new < d_common.in_elem){
  1729. // figure out row/col location in new matrix
  1730. row = (ei_new+1) % d_common.in_rows - 1; // (0-n) row
  1731. col = (ei_new+1) / d_common.in_rows + 1 - 1; // (0-n) column
  1732. if((ei_new+1) % d_common.in_rows == 0){
  1733. row = d_common.in_rows - 1;
  1734. col = col-1;
  1735. }
  1736. // figure out row/col location in corresponding new template area in image and give to every thread (get top left corner and progress down and right)
  1737. ori_row = d_unique_d_Row[d_unique_point_no] - 25 + row - 1;
  1738. ori_col = d_unique_d_Col[d_unique_point_no] - 25 + col - 1;
  1739. ori_pointer = ori_col*d_common.frame_rows+ori_row;
  1740. // update template
  1741. d_in[ei_new] = d_common.alpha*d_in[ei_new] + (1-d_common.alpha)*d_common_change_d_frame[ori_pointer];
  1742. // go for second round
  1743. ei_new = ei_new + NUMBER_THREADS;
  1744. }
  1745. //==================================================50
  1746. // SYNCHRONIZE THREADS
  1747. //==================================================50
  1748. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1749. //==================================================50
  1750. // CHECKSUM
  1751. //==================================================50
  1752. #ifdef TEST_CHECKSUM
  1753. if(bx==0 && tx==0){
  1754. for(i=0; i<d_common.in_elem; i++){
  1755. checksum[36] = checksum[36]+d_in[i];
  1756. }
  1757. }
  1758. //==================================================50
  1759. // SYNCHRONIZE THREADS
  1760. //==================================================50
  1761. barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
  1762. #endif
  1763. //==================================================50
  1764. // End
  1765. //==================================================50
  1766. //====================================================================================================100
  1767. // End
  1768. //====================================================================================================100
  1769. }
  1770. //======================================================================================================================================================150
  1771. // End
  1772. //======================================================================================================================================================150
  1773. }
  1774. //========================================================================================================================================================================================================200
  1775. // END
  1776. //========================================================================================================================================================================================================200