//========================================================================================================================================================================================================200 // DEFINE / INCLUDE //========================================================================================================================================================================================================200 //======================================================================================================================================================150 // MAIN FUNCTION HEADER //======================================================================================================================================================150 #include "./main.h" //======================================================================================================================================================150 // End //======================================================================================================================================================150 //========================================================================================================================================================================================================200 // KERNEL //========================================================================================================================================================================================================200 __kernel void kernel_gpu_opencl( // structures params_common d_common, // 0 // common_change __global fp* d_frame, // 1 INPUT int d_frame_no, // 2 INPUT // common __global int* d_endoRow, // 3 INPUT __global int* d_endoCol, // 4 INPUT __global int* d_tEndoRowLoc, // 5 OUTPUT common.endoPoints * common.no_frames __global int* d_tEndoColLoc, // 6 OUTPUT common.endoPoints * common.no_frames __global int* d_epiRow, // 7 INPUT __global int* d_epiCol, // 8 INPUT __global int* d_tEpiRowLoc, // 9 OUTPUT common.epiPoints * common.no_frames __global int* d_tEpiColLoc, // 10 OUTPUT common.epiPoints * common.no_frames // common_unique __global fp* d_endoT, // 11 OUTPUT common.in_elem * common.endoPoints __global fp* d_epiT, // 12 OUTPUT common.in_elem * common.epiPoints __global fp* d_in2_all, // 13 OUTPUT common.in2_elem * common.allPoints __global fp* d_conv_all, // 14 OUTPUT common.conv_elem * common.allPoints __global fp* d_in2_pad_cumv_all, // 15 OUTPUT common.in2_pad_cumv_elem * common.allPoints __global fp* d_in2_pad_cumv_sel_all, // 16 OUTPUT common.in2_pad_cumv_sel_elem * common.allPoints __global fp* d_in2_sub_cumh_all, // 17 OUTPUT common.in2_sub_cumh_elem * common.allPoints __global fp* d_in2_sub_cumh_sel_all, // 18 OUTPUT common.in2_sub_cumh_sel_elem * common.allPoints __global fp* d_in2_sub2_all, // 19 OUTPUT common.in2_sub2_elem * common.allPoints __global fp* d_in2_sqr_all, // 20 OUTPUT common.in2_elem * common.allPoints __global fp* d_in2_sqr_sub2_all, // 21 OUTPUT common.in2_sub2_elem * common.allPoints __global fp* d_in_sqr_all, // 22 OUTPUT common.in_elem * common.allPoints __global fp* d_tMask_all, // 23 OUTPUT common.tMask_elem * common.allPoints __global fp* d_mask_conv_all, // 24 OUTPUT common.mask_conv_elem * common.allPoints // // local // __local fp* d_in_mod_temp, // 25 OUTPUT common.in_elem // __local fp* in_partial_sum, // 26 OUTPUT common.in_cols // __local fp* in_sqr_partial_sum, // 27 OUTPUT common.in_sqr_rows // __local fp* par_max_val, // 28 OUTPUT common.mask_conv_rows // __local int* par_max_coo) // 29 OUTPUT common.mask_conv_rows // local __global fp* d_in_mod_temp_all, // 25 OUTPUT common.in_elem * common.allPoints __global fp* in_partial_sum_all, // 26 OUTPUT common.in_cols * common.allPoints __global fp* in_sqr_partial_sum_all, // 27 OUTPUT common.in_sqr_rows * common.allPoints __global fp* par_max_val_all, // 28 OUTPUT common.mask_conv_rows * common.allPoints __global int* par_max_coo_all, // 29 OUTPUT common.mask_conv_rows * common.allPoints __global fp* in_final_sum_all, // 30 OUTPUT common.allPoints __global fp* in_sqr_final_sum_all, // 31 OUTPUT common.allPoints __global fp* denomT_all, // 32 OUTPUT common.allPoints __global fp* checksum) // 33 OUTPUT 100 { //======================================================================================================================================================150 // COMMON VARIABLES //======================================================================================================================================================150 // __global fp* d_in; int rot_row; int rot_col; int in2_rowlow; int in2_collow; int ic; int jc; int jp1; int ja1, ja2; int ip1; int ia1, ia2; int ja, jb; int ia, ib; fp s; int i; int j; int row; int col; int ori_row; int ori_col; int position; fp sum; int pos_ori; fp temp; fp temp2; int location; int cent; int tMask_row; int tMask_col; fp largest_value_current = 0; fp largest_value = 0; int largest_coordinate_current = 0; int largest_coordinate = 0; fp fin_max_val = 0; int fin_max_coo = 0; int largest_row; int largest_col; int offset_row; int offset_col; fp mean; fp mean_sqr; fp variance; fp deviation; int pointer; int ori_pointer; int loc_pointer; // __local fp in_final_sum; // __local fp in_sqr_final_sum; // __local fp denomT; //======================================================================================================================================================150 // BLOCK/THREAD IDs //======================================================================================================================================================150 int bx = get_group_id(0); // get current horizontal block index (0-n) int tx = get_local_id(0); // get current horizontal thread index (0-n) int ei_new; //======================================================================================================================================================150 // UNIQUE STRUCTURE RECONSTRUCTED HERE //======================================================================================================================================================150 // common __global fp* d_common_change_d_frame = &d_frame[0]; // offsets for either endo or epi points (separate arrays for endo and epi points) int d_unique_point_no; __global int* d_unique_d_Row; __global int* d_unique_d_Col; __global int* d_unique_d_tRowLoc; __global int* d_unique_d_tColLoc; __global fp* d_in; if(bx < d_common.endoPoints){ d_unique_point_no = bx; // endo point number 0-??? d_unique_d_Row = d_endoRow; // initial endo row coordinates d_unique_d_Col = d_endoCol; // initial endo col coordinates d_unique_d_tRowLoc = d_tEndoRowLoc; // all endo row coordinates d_unique_d_tColLoc = d_tEndoColLoc; // all endo col coordinates d_in = &d_endoT[d_unique_point_no * d_common.in_elem]; // endo templates } else{ d_unique_point_no = bx-d_common.endoPoints; // epi point number 0-??? d_unique_d_Row = d_epiRow; // initial epi row coordinates d_unique_d_Col = d_epiCol; // initial epi col coordinates d_unique_d_tRowLoc = d_tEpiRowLoc; // all epi row coordinates d_unique_d_tColLoc = d_tEpiColLoc; // all epi col coordinates d_in = &d_epiT[d_unique_point_no * d_common.in_elem]; // epi templates } // offsets for all points (one array for all points) __global fp* d_unique_d_in2 = &d_in2_all[bx*d_common.in2_elem]; __global fp* d_unique_d_conv = &d_conv_all[bx*d_common.conv_elem]; __global fp* d_unique_d_in2_pad_cumv = &d_in2_pad_cumv_all[bx*d_common.in2_pad_cumv_elem]; __global fp* d_unique_d_in2_pad_cumv_sel = &d_in2_pad_cumv_sel_all[bx*d_common.in2_pad_cumv_sel_elem]; __global fp* d_unique_d_in2_sub_cumh = &d_in2_sub_cumh_all[bx*d_common.in2_sub_cumh_elem]; __global fp* d_unique_d_in2_sub_cumh_sel = &d_in2_sub_cumh_sel_all[bx*d_common.in2_sub_cumh_sel_elem]; __global fp* d_unique_d_in2_sub2 = &d_in2_sub2_all[bx*d_common.in2_sub2_elem]; __global fp* d_unique_d_in2_sqr = &d_in2_sqr_all[bx*d_common.in2_sqr_elem]; __global fp* d_unique_d_in2_sqr_sub2 = &d_in2_sqr_sub2_all[bx*d_common.in2_sqr_sub2_elem]; __global fp* d_unique_d_in_sqr = &d_in_sqr_all[bx*d_common.in_sqr_elem]; __global fp* d_unique_d_tMask = &d_tMask_all[bx*d_common.tMask_elem]; __global fp* d_unique_d_mask_conv = &d_mask_conv_all[bx*d_common.mask_conv_elem]; // used to be local __global fp* d_in_mod_temp = &d_in_mod_temp_all[bx*d_common.in_elem]; __global fp* in_partial_sum = &in_partial_sum_all[bx*d_common.in_cols]; __global fp* in_sqr_partial_sum = &in_sqr_partial_sum_all[bx*d_common.in_sqr_rows]; __global fp* par_max_val = &par_max_val_all[bx*d_common.mask_conv_rows]; __global int* par_max_coo = &par_max_coo_all[bx*d_common.mask_conv_rows]; __global fp* in_final_sum = &in_final_sum_all[bx]; __global fp* in_sqr_final_sum = &in_sqr_final_sum_all[bx]; __global fp* denomT = &denomT_all[bx]; //======================================================================================================================================================150 // END //======================================================================================================================================================150 //======================================================================================================================================================150 // Initialize checksum //======================================================================================================================================================150 #ifdef TEST_CHECKSUM if(bx==0 && tx==0){ for(i=0; i (2601*51 // printf("frame_no IS %d\n", d_common_change[0].frame_no); // d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_point_no]; // d_in_mod_temp[ei_new] = 1; // kot = d_in[rot_col*d_common.in_rows+rot_row]; // d_in_mod_temp[ei_new] = kot; // d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_in_pointer+rot_col*d_common.in_rows+rot_row]; // 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]; //d_in_mod_temp[ei_new] = d_unique_d_T[d_unique_point_no]; // d_unique_d_T[d_unique_in_pointer+rot_col*d_common.in_rows+rot_row] = 1; // d_unique_d_T[d_unique_in_pointer] = 1; // d_endoT[d_unique_in_pointer] = 1; // d_in_mod_temp[ei_new] = 1; // go for second round ei_new = ei_new + NUMBER_THREADS; } //==================================================50 // SYNCHRONIZE THREADS //==================================================50 barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); //==================================================50 // CHECKSUM //==================================================50 #ifdef TEST_CHECKSUM if(bx==0 && tx==0){ for(i=0; i (d_common.in2_pad_add_rows-1) && // do if has numbers in original array row < (d_common.in2_pad_add_rows+d_common.in2_rows) && col > (d_common.in2_pad_add_cols-1) && col < (d_common.in2_pad_add_cols+d_common.in2_cols)){ ori_row = row - d_common.in2_pad_add_rows; ori_col = col - d_common.in2_pad_add_cols; d_unique_d_in2_pad_cumv[ei_new] = d_unique_d_in2[ori_col*d_common.in2_rows+ori_row]; } else{ // do if otherwise d_unique_d_in2_pad_cumv[ei_new] = 0; } // go for second round ei_new = ei_new + NUMBER_THREADS; } //==================================================50 // SYNCHRONIZE THREADS //==================================================50 barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); //==================================================50 // CHECKSUM //==================================================50 #ifdef TEST_CHECKSUM if(bx==0 && tx==0){ for(i=0; i (d_common.in2_pad_add_rows-1) && // do if has numbers in original array row < (d_common.in2_pad_add_rows+d_common.in2_sqr_rows) && col > (d_common.in2_pad_add_cols-1) && col < (d_common.in2_pad_add_cols+d_common.in2_sqr_cols)){ ori_row = row - d_common.in2_pad_add_rows; ori_col = col - d_common.in2_pad_add_cols; d_unique_d_in2_pad_cumv[ei_new] = d_unique_d_in2_sqr[ori_col*d_common.in2_sqr_rows+ori_row]; } else{ // do if otherwise d_unique_d_in2_pad_cumv[ei_new] = 0; } // go for second round ei_new = ei_new + NUMBER_THREADS; } //==================================================50 // SYNCHRONIZE THREADS //==================================================50 barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); //==================================================50 // CHECKSUM //==================================================50 #ifdef TEST_CHECKSUM if(bx==0 && tx==0){ for(i=0; i largest_value){ largest_coordinate = largest_coordinate_current; largest_value = largest_value_current; } } par_max_coo[ei_new] = largest_coordinate; par_max_val[ei_new] = largest_value; // go for second round ei_new = ei_new + NUMBER_THREADS; } //==================================================50 // SYNCHRONIZE THREADS //==================================================50 barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); //==================================================50 // CHECKSUM //==================================================50 #ifdef TEST_CHECKSUM if(bx==0 && tx==0){ for(i=0; i fin_max_val){ fin_max_val = par_max_val[i]; fin_max_coo = par_max_coo[i]; } } // convert coordinate to row/col form largest_row = (fin_max_coo+1) % d_common.mask_conv_rows - 1; // (0-n) row largest_col = (fin_max_coo+1) / d_common.mask_conv_rows; // (0-n) column if((fin_max_coo+1) % d_common.mask_conv_rows == 0){ largest_row = d_common.mask_conv_rows - 1; largest_col = largest_col - 1; } // calculate offset largest_row = largest_row + 1; // compensate to match MATLAB format (1-n) largest_col = largest_col + 1; // compensate to match MATLAB format (1-n) offset_row = largest_row - d_common.in_rows - (d_common.sSize - d_common.tSize); offset_col = largest_col - d_common.in_cols - (d_common.sSize - d_common.tSize); pointer = d_unique_point_no*d_common.no_frames+d_frame_no; d_unique_d_tRowLoc[pointer] = d_unique_d_Row[d_unique_point_no] + offset_row; d_unique_d_tColLoc[pointer] = d_unique_d_Col[d_unique_point_no] + offset_col; } //==================================================50 // SYNCHRONIZE THREADS //==================================================50 barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); //==================================================50 // CHECKSUM //==================================================50 #ifdef TEST_CHECKSUM if(bx==0 && tx==0){ checksum[35] = checksum[35]+d_unique_d_tRowLoc[pointer]+d_unique_d_tColLoc[pointer]; } //==================================================50 // SYNCHRONIZE THREADS //==================================================50 barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); #endif //==================================================50 // End //==================================================50 //====================================================================================================100 // End //====================================================================================================100 } //======================================================================================================================================================150 // PERIODIC COORDINATE AND TEMPLATE UPDATE //======================================================================================================================================================150 if(d_frame_no != 0 && (d_frame_no)%10 == 0){ //====================================================================================================100 // initialize cross-frame variables //====================================================================================================100 #ifdef INIT // only the first thread initializes if(tx==0){ // this block for(i=0; i