ELF>f@@features2dorbfastbrute_force_matchakaze#define LAYERINFO_SIZE 1 #define LAYERINFO_OFS 0 #define KEYPOINT_SIZE 3 #define ORIENTED_KEYPOINT_SIZE 4 #define KEYPOINT_X 0 #define KEYPOINT_Y 1 #define KEYPOINT_Z 2 #define KEYPOINT_ANGLE 3 #ifdef ORB_RESPONSES __kernel void ORB_HarrisResponses(__global const uchar* imgbuf, int imgstep, int imgoffset0, __global const int* layerinfo, __global const int* keypoints, __global float* responses, int nkeypoints ) { int idx = get_global_id(0); if( idx < nkeypoints ) { __global const int* kpt = keypoints + idx*KEYPOINT_SIZE; __global const int* layer = layerinfo + kpt[KEYPOINT_Z]*LAYERINFO_SIZE; __global const uchar* img = imgbuf + imgoffset0 + layer[LAYERINFO_OFS] + (kpt[KEYPOINT_Y] - blockSize/2)*imgstep + (kpt[KEYPOINT_X] - blockSize/2); int i, j; int a = 0, b = 0, c = 0; for( i = 0; i < blockSize; i++, img += imgstep-blockSize ) { for( j = 0; j < blockSize; j++, img++ ) { int Ix = (img[1] - img[-1])*2 + img[-imgstep+1] - img[-imgstep-1] + img[imgstep+1] - img[imgstep-1]; int Iy = (img[imgstep] - img[-imgstep])*2 + img[imgstep-1] - img[-imgstep-1] + img[imgstep+1] - img[-imgstep+1]; a += Ix*Ix; b += Iy*Iy; c += Ix*Iy; } } responses[idx] = ((float)a * b - (float)c * c - HARRIS_K * (float)(a + b) * (a + b))*scale_sq_sq; } } #endif #ifdef ORB_ANGLES #define _DBL_EPSILON 2.2204460492503131e-16f #define atan2_p1 (0.9997878412794807f*57.29577951308232f) #define atan2_p3 (-0.3258083974640975f*57.29577951308232f) #define atan2_p5 (0.1555786518463281f*57.29577951308232f) #define atan2_p7 (-0.04432655554792128f*57.29577951308232f) inline float fastAtan2( float y, float x ) { float ax = fabs(x), ay = fabs(y); float a, c, c2; if( ax >= ay ) { c = ay/(ax + _DBL_EPSILON); c2 = c*c; a = (((atan2_p7*c2 + atan2_p5)*c2 + atan2_p3)*c2 + atan2_p1)*c; } else { c = ax/(ay + _DBL_EPSILON); c2 = c*c; a = 90.f - (((atan2_p7*c2 + atan2_p5)*c2 + atan2_p3)*c2 + atan2_p1)*c; } if( x < 0 ) a = 180.f - a; if( y < 0 ) a = 360.f - a; return a; } __kernel void ORB_ICAngle(__global const uchar* imgbuf, int imgstep, int imgoffset0, __global const int* layerinfo, __global const int* keypoints, __global float* responses, const __global int* u_max, int nkeypoints, int half_k ) { int idx = get_global_id(0); if( idx < nkeypoints ) { __global const int* kpt = keypoints + idx*KEYPOINT_SIZE; __global const int* layer = layerinfo + kpt[KEYPOINT_Z]*LAYERINFO_SIZE; __global const uchar* center = imgbuf + imgoffset0 + layer[LAYERINFO_OFS] + kpt[KEYPOINT_Y]*imgstep + kpt[KEYPOINT_X]; int u, v, m_01 = 0, m_10 = 0; for( u = -half_k; u <= half_k; u++ ) m_10 += u * center[u]; for( v = 1; v <= half_k; v++ ) { int v_sum = 0; int d = u_max[v]; for( u = -d; u <= d; u++ ) { int val_plus = center[u + v*imgstep], val_minus = center[u - v*imgstep]; v_sum += (val_plus - val_minus); m_10 += u * (val_plus + val_minus); } m_01 += v * v_sum; } responses[idx] = fastAtan2((float)m_01, (float)m_10); } } #endif #ifdef ORB_DESCRIPTORS __kernel void ORB_computeDescriptor(__global const uchar* imgbuf, int imgstep, int imgoffset0, __global const int* layerinfo, __global const int* keypoints, __global uchar* _desc, const __global int* pattern, int nkeypoints, int dsize ) { int idx = get_global_id(0); if( idx < nkeypoints ) { int i; __global const int* kpt = keypoints + idx*ORIENTED_KEYPOINT_SIZE; __global const int* layer = layerinfo + kpt[KEYPOINT_Z]*LAYERINFO_SIZE; __global const uchar* center = imgbuf + imgoffset0 + layer[LAYERINFO_OFS] + kpt[KEYPOINT_Y]*imgstep + kpt[KEYPOINT_X]; float angle = as_float(kpt[KEYPOINT_ANGLE]); angle *= 0.01745329251994329547f; float cosa; float sina = sincos(angle, &cosa); __global uchar* desc = _desc + idx*dsize; #define GET_VALUE(idx) \ center[mad24(convert_int_rte(pattern[(idx)*2] * sina + pattern[(idx)*2+1] * cosa), imgstep, \ convert_int_rte(pattern[(idx)*2] * cosa - pattern[(idx)*2+1] * sina))] for( i = 0; i < dsize; i++ ) { int val; #if WTA_K == 2 int t0, t1; t0 = GET_VALUE(0); t1 = GET_VALUE(1); val = t0 < t1; t0 = GET_VALUE(2); t1 = GET_VALUE(3); val |= (t0 < t1) << 1; t0 = GET_VALUE(4); t1 = GET_VALUE(5); val |= (t0 < t1) << 2; t0 = GET_VALUE(6); t1 = GET_VALUE(7); val |= (t0 < t1) << 3; t0 = GET_VALUE(8); t1 = GET_VALUE(9); val |= (t0 < t1) << 4; t0 = GET_VALUE(10); t1 = GET_VALUE(11); val |= (t0 < t1) << 5; t0 = GET_VALUE(12); t1 = GET_VALUE(13); val |= (t0 < t1) << 6; t0 = GET_VALUE(14); t1 = GET_VALUE(15); val |= (t0 < t1) << 7; pattern += 16*2; #elif WTA_K == 3 int t0, t1, t2; t0 = GET_VALUE(0); t1 = GET_VALUE(1); t2 = GET_VALUE(2); val = t2 > t1 ? (t2 > t0 ? 2 : 0) : (t1 > t0); t0 = GET_VALUE(3); t1 = GET_VALUE(4); t2 = GET_VALUE(5); val |= (t2 > t1 ? (t2 > t0 ? 2 : 0) : (t1 > t0)) << 2; t0 = GET_VALUE(6); t1 = GET_VALUE(7); t2 = GET_VALUE(8); val |= (t2 > t1 ? (t2 > t0 ? 2 : 0) : (t1 > t0)) << 4; t0 = GET_VALUE(9); t1 = GET_VALUE(10); t2 = GET_VALUE(11); val |= (t2 > t1 ? (t2 > t0 ? 2 : 0) : (t1 > t0)) << 6; pattern += 12*2; #elif WTA_K == 4 int t0, t1, t2, t3, k; int a, b; t0 = GET_VALUE(0); t1 = GET_VALUE(1); t2 = GET_VALUE(2); t3 = GET_VALUE(3); a = 0, b = 2; if( t1 > t0 ) t0 = t1, a = 1; if( t3 > t2 ) t2 = t3, b = 3; k = t0 > t2 ? a : b; val = k; t0 = GET_VALUE(4); t1 = GET_VALUE(5); t2 = GET_VALUE(6); t3 = GET_VALUE(7); a = 0, b = 2; if( t1 > t0 ) t0 = t1, a = 1; if( t3 > t2 ) t2 = t3, b = 3; k = t0 > t2 ? a : b; val |= k << 2; t0 = GET_VALUE(8); t1 = GET_VALUE(9); t2 = GET_VALUE(10); t3 = GET_VALUE(11); a = 0, b = 2; if( t1 > t0 ) t0 = t1, a = 1; if( t3 > t2 ) t2 = t3, b = 3; k = t0 > t2 ? a : b; val |= k << 4; t0 = GET_VALUE(12); t1 = GET_VALUE(13); t2 = GET_VALUE(14); t3 = GET_VALUE(15); a = 0, b = 2; if( t1 > t0 ) t0 = t1, a = 1; if( t3 > t2 ) t2 = t3, b = 3; k = t0 > t2 ? a : b; val |= k << 6; pattern += 16*2; #else #error "unknown/undefined WTA_K value; should be 2, 3 or 4" #endif desc[i] = (uchar)val; } } } #endif a7c2cfaeda19907b637211b1cc91d253inline int cornerScore(__global const uchar* img, int step) { int k, tofs, v = img[0], a0 = 0, b0; int d[16]; #define LOAD2(idx, ofs) \ tofs = ofs; d[idx] = (short)(v - img[tofs]); d[idx+8] = (short)(v - img[-tofs]) LOAD2(0, 3); LOAD2(1, -step+3); LOAD2(2, -step*2+2); LOAD2(3, -step*3+1); LOAD2(4, -step*3); LOAD2(5, -step*3-1); LOAD2(6, -step*2-2); LOAD2(7, -step-3); #pragma unroll for( k = 0; k < 16; k += 2 ) { int a = min((int)d[(k+1)&15], (int)d[(k+2)&15]); a = min(a, (int)d[(k+3)&15]); a = min(a, (int)d[(k+4)&15]); a = min(a, (int)d[(k+5)&15]); a = min(a, (int)d[(k+6)&15]); a = min(a, (int)d[(k+7)&15]); a = min(a, (int)d[(k+8)&15]); a0 = max(a0, min(a, (int)d[k&15])); a0 = max(a0, min(a, (int)d[(k+9)&15])); } b0 = -a0; #pragma unroll for( k = 0; k < 16; k += 2 ) { int b = max((int)d[(k+1)&15], (int)d[(k+2)&15]); b = max(b, (int)d[(k+3)&15]); b = max(b, (int)d[(k+4)&15]); b = max(b, (int)d[(k+5)&15]); b = max(b, (int)d[(k+6)&15]); b = max(b, (int)d[(k+7)&15]); b = max(b, (int)d[(k+8)&15]); b0 = min(b0, max(b, (int)d[k])); b0 = min(b0, max(b, (int)d[(k+9)&15])); } return -b0-1; } __kernel void FAST_findKeypoints( __global const uchar * _img, int step, int img_offset, int img_rows, int img_cols, volatile __global int* kp_loc, int max_keypoints, int threshold ) { int j = get_global_id(0) + 3; int i = get_global_id(1) + 3; if (i < img_rows - 3 && j < img_cols - 3) { __global const uchar* img = _img + mad24(i, step, j + img_offset); int v = img[0], t0 = v - threshold, t1 = v + threshold; int k, tofs, v0, v1; int m0 = 0, m1 = 0; #define UPDATE_MASK(idx, ofs) \ tofs = ofs; v0 = img[tofs]; v1 = img[-tofs]; \ m0 |= ((v0 < t0) << idx) | ((v1 < t0) << (8 + idx)); \ m1 |= ((v0 > t1) << idx) | ((v1 > t1) << (8 + idx)) UPDATE_MASK(0, 3); if( (m0 | m1) == 0 ) return; UPDATE_MASK(2, -step*2+2); UPDATE_MASK(4, -step*3); UPDATE_MASK(6, -step*2-2); #define EVEN_MASK (1+4+16+64) if( ((m0 | (m0 >> 8)) & EVEN_MASK) != EVEN_MASK && ((m1 | (m1 >> 8)) & EVEN_MASK) != EVEN_MASK ) return; UPDATE_MASK(1, -step+3); UPDATE_MASK(3, -step*3+1); UPDATE_MASK(5, -step*3-1); UPDATE_MASK(7, -step-3); if( ((m0 | (m0 >> 8)) & 255) != 255 && ((m1 | (m1 >> 8)) & 255) != 255 ) return; m0 |= m0 << 16; m1 |= m1 << 16; #define CHECK0(i) ((m0 & (511 << i)) == (511 << i)) #define CHECK1(i) ((m1 & (511 << i)) == (511 << i)) if( CHECK0(0) + CHECK0(1) + CHECK0(2) + CHECK0(3) + CHECK0(4) + CHECK0(5) + CHECK0(6) + CHECK0(7) + CHECK0(8) + CHECK0(9) + CHECK0(10) + CHECK0(11) + CHECK0(12) + CHECK0(13) + CHECK0(14) + CHECK0(15) + CHECK1(0) + CHECK1(1) + CHECK1(2) + CHECK1(3) + CHECK1(4) + CHECK1(5) + CHECK1(6) + CHECK1(7) + CHECK1(8) + CHECK1(9) + CHECK1(10) + CHECK1(11) + CHECK1(12) + CHECK1(13) + CHECK1(14) + CHECK1(15) == 0 ) return; { int idx = atomic_inc(kp_loc); if( idx < max_keypoints ) { kp_loc[1 + 2*idx] = j; kp_loc[2 + 2*idx] = i; } } } } __kernel void FAST_nonmaxSupression( __global const int* kp_in, volatile __global int* kp_out, __global const uchar * _img, int step, int img_offset, int rows, int cols, int counter, int max_keypoints) { const int idx = get_global_id(0); if (idx < counter) { int x = kp_in[1 + 2*idx]; int y = kp_in[2 + 2*idx]; __global const uchar* img = _img + mad24(y, step, x + img_offset); int s = cornerScore(img, step); if( (x < 4 || s > cornerScore(img-1, step)) + (y < 4 || s > cornerScore(img-step, step)) != 2 ) return; if( (x >= cols - 4 || s > cornerScore(img+1, step)) + (y >= rows - 4 || s > cornerScore(img+step, step)) + (x < 4 || y < 4 || s > cornerScore(img-step-1, step)) + (x >= cols - 4 || y < 4 || s > cornerScore(img-step+1, step)) + (x < 4 || y >= rows - 4 || s > cornerScore(img+step-1, step)) + (x >= cols - 4 || y >= rows - 4 || s > cornerScore(img+step+1, step)) == 6) { int new_idx = atomic_inc(kp_out); if( new_idx < max_keypoints ) { kp_out[1 + 3*new_idx] = x; kp_out[2 + 3*new_idx] = y; kp_out[3 + 3*new_idx] = s; } } } } f5e6f463f21a7ed77bd4d2c753478305#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics:enable #define MAX_FLOAT 3.40282e+038f #ifndef T #define T float #endif #ifndef BLOCK_SIZE #define BLOCK_SIZE 16 #endif #ifndef MAX_DESC_LEN #define MAX_DESC_LEN 64 #endif #define BLOCK_SIZE_ODD (BLOCK_SIZE + 1) #ifndef SHARED_MEM_SZ # if (BLOCK_SIZE < MAX_DESC_LEN) # define SHARED_MEM_SZ (kercn * (BLOCK_SIZE * MAX_DESC_LEN + BLOCK_SIZE * BLOCK_SIZE)) # else # define SHARED_MEM_SZ (kercn * 2 * BLOCK_SIZE_ODD * BLOCK_SIZE) # endif #endif #ifndef DIST_TYPE #define DIST_TYPE 2 #endif #if (DIST_TYPE == 2) # ifdef T_FLOAT typedef float result_type; # if (8 == kercn) typedef float8 value_type; # define DIST(x, y) {value_type d = fabs((x) - (y)); result += d.s0 + d.s1 + d.s2 + d.s3 + d.s4 + d.s5 + d.s6 + d.s7;} # elif (4 == kercn) typedef float4 value_type; # define DIST(x, y) {value_type d = fabs((x) - (y)); result += d.s0 + d.s1 + d.s2 + d.s3;} # else typedef float value_type; # define DIST(x, y) result += fabs((x) - (y)) # endif # else typedef int result_type; # if (8 == kercn) typedef int8 value_type; # define DIST(x, y) {value_type d = abs((x) - (y)); result += d.s0 + d.s1 + d.s2 + d.s3 + d.s4 + d.s5 + d.s6 + d.s7;} # elif (4 == kercn) typedef int4 value_type; # define DIST(x, y) {value_type d = abs((x) - (y)); result += d.s0 + d.s1 + d.s2 + d.s3;} # else typedef int value_type; # define DIST(x, y) result += abs((x) - (y)) # endif # endif # define DIST_RES(x) (x) #elif (DIST_TYPE == 4) typedef float result_type; # if (8 == kercn) typedef float8 value_type; # define DIST(x, y) {value_type d = ((x) - (y)); result += dot(d.s0123, d.s0123) + dot(d.s4567, d.s4567);} # elif (4 == kercn) typedef float4 value_type; # define DIST(x, y) {value_type d = ((x) - (y)); result += dot(d, d);} # else typedef float value_type; # define DIST(x, y) {value_type d = ((x) - (y)); result = mad(d, d, result);} # endif # define DIST_RES(x) sqrt(x) #elif (DIST_TYPE == 6) # if (8 == kercn) typedef int8 value_type; # elif (4 == kercn) typedef int4 value_type; # else typedef int value_type; # endif typedef int result_type; # define DIST(x, y) result += popcount( (x) ^ (y) ) # define DIST_RES(x) (x) #endif inline result_type reduce_block( __local value_type *s_query, __local value_type *s_train, int lidx, int lidy ) { result_type result = 0; #pragma unroll for (int j = 0 ; j < BLOCK_SIZE ; j++) { DIST(s_query[lidy * BLOCK_SIZE_ODD + j], s_train[j * BLOCK_SIZE_ODD + lidx]); } return DIST_RES(result); } inline result_type reduce_block_match( __local value_type *s_query, __local value_type *s_train, int lidx, int lidy ) { result_type result = 0; #pragma unroll for (int j = 0 ; j < BLOCK_SIZE ; j++) { DIST(s_query[lidy * BLOCK_SIZE_ODD + j], s_train[j * BLOCK_SIZE_ODD + lidx]); } return result; } inline result_type reduce_multi_block( __local value_type *s_query, __local value_type *s_train, int block_index, int lidx, int lidy ) { result_type result = 0; #pragma unroll for (int j = 0 ; j < BLOCK_SIZE ; j++) { DIST(s_query[lidy * MAX_DESC_LEN + block_index * BLOCK_SIZE + j], s_train[j * BLOCK_SIZE + lidx]); } return result; } __kernel void BruteForceMatch_Match( __global T *query, __global T *train, __global int *bestTrainIdx, __global float *bestDistance, int query_rows, int query_cols, int train_rows, int train_cols, int step ) { const int lidx = get_local_id(0); const int lidy = get_local_id(1); const int groupidx = get_group_id(0); const int queryIdx = mad24(BLOCK_SIZE, groupidx, lidy); const int queryOffset = min(queryIdx, query_rows - 1) * step; __global TN *query_vec = (__global TN *)(query + queryOffset); query_cols /= kercn; __local float sharebuffer[SHARED_MEM_SZ]; __local value_type *s_query = (__local value_type *)sharebuffer; #if 0 < MAX_DESC_LEN __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE * MAX_DESC_LEN; #pragma unroll for (int i = 0; i < MAX_DESC_LEN / BLOCK_SIZE; i++) { const int loadx = mad24(BLOCK_SIZE, i, lidx); s_query[mad24(MAX_DESC_LEN, lidy, loadx)] = loadx < query_cols ? query_vec[loadx] : 0; } #else __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE_ODD * BLOCK_SIZE; const int s_query_i = mad24(BLOCK_SIZE_ODD, lidy, lidx); const int s_train_i = mad24(BLOCK_SIZE_ODD, lidx, lidy); #endif float myBestDistance = MAX_FLOAT; int myBestTrainIdx = -1; for (int t = 0, endt = (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE; t < endt; t++) { result_type result = 0; const int trainOffset = min(mad24(BLOCK_SIZE, t, lidy), train_rows - 1) * step; __global TN *train_vec = (__global TN *)(train + trainOffset); #if 0 < MAX_DESC_LEN #pragma unroll for (int i = 0; i < MAX_DESC_LEN / BLOCK_SIZE; i++) { const int loadx = mad24(BLOCK_SIZE, i, lidx); s_train[mad24(BLOCK_SIZE, lidx, lidy)] = loadx < train_cols ? train_vec[loadx] : 0; barrier(CLK_LOCAL_MEM_FENCE); result += reduce_multi_block(s_query, s_train, i, lidx, lidy); barrier(CLK_LOCAL_MEM_FENCE); } #else for (int i = 0, endq = (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE; i < endq; i++) { const int loadx = mad24(i, BLOCK_SIZE, lidx); if (loadx < query_cols) { s_query[s_query_i] = query_vec[loadx]; s_train[s_train_i] = train_vec[loadx]; } else { s_query[s_query_i] = 0; s_train[s_train_i] = 0; } barrier(CLK_LOCAL_MEM_FENCE); result += reduce_block_match(s_query, s_train, lidx, lidy); barrier(CLK_LOCAL_MEM_FENCE); } #endif result = DIST_RES(result); const int trainIdx = mad24(BLOCK_SIZE, t, lidx); if (queryIdx < query_rows && trainIdx < train_rows && result < myBestDistance ) { myBestDistance = result; myBestTrainIdx = trainIdx; } } barrier(CLK_LOCAL_MEM_FENCE); __local float *s_distance = (__local float *)sharebuffer; __local int *s_trainIdx = (__local int *)(sharebuffer + BLOCK_SIZE_ODD * BLOCK_SIZE); s_distance += lidy * BLOCK_SIZE_ODD; s_trainIdx += lidy * BLOCK_SIZE_ODD; s_distance[lidx] = myBestDistance; s_trainIdx[lidx] = myBestTrainIdx; barrier(CLK_LOCAL_MEM_FENCE); #pragma unroll for (int k = 0 ; k < BLOCK_SIZE; k++) { if (myBestDistance > s_distance[k]) { myBestDistance = s_distance[k]; myBestTrainIdx = s_trainIdx[k]; } } if (queryIdx < query_rows && lidx == 0) { bestTrainIdx[queryIdx] = myBestTrainIdx; bestDistance[queryIdx] = myBestDistance; } } __kernel void BruteForceMatch_RadiusMatch( __global T *query, __global T *train, float maxDistance, __global int *bestTrainIdx, __global float *bestDistance, __global int *nMatches, int query_rows, int query_cols, int train_rows, int train_cols, int bestTrainIdx_cols, int step, int ostep ) { const int lidx = get_local_id(0); const int lidy = get_local_id(1); const int groupidx = get_group_id(0); const int groupidy = get_group_id(1); const int queryIdx = mad24(BLOCK_SIZE, groupidy, lidy); const int queryOffset = min(queryIdx, query_rows - 1) * step; __global TN *query_vec = (__global TN *)(query + queryOffset); const int trainIdx = mad24(BLOCK_SIZE, groupidx, lidx); const int trainOffset = min(mad24(BLOCK_SIZE, groupidx, lidy), train_rows - 1) * step; __global TN *train_vec = (__global TN *)(train + trainOffset); query_cols /= kercn; __local float sharebuffer[SHARED_MEM_SZ]; __local value_type *s_query = (__local value_type *)sharebuffer; __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE_ODD * BLOCK_SIZE; result_type result = 0; const int s_query_i = mad24(BLOCK_SIZE_ODD, lidy, lidx); const int s_train_i = mad24(BLOCK_SIZE_ODD, lidx, lidy); for (int i = 0 ; i < (query_cols + BLOCK_SIZE - 1) / BLOCK_SIZE ; ++i) { const int loadx = mad24(BLOCK_SIZE, i, lidx); if (loadx < query_cols) { s_query[s_query_i] = query_vec[loadx]; s_train[s_train_i] = train_vec[loadx]; } else { s_query[s_query_i] = 0; s_train[s_train_i] = 0; } barrier(CLK_LOCAL_MEM_FENCE); result += reduce_block(s_query, s_train, lidx, lidy); barrier(CLK_LOCAL_MEM_FENCE); } if (queryIdx < query_rows && trainIdx < train_rows && convert_float(result) < maxDistance) { int ind = atom_inc(nMatches + queryIdx); if(ind < bestTrainIdx_cols) { bestTrainIdx[mad24(queryIdx, ostep, ind)] = trainIdx; bestDistance[mad24(queryIdx, ostep, ind)] = result; } } } __kernel void BruteForceMatch_knnMatch( __global T *query, __global T *train, __global int2 *bestTrainIdx, __global float2 *bestDistance, int query_rows, int query_cols, int train_rows, int train_cols, int step ) { const int lidx = get_local_id(0); const int lidy = get_local_id(1); const int groupidx = get_group_id(0); const int queryIdx = mad24(BLOCK_SIZE, groupidx, lidy); const int queryOffset = min(queryIdx, query_rows - 1) * step; __global TN *query_vec = (__global TN *)(query + queryOffset); query_cols /= kercn; __local float sharebuffer[SHARED_MEM_SZ]; __local value_type *s_query = (__local value_type *)sharebuffer; #if 0 < MAX_DESC_LEN __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE * MAX_DESC_LEN; #pragma unroll for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE; i ++) { int loadx = mad24(BLOCK_SIZE, i, lidx); s_query[mad24(MAX_DESC_LEN, lidy, loadx)] = loadx < query_cols ? query_vec[loadx] : 0; } #else __local value_type *s_train = (__local value_type *)sharebuffer + BLOCK_SIZE_ODD * BLOCK_SIZE; const int s_query_i = mad24(BLOCK_SIZE_ODD, lidy, lidx); const int s_train_i = mad24(BLOCK_SIZE_ODD, lidx, lidy); #endif float myBestDistance1 = MAX_FLOAT; float myBestDistance2 = MAX_FLOAT; int myBestTrainIdx1 = -1; int myBestTrainIdx2 = -1; for (int t = 0, endt = (train_rows + BLOCK_SIZE - 1) / BLOCK_SIZE; t < endt ; t++) { result_type result = 0; int trainOffset = min(mad24(BLOCK_SIZE, t, lidy), train_rows - 1) * step; __global TN *train_vec = (__global TN *)(train + trainOffset); #if 0 < MAX_DESC_LEN #pragma unroll for (int i = 0 ; i < MAX_DESC_LEN / BLOCK_SIZE ; i++) { const int loadx = mad24(BLOCK_SIZE, i, lidx); s_train[mad24(BLOCK_SIZE, lidx, lidy)] = loadx < train_cols ? train_vec[loadx] : 0; barrier(CLK_LOCAL_MEM_FENCE); result += reduce_multi_block(s_query, s_train, i, lidx, lidy); barrier(CLK_LOCAL_MEM_FENCE); } #else for (int i = 0, endq = (query_cols + BLOCK_SIZE -1) / BLOCK_SIZE; i < endq ; i++) { const int loadx = mad24(BLOCK_SIZE, i, lidx); if (loadx < query_cols) { s_query[s_query_i] = query_vec[loadx]; s_train[s_train_i] = train_vec[loadx]; } else { s_query[s_query_i] = 0; s_train[s_train_i] = 0; } barrier(CLK_LOCAL_MEM_FENCE); result += reduce_block_match(s_query, s_train, lidx, lidy); barrier(CLK_LOCAL_MEM_FENCE); } #endif result = DIST_RES(result); const int trainIdx = mad24(BLOCK_SIZE, t, lidx); if (queryIdx < query_rows && trainIdx < train_rows) { if (result < myBestDistance1) { myBestDistance2 = myBestDistance1; myBestTrainIdx2 = myBestTrainIdx1; myBestDistance1 = result; myBestTrainIdx1 = trainIdx; } else if (result < myBestDistance2) { myBestDistance2 = result; myBestTrainIdx2 = trainIdx; } } } barrier(CLK_LOCAL_MEM_FENCE); __local float *s_distance = (__local float *)sharebuffer; __local int *s_trainIdx = (__local int *)(sharebuffer + BLOCK_SIZE_ODD * BLOCK_SIZE); s_distance += lidy * BLOCK_SIZE_ODD; s_trainIdx += lidy * BLOCK_SIZE_ODD; s_distance[lidx] = myBestDistance1; s_trainIdx[lidx] = myBestTrainIdx1; float bestDistance1 = MAX_FLOAT; float bestDistance2 = MAX_FLOAT; int bestTrainIdx1 = -1; int bestTrainIdx2 = -1; barrier(CLK_LOCAL_MEM_FENCE); if (lidx == 0) { for (int i = 0 ; i < BLOCK_SIZE ; i++) { float val = s_distance[i]; if (val < bestDistance1) { bestDistance2 = bestDistance1; bestTrainIdx2 = bestTrainIdx1; bestDistance1 = val; bestTrainIdx1 = s_trainIdx[i]; } else if (val < bestDistance2) { bestDistance2 = val; bestTrainIdx2 = s_trainIdx[i]; } } } barrier(CLK_LOCAL_MEM_FENCE); s_distance[lidx] = myBestDistance2; s_trainIdx[lidx] = myBestTrainIdx2; barrier(CLK_LOCAL_MEM_FENCE); if (lidx == 0) { for (int i = 0 ; i < BLOCK_SIZE ; i++) { float val = s_distance[i]; if (val < bestDistance2) { bestDistance2 = val; bestTrainIdx2 = s_trainIdx[i]; } } } myBestDistance1 = bestDistance1; myBestDistance2 = bestDistance2; myBestTrainIdx1 = bestTrainIdx1; myBestTrainIdx2 = bestTrainIdx2; if (queryIdx < query_rows && lidx == 0) { bestTrainIdx[queryIdx] = (int2)(myBestTrainIdx1, myBestTrainIdx2); bestDistance[queryIdx] = (float2)(myBestDistance1, myBestDistance2); } } 35c3a1e231d446e4088561e3604fb94f__kernel void AKAZE_pm_g2(__global const float* lx, __global const float* ly, __global float* dst, float k, int size) { int i = get_global_id(0); if (!(i < size)) { return; } const float k2inv = 1.0f / (k * k); dst[i] = 1.0f / (1.0f + ((lx[i] * lx[i] + ly[i] * ly[i]) * k2inv)); } __kernel void AKAZE_nld_step_scalar(__global const float* lt, int lt_step, int lt_offset, int rows, int cols, __global const float* lf, __global float* dst, float step_size) { int i = get_global_id(1); int j = get_global_id(0); if (!(i < rows && j < cols)) { return; } int a = (i - 1) * cols; int c = (i ) * cols; int b = (i + 1) * cols; float res = 0.0f; if (i == 0) { if (j == 0 || j == (cols - 1)) { res = 0.0f; } else { res = (lf[c + j] + lf[c + j + 1])*(lt[c + j + 1] - lt[c + j]) + (lf[c + j] + lf[c + j - 1])*(lt[c + j - 1] - lt[c + j]) + (lf[c + j] + lf[b + j ])*(lt[b + j ] - lt[c + j]); } } else if (i == (rows - 1)) { if (j == 0 || j == (cols - 1)) { res = 0.0f; } else { res = (lf[c + j] + lf[c + j + 1])*(lt[c + j + 1] - lt[c + j]) + (lf[c + j] + lf[c + j - 1])*(lt[c + j - 1] - lt[c + j]) + (lf[c + j] + lf[a + j ])*(lt[a + j ] - lt[c + j]); } } else { if (j == 0) { res = (lf[c + 0] + lf[c + 1])*(lt[c + 1] - lt[c + 0]) + (lf[c + 0] + lf[b + 0])*(lt[b + 0] - lt[c + 0]) + (lf[c + 0] + lf[a + 0])*(lt[a + 0] - lt[c + 0]); } else if (j == (cols - 1)) { res = (lf[c + j] + lf[c + j - 1])*(lt[c + j - 1] - lt[c + j]) + (lf[c + j] + lf[b + j ])*(lt[b + j ] - lt[c + j]) + (lf[c + j] + lf[a + j ])*(lt[a + j ] - lt[c + j]); } else { res = (lf[c + j] + lf[c + j + 1])*(lt[c + j + 1] - lt[c + j]) + (lf[c + j] + lf[c + j - 1])*(lt[c + j - 1] - lt[c + j]) + (lf[c + j] + lf[b + j ])*(lt[b + j ] - lt[c + j]) + (lf[c + j] + lf[a + j ])*(lt[a + j ] - lt[c + j]); } } dst[c + j] = res * step_size; } __kernel void AKAZE_compute_determinant(__global const float* lxx, __global const float* lxy, __global const float* lyy, __global float* dst, float sigma, int size) { int i = get_global_id(0); if (!(i < size)) { return; } dst[i] = (lxx[i] * lyy[i] - lxy[i] * lxy[i]) * sigma; } 80f6cd2f334b70062ed64a0a1a866593GCC: (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0GNUÀñÿ  (C(h (š (opencl_kernels_features2d.cpp_ZN2cv3ocl10features2d10orb_oclsrcE_ZN2cv3ocl10features2d11fast_oclsrcE_ZN2cv3ocl10features2d24brute_force_match_oclsrcE_ZN2cv3ocl10features2d12akaze_oclsrcE ¸à&8&€V&¨Và^.symtab.strtab.shstrtab.text.data.bss.rodata.str1.1.rodata.str1.8.rela.data.rel.local._ZN2cv3ocl10features2d10orb_oclsrcE.rela.data.rel.local._ZN2cv3ocl10features2d11fast_oclsrcE.rela.data.rel.local._ZN2cv3ocl10features2d24brute_force_match_oclsrcE.rela.data.rel.local._ZN2cv3ocl10features2d12akaze_oclsrcE.comment.note.GNU-stack.note.gnu.property@!@'@,2@,;2p_O€_( J@(c`ˆÀ_( ƒ@ˆc`Â`( ½@èc`  @`( @Hd` ?0h`,H”`X˜` ¸`° hbÀ¨dk