ELF> @@trackingtmmtldDetector#define LOCAL_SIZE_X 64 #define BLOCK_SIZE_X 3 __kernel void tmm(__global float *A, int m, int n, float alpha, __global float *D) { int lidX = get_local_id(0); uint lsizeX = get_local_size(0); uint matI = get_group_id(1); uint matJ = get_group_id(0); if (matI < matJ) return; __local float4 a[LOCAL_SIZE_X], b[LOCAL_SIZE_X]; float4 result; __local uint cnt; result = 0; cnt = 0; barrier(CLK_LOCAL_MEM_FENCE); do { int global_block_base = (lidX + cnt * lsizeX) * BLOCK_SIZE_X; float4 pa[BLOCK_SIZE_X], pb[BLOCK_SIZE_X]; #pragma unroll for(uint j = 0; j < BLOCK_SIZE_X && (cnt * lsizeX + lidX) * BLOCK_SIZE_X < n / 4; j++) { pa[j] = *(__global float4*)&A[matI * n + (global_block_base + j) * 4]; if (matI != matJ) pb[j] = *(__global float4*)&A[matJ * n + (global_block_base + j) * 4]; else pb[j] = pa[j]; } if (global_block_base + BLOCK_SIZE_X - 1 >= n/4) { #pragma unroll for(int i = 0; i < BLOCK_SIZE_X; i++) { if (global_block_base + i >= n/4) pb[i] = 0; } } pb[0] *= pa[0]; for(int j = 1; j < BLOCK_SIZE_X; j++) pb[0] = fma(pb[j], pa[j], pb[0]); b[lidX] = pb[0]; barrier(CLK_LOCAL_MEM_FENCE); for(int offset = LOCAL_SIZE_X / 2; offset > 0; offset >>= 1) { if (lidX < offset) b[lidX] += b[(lidX + offset)]; barrier(CLK_LOCAL_MEM_FENCE); } if (lidX == 0) { result += b[0]; cnt++; } barrier(CLK_LOCAL_MEM_FENCE); } while(cnt * BLOCK_SIZE_X * lsizeX < n / 4); if (lidX == 0) { float ret = (result.s0 + result.s1 + result.s2 + result.s3) * alpha; D[matI * m + matJ] = ret; if (matI != matJ) D[matJ * m + matI] = ret; } } 1968d33b7c3068a6469896acdba60dee__kernel void NCC(__global const uchar *patch, __global const uchar *positiveSamples, __global const uchar *negativeSamples, __global float *ncc, int posNum, int negNum) { int id = get_global_id(0); if (id >= 1000) return; bool posFlg; if (id < 500) posFlg = true; if (id >= 500) { id = id - 500; posFlg = false; } int s1 = 0, s2 = 0, n1 = 0, n2 = 0, prod = 0; float sq1 = 0, sq2 = 0, ares = 0; int N = 225; if (posFlg && id < posNum) { for (int i = 0; i < N; i++) { s1 += positiveSamples[id * N + i]; s2 += patch[i]; n1 += positiveSamples[id * N + i] * positiveSamples[id * N + i]; n2 += patch[i] * patch[i]; prod += positiveSamples[id * N + i] * patch[i]; } sq1 = sqrt(max(0.0, n1 - 1.0 * s1 * s1 / N)); sq2 = sqrt(max(0.0, n2 - 1.0 * s2 * s2 / N)); ares = (sq2 == 0) ? sq1 / fabs(sq1) : (prod - s1 * s2 / N) / sq1 / sq2; ncc[id] = ares; } if (!posFlg && id < negNum) { for (int i = 0; i < N; i++) { s1 += negativeSamples[id * N + i]; s2 += patch[i]; n1 += negativeSamples[id * N + i] * negativeSamples[id * N + i]; n2 += patch[i] * patch[i]; prod += negativeSamples[id * N + i] * patch[i]; } sq1 = sqrt(max(0.0, n1 - 1.0 * s1 * s1 / N)); sq2 = sqrt(max(0.0, n2 - 1.0 * s2 * s2 / N)); ares = (sq2 == 0) ? sq1 / fabs(sq1) : (prod - s1 * s2 / N) / sq1 / sq2; ncc[id+500] = ares; } } __kernel void batchNCC(__global const uchar *patches, __global const uchar *positiveSamples, __global const uchar *negativeSamples, __global float *posNcc, __global float *negNcc, int posNum, int negNum, int patchNum) { int id = get_global_id(0); bool posFlg; if (id < 500*patchNum) posFlg = true; if (id >= 500*patchNum) { id = id - 500*patchNum; posFlg = false; } int modelSampleID = id % 500; int patchID = id / 500; int s1 = 0, s2 = 0, n1 = 0, n2 = 0, prod = 0; float sq1 = 0, sq2 = 0, ares = 0; int N = 225; if (posFlg && modelSampleID < posNum) { for (int i = 0; i < N; i++) { s1 += positiveSamples[modelSampleID * N + i]; s2 += patches[patchID*N + i]; n1 += positiveSamples[modelSampleID * N + i] * positiveSamples[modelSampleID * N + i]; n2 += patches[patchID*N + i] * patches[patchID*N + i]; prod += positiveSamples[modelSampleID * N + i] * patches[patchID*N + i]; } sq1 = sqrt(max(0.0, n1 - 1.0 * s1 * s1 / N)); sq2 = sqrt(max(0.0, n2 - 1.0 * s2 * s2 / N)); ares = (sq2 == 0) ? sq1 / fabs(sq1) : (prod - s1 * s2 / N) / sq1 / sq2; posNcc[id] = ares; } if (!posFlg && modelSampleID < negNum) { for (int i = 0; i < N; i++) { s1 += negativeSamples[modelSampleID * N + i]; s2 += patches[patchID*N + i]; n1 += negativeSamples[modelSampleID * N + i] * negativeSamples[modelSampleID * N + i]; n2 += patches[patchID*N + i] * patches[patchID*N + i]; prod += negativeSamples[modelSampleID * N + i] * patches[patchID*N + i]; } sq1 = sqrt(max(0.0, n1 - 1.0 * s1 * s1 / N)); sq2 = sqrt(max(0.0, n2 - 1.0 * s2 * s2 / N)); ares = (sq2 == 0) ? sq1 / fabs(sq1) : (prod - s1 * s2 / N) / sq1 / sq2; negNcc[id] = ares; } } f15ba33f44bb0051df83a2b52cf9d67fGCC: (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0GNUÀñÿ   (>(opencl_kernels_tracking.cpp_ZN2cv3ocl8tracking10tmm_oclsrcE_ZN2cv3ocl8tracking18tldDetector_oclsrcE ø  p.symtab.strtab.shstrtab.text.data.bss.rodata.str1.1.rodata.str1.8.rela.data.rel.local._ZN2cv3ocl8tracking10tmm_oclsrcE.rela.data.rel.local._ZN2cv3ocl8tracking18tldDetector_oclsrcE.comment.note.GNU-stack.note.gnu.property@!@'@,2@;2`‘O( J@p` …@( €@Ð` ¾0h,ǔט ¸P  g0ê