/external/opencv3/modules/cudev/include/opencv2/cudev/warp/ |
scan.hpp | 59 __device__ T warpScanInclusive(T data, volatile T* smem, uint tid) 62 (void) smem; 79 smem[pos] = 0; 82 smem[pos] = data; 84 smem[pos] += smem[pos - 1]; 85 smem[pos] += smem[pos - 2]; 86 smem[pos] += smem[pos - 4] [all...] |
/external/opencv3/modules/cudev/include/opencv2/cudev/block/ |
scan.hpp | 58 __device__ T blockScanInclusive(T data, volatile T* smem, uint tid) 63 T warpResult = warpScanInclusive(data, smem, tid); 71 smem[tid >> LOG_WARP_SIZE] = warpResult; 79 T val = smem[tid]; 82 smem[tid] = warpScanExclusive(val, smem, tid); 88 return warpResult + smem[tid >> LOG_WARP_SIZE]; 92 return warpScanInclusive(data, smem, tid); 97 __device__ __forceinline__ T blockScanExclusive(T data, volatile T* smem, uint tid) 99 return blockScanInclusive<THREADS_NUM>(data, smem, tid) - data [all...] |
vec_distance.hpp | 75 __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid) 77 warpReduce(smem, mySum, tid, plus<result_type>()); 80 template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid) 82 blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>()); 104 __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid) 106 warpReduce(smem, mySum, tid, plus<result_type>()); 109 template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid) 111 blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>()); 137 __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid) 139 warpReduce(smem, mySum, tid, plus<result_type>()) [all...] |
/external/opencv3/modules/imgproc/src/opencl/ |
clahe.cl | 50 inline int calc_lut(__local int* smem, int val, int tid) 52 smem[tid] = val; 57 smem[i] += smem[i - 1]; 60 return smem[tid]; 64 inline void reduce(volatile __local int* smem, int val, int tid) 66 smem[tid] = val; 70 smem[tid] = val += smem[tid + 128]; 74 smem[tid] = val += smem[tid + 64] [all...] |
pyr_down.cl | 104 smem[0][col_lcl] = sum0; \ 107 smem[1][col_lcl] = sum1; 123 vstore4(sum40, col_lcl, (__local float*) &smem[0][2]); \ 126 vstore4(sum41, col_lcl, (__local float*) &smem[1][2]); 137 __local FT smem[2][LOCAL_SIZE + 4]; 229 FT sum = dot(vload4(0, (__local float*) (&smem) + tid2 + (yin - y) * (LOCAL_SIZE + 4)), (float4)(co3, co2, co1, co2)); 231 FT sum = dot(vload4(0, (__local double*) (&smem) + tid2 + (yin - y) * (LOCAL_SIZE + 4)), (double4)(co3, co2, co1, co2)); 234 FT sum = co3 * smem[yin - y][2 + tid2 - 2]; 235 sum = MAD(co2, smem[yin - y][2 + tid2 - 1], sum); 236 sum = MAD(co1, smem[yin - y][2 + tid2 ], sum) [all...] |
canny.cl | 80 inline float3 sobel(int idx, __local const floatN *smem) 85 floatN dx = fma(2, smem[idx + GRP_SIZEX + 6] - smem[idx + GRP_SIZEX + 4], 86 smem[idx + 2] - smem[idx] + smem[idx + 2 * GRP_SIZEX + 10] - smem[idx + 2 * GRP_SIZEX + 8]); 88 floatN dy = fma(2, smem[idx + 1] - smem[idx + 2 * GRP_SIZEX + 9], 89 smem[idx + 2] - smem[idx + 2 * GRP_SIZEX + 10] + smem[idx] - smem[idx + 2 * GRP_SIZEX + 8]) [all...] |
/external/opencv3/modules/objdetect/src/opencl/ |
objdetect_hog.cl | 72 __global float* block_hists, __local float* smem) 87 __local float* hists = smem + lp * cnbins * (CELLS_PER_BLOCK_X * 172 __local float* smem = squares + boffset; 173 float sum = smem[hid]; 175 smem[hid] = sum = sum + smem[hid + 18]; 178 smem[hid] = sum = sum + smem[hid + 9]; 181 smem[hid] = sum + smem[hid + 4] [all...] |
/external/opencv3/modules/cudev/include/opencv2/cudev/grid/detail/ |
reduce_to_column.hpp | 61 __device__ __forceinline__ static void call(work_elem_type smem[1][BLOCK_SIZE], work_type& myVal) 64 blockReduce<BLOCK_SIZE>(smem[0], myVal, threadIdx.x, op); 70 __device__ __forceinline__ static void call(work_elem_type smem[2][BLOCK_SIZE], work_type& myVal) 73 blockReduce<BLOCK_SIZE>(smem_tuple(smem[0], smem[1]), tie(myVal.x, myVal.y), threadIdx.x, make_tuple(op, op)); 79 __device__ __forceinline__ static void call(work_elem_type smem[3][BLOCK_SIZE], work_type& myVal) 82 blockReduce<BLOCK_SIZE>(smem_tuple(smem[0], smem[1], smem[2]), tie(myVal.x, myVal.y, myVal.z), threadIdx.x, make_tuple(op, op, op)); 88 __device__ __forceinline__ static void call(work_elem_type smem[4][BLOCK_SIZE], work_type& myVal 102 __shared__ work_elem_type smem[cn][BLOCK_SIZE]; local [all...] |
histogram.hpp | 60 __shared__ ResType smem[BIN_COUNT]; local 66 smem[i] = 0; 77 atomicAdd(&smem[data % BIN_COUNT], 1); 86 const ResType histVal = smem[i];
|
reduce_to_row.hpp | 62 __shared__ work_type smem[BLOCK_SIZE_X * BLOCK_SIZE_Y]; local 81 smem[threadIdx.x * BLOCK_SIZE_Y + threadIdx.y] = myVal; 85 volatile work_type* srow = smem + threadIdx.y * BLOCK_SIZE_X; 96 dst[x] = saturate_cast<ResType>(Reductor::result(smem[threadIdx.x * BLOCK_SIZE_X], rows));
|
pyr_down.hpp | 68 __shared__ work_type smem[256 + 4]; local 86 smem[2 + threadIdx.x] = sum; 101 smem[threadIdx.x] = sum; 116 smem[4 + threadIdx.x] = sum; 130 smem[2 + threadIdx.x] = sum; 145 smem[threadIdx.x] = sum; 160 smem[4 + threadIdx.x] = sum; 172 sum = 0.0625f * smem[2 + tid2 - 2]; 173 sum = sum + 0.25f * smem[2 + tid2 - 1]; 174 sum = sum + 0.375f * smem[2 + tid2 ] [all...] |
integral.hpp | 63 __shared__ D smem[NUM_SCAN_THREADS * 2]; local 84 const D curScanElem = blockScanInclusive<NUM_SCAN_THREADS>(curElem, smem, threadIdx.x); 105 __shared__ D smem[NUM_SCAN_THREADS * 2]; local 127 const D curScanElem = blockScanInclusive<NUM_SCAN_THREADS>(curElem, smem, threadIdx.x); 445 // place into SMEM 446 // shfl scan reduce the SMEM, reformating so the column 481 __shared__ T smem[32][32]; 484 volatile T* smem_row = &smem[0][0] + 64 * threadIdx.y; 505 smem[threadIdx.y + 0][threadIdx.x] = 0.0f; 506 smem[threadIdx.y + 8][threadIdx.x] = 0.0f [all...] |
/external/opencv3/modules/core/include/opencv2/core/cuda/detail/ |
reduce.hpp | 74 static __device__ void loadToSmem(const PointerTuple& smem, const ValTuple& val, unsigned int tid) 76 thrust::get<I>(smem)[tid] = thrust::get<I>(val); 78 For<I + 1, N>::loadToSmem(smem, val, tid); 81 static __device__ void loadFromSmem(const PointerTuple& smem, const ValTuple& val, unsigned int tid) 83 thrust::get<I>(val) = thrust::get<I>(smem)[tid]; 85 For<I + 1, N>::loadFromSmem(smem, val, tid); 89 static __device__ void merge(const PointerTuple& smem, const ValTuple& val, unsigned int tid, unsigned int delta, const OpTuple& op) 91 typename GetType<typename thrust::tuple_element<I, PointerTuple>::type>::type reg = thrust::get<I>(smem)[tid + delta]; 92 thrust::get<I>(smem)[tid] = thrust::get<I>(val) = thrust::get<I>(op)(thrust::get<I>(val), reg); 94 For<I + 1, N>::merge(smem, val, tid, delta, op) [all...] |
/external/opencv3/modules/cudev/include/opencv2/cudev/block/detail/ |
reduce.hpp | 83 __device__ static void loadToSmem(const PointerTuple& smem, const ValTuple& val, uint tid) 85 get<I>(smem)[tid] = get<I>(val); 87 For<I + 1, N>::loadToSmem(smem, val, tid); 91 __device__ static void loadFromSmem(const PointerTuple& smem, const ValTuple& val, uint tid) 93 get<I>(val) = get<I>(smem)[tid]; 95 For<I + 1, N>::loadFromSmem(smem, val, tid); 99 __device__ static void merge(const PointerTuple& smem, const ValTuple& val, uint tid, uint delta, const OpTuple& op) 101 typename GetType<typename tuple_element<I, PointerTuple>::type>::type reg = get<I>(smem)[tid + delta]; 102 get<I>(smem)[tid] = get<I>(val) = get<I>(op)(get<I>(val), reg); 104 For<I + 1, N>::merge(smem, val, tid, delta, op) [all...] |
/external/opencv3/modules/core/include/opencv2/core/cuda/ |
vec_distance.hpp | 70 template <int THREAD_DIM> __device__ __forceinline__ void reduceAll(int* smem, int tid) 72 reduce<THREAD_DIM>(smem, mySum, tid, plus<int>()); 94 template <int THREAD_DIM> __device__ __forceinline__ void reduceAll(float* smem, int tid) 96 reduce<THREAD_DIM>(smem, mySum, tid, plus<float>()); 120 template <int THREAD_DIM> __device__ __forceinline__ void reduceAll(float* smem, int tid) 122 reduce<THREAD_DIM>(smem, mySum, tid, plus<float>()); 145 template <int THREAD_DIM> __device__ __forceinline__ void reduceAll(int* smem, int tid) 147 reduce<THREAD_DIM>(smem, mySum, tid, plus<int>()); 160 __device__ void calcVecDiffGlobal(const T1* vec1, const T2* vec2, int len, Dist& dist, typename Dist::result_type* smem, int tid) 173 dist.reduceAll<THREAD_DIM>(smem, tid) [all...] |
/external/opencv3/modules/cudev/include/opencv2/cudev/warp/detail/ |
reduce.hpp | 81 __device__ static void loadToSmem(const PointerTuple& smem, const ValTuple& val, uint tid) 83 get<I>(smem)[tid] = get<I>(val); 85 For<I + 1, N>::loadToSmem(smem, val, tid); 89 __device__ static void merge(const PointerTuple& smem, const ValTuple& val, uint tid, uint delta, const OpTuple& op) 91 typename GetType<typename tuple_element<I, PointerTuple>::type>::type reg = get<I>(smem)[tid + delta]; 92 get<I>(smem)[tid] = get<I>(val) = get<I>(op)(get<I>(val), reg); 94 For<I + 1, N>::merge(smem, val, tid, delta, op); 132 __device__ __forceinline__ void loadToSmem(volatile T* smem, T& val, uint tid) 134 smem[tid] = val; 139 __device__ __forceinline__ void loadToSmem(const tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>& smem, [all...] |
/hardware/intel/common/libmix/videoencoder/ |
IntelMetadataBuffer.cpp | 127 static status_t WriteMemObjToBinder(Parcel& data, ShareMemMap* smem) { 129 if (smem->type >= ST_MAX) 134 data.writeInt32(smem->type); 136 if (smem->type == ST_MEMBASE) /*offset, size, heap*/ 140 sp<IMemoryHeap> heap = smem->membase->getMemory(&offset, &size); 146 heap->getHeapID(), smem->membase->pointer(), *((int *)(smem->membase->pointer()))); 149 else if (smem->type == ST_GFX) /*graphicbuffer*/ 150 data.write(*(smem->gbuffer)); 222 static void PushShareMem(ShareMemMap* &smem) 311 ShareMemMap* smem = FindShareMem(sessionflag, value, false); local 706 ShareMemMap smem; local 722 ShareMemMap* smem = new ShareMemMap; local 758 ShareMemMap smem; local 774 ShareMemMap* smem = new ShareMemMap; local [all...] |
/external/opencv3/modules/cudaimgproc/src/cuda/ |
canny.cu | 249 __shared__ volatile int smem[18][18]; 254 smem[threadIdx.y + 1][threadIdx.x + 1] = checkIdx(y, x, map.rows, map.cols) ? map(y, x) : 0; 256 smem[0][threadIdx.x + 1] = checkIdx(y - 1, x, map.rows, map.cols) ? map(y - 1, x) : 0; 258 smem[blockDim.y + 1][threadIdx.x + 1] = checkIdx(y + 1, x, map.rows, map.cols) ? map(y + 1, x) : 0; 260 smem[threadIdx.y + 1][0] = checkIdx(y, x - 1, map.rows, map.cols) ? map(y, x - 1) : 0; 262 smem[threadIdx.y + 1][blockDim.x + 1] = checkIdx(y, x + 1, map.rows, map.cols) ? map(y, x + 1) : 0; 264 smem[0][0] = checkIdx(y - 1, x - 1, map.rows, map.cols) ? map(y - 1, x - 1) : 0; 266 smem[0][blockDim.x + 1] = checkIdx(y - 1, x + 1, map.rows, map.cols) ? map(y - 1, x + 1) : 0; 268 smem[blockDim.y + 1][0] = checkIdx(y + 1, x - 1, map.rows, map.cols) ? map(y + 1, x - 1) : 0; 270 smem[blockDim.y + 1][blockDim.x + 1] = checkIdx(y + 1, x + 1, map.rows, map.cols) ? map(y + 1, x + 1) : (…) [all...] |
clahe.cu | 61 __shared__ int smem[512]; 67 smem[tid] = 0; 76 Emulation::smem::atomicAdd(&smem[data], 1); 82 int tHistVal = smem[tid]; 99 reduce<256>(smem, clipped, tid, plus<int>()); 119 const int lutVal = blockScanInclusive<256>(tHistVal, smem, tid);
|
hist.cu | 74 Emulation::smem::atomicAdd(&shist[(data >> 0) & 0xFFU], 1); 75 Emulation::smem::atomicAdd(&shist[(data >> 8) & 0xFFU], 1); 76 Emulation::smem::atomicAdd(&shist[(data >> 16) & 0xFFU], 1); 77 Emulation::smem::atomicAdd(&shist[(data >> 24) & 0xFFU], 1); 85 Emulation::smem::atomicAdd(&shist[data], 1); 119 Emulation::smem::atomicAdd(shist + ind, 1);
|
hough_lines.cu | 91 int* smem = DynamicSharedMem<int>(); 94 smem[i] = 0; 119 Emulation::smem::atomicAdd(&smem[r + 1], 1); 126 accumRow[i] = smem[i];
|
/external/opencv3/modules/cudawarping/src/cuda/ |
pyr_down.cu | 59 __shared__ work_t smem[256 + 4]; 77 smem[2 + threadIdx.x] = sum; 92 smem[threadIdx.x] = sum; 107 smem[4 + threadIdx.x] = sum; 121 smem[2 + threadIdx.x] = sum; 136 smem[threadIdx.x] = sum; 151 smem[4 + threadIdx.x] = sum; 163 sum = 0.0625f * smem[2 + tid2 - 2]; 164 sum = sum + 0.25f * smem[2 + tid2 - 1]; 165 sum = sum + 0.375f * smem[2 + tid2 ] [all...] |
/external/opencv3/modules/cudafeatures2d/src/cuda/ |
bf_match.cu | 141 extern __shared__ int smem[]; 145 typename Dist::value_type* s_query = (typename Dist::value_type*)(smem); 146 typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * MAX_DESC_LEN); 157 float* s_distance = (float*)(smem); 158 int* s_trainIdx = (int*)(smem + BLOCK_SIZE * BLOCK_SIZE); 190 extern __shared__ int smem[]; 194 typename Dist::value_type* s_query = (typename Dist::value_type*)(smem); 195 typename Dist::value_type* s_train = (typename Dist::value_type*)(smem + BLOCK_SIZE * MAX_DESC_LEN); 214 float* s_distance = (float*)(smem); 215 int* s_trainIdx = (int*)(smem + BLOCK_SIZE * BLOCK_SIZE) [all...] |
/external/opencv3/modules/cudaobjdetect/src/cuda/ |
lbp.cu | 150 Emulation::smem::atomicAdd((rrects + cls * 4 + 0), candidates[tid].x); 151 Emulation::smem::atomicAdd((rrects + cls * 4 + 1), candidates[tid].y); 152 Emulation::smem::atomicAdd((rrects + cls * 4 + 2), candidates[tid].z); 153 Emulation::smem::atomicAdd((rrects + cls * 4 + 3), candidates[tid].w); 159 Emulation::smem::atomicInc((unsigned int*)labels + cls, n); 181 int aidx = Emulation::smem::atomicInc(nclasses, n); 190 int smem = block * ( sizeof(int) + sizeof(int4) ); 191 disjoin<InSameComponint><<<1, block, smem>>>(candidates, objects, ncandidates, groupThreshold, grouping_eps, nclasses);
|
lbp.hpp | 97 Emulation::smem::atomicMin(labels + id, p); 101 Emulation::smem::atomicMin(labels + tid, q);
|