{ \r
return mask.ptr(queryIdx)[trainIdx] != 0;\r
}\r
+\r
private:\r
PtrStep mask;\r
};\r
{ \r
return curMask.data == 0 || curMask.ptr(queryIdx)[trainIdx] != 0;\r
}\r
+\r
private:\r
PtrStep* maskCollection;\r
PtrStep curMask;\r
///////////////////////////////////////////////////////////////////////////////\r
// Reduce Sum\r
\r
- template <int BLOCK_DIM_X>\r
- __device__ void reduceSum(float* sdiff, float mySum, int tid)\r
- {\r
- sdiff[tid] = mySum;\r
- __syncthreads();\r
+ template <int BLOCK_DIM_X> __device__ void reduceSum(float* sdiff_row, float& mySum);\r
\r
- if (BLOCK_DIM_X == 512) \r
- {\r
- if (tid < 256) \r
- { \r
- sdiff[tid] = mySum += sdiff[tid + 256]; __syncthreads(); \r
- sdiff[tid] = mySum += sdiff[tid + 128]; __syncthreads();\r
- sdiff[tid] = mySum += sdiff[tid + 64]; __syncthreads();\r
- }\r
- volatile float* smem = sdiff;\r
- smem[tid] = mySum += smem[tid + 32]; \r
- smem[tid] = mySum += smem[tid + 16]; \r
- smem[tid] = mySum += smem[tid + 8]; \r
- smem[tid] = mySum += smem[tid + 4]; \r
- smem[tid] = mySum += smem[tid + 2];\r
- smem[tid] = mySum += smem[tid + 1]; \r
- }\r
- if (BLOCK_DIM_X == 256)\r
- {\r
- if (tid < 128) \r
- { \r
- sdiff[tid] = mySum += sdiff[tid + 128]; __syncthreads(); \r
- sdiff[tid] = mySum += sdiff[tid + 64]; __syncthreads();\r
- }\r
- volatile float* smem = sdiff;\r
- smem[tid] = mySum += smem[tid + 32]; \r
- smem[tid] = mySum += smem[tid + 16]; \r
- smem[tid] = mySum += smem[tid + 8]; \r
- smem[tid] = mySum += smem[tid + 4]; \r
- smem[tid] = mySum += smem[tid + 2];\r
- smem[tid] = mySum += smem[tid + 1];\r
- }\r
- if (BLOCK_DIM_X == 128)\r
- {\r
- if (tid < 64) \r
- { \r
- sdiff[tid] = mySum += sdiff[tid + 64]; __syncthreads(); \r
- }\r
- volatile float* smem = sdiff;\r
- smem[tid] = mySum += smem[tid + 32]; \r
- smem[tid] = mySum += smem[tid + 16]; \r
- smem[tid] = mySum += smem[tid + 8]; \r
- smem[tid] = mySum += smem[tid + 4]; \r
- smem[tid] = mySum += smem[tid + 2];\r
- smem[tid] = mySum += smem[tid + 1];\r
- }\r
- \r
- volatile float* smem = sdiff;\r
- if (BLOCK_DIM_X == 64) \r
- {\r
- if (tid < 32) \r
- {\r
- smem[tid] = mySum += smem[tid + 32]; \r
- smem[tid] = mySum += smem[tid + 16]; \r
- smem[tid] = mySum += smem[tid + 8]; \r
- smem[tid] = mySum += smem[tid + 4]; \r
- smem[tid] = mySum += smem[tid + 2];\r
- smem[tid] = mySum += smem[tid + 1]; \r
- }\r
- }\r
- if (BLOCK_DIM_X == 32) \r
- {\r
- if (tid < 16) \r
- {\r
- smem[tid] = mySum += smem[tid + 16]; \r
- smem[tid] = mySum += smem[tid + 8]; \r
- smem[tid] = mySum += smem[tid + 4]; \r
- smem[tid] = mySum += smem[tid + 2];\r
- smem[tid] = mySum += smem[tid + 1]; \r
- }\r
- }\r
- if (BLOCK_DIM_X == 16) \r
- {\r
- if (tid < 8) \r
- {\r
- smem[tid] = mySum += smem[tid + 8]; \r
- smem[tid] = mySum += smem[tid + 4]; \r
- smem[tid] = mySum += smem[tid + 2];\r
- smem[tid] = mySum += smem[tid + 1]; \r
- }\r
- }\r
- if (BLOCK_DIM_X == 8) \r
- {\r
- if (tid < 4) \r
- {\r
- smem[tid] = mySum += smem[tid + 4]; \r
- smem[tid] = mySum += smem[tid + 2];\r
- smem[tid] = mySum += smem[tid + 1]; \r
- }\r
- }\r
- if (BLOCK_DIM_X == 4) \r
- {\r
- if (tid < 2) \r
- {\r
- smem[tid] = mySum += smem[tid + 2];\r
- smem[tid] = mySum += smem[tid + 1]; \r
- }\r
- }\r
- if (BLOCK_DIM_X == 2) \r
- {\r
- if (tid < 1) \r
- {\r
- smem[tid] = mySum += smem[tid + 1]; \r
- }\r
- }\r
- }\r
-\r
- ///////////////////////////////////////////////////////////////////////////////\r
- // loadDescsVals\r
-\r
- template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int MAX_DESCRIPTORS_LEN, typename T> \r
- __device__ void loadDescsVals(const T* descs, int desc_len, float* smem, float* queryVals)\r
+ template <> __device__ void reduceSum<16>(float* sdiff_row, float& mySum)\r
{\r
- const int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
+ volatile float* smem = sdiff_row;\r
\r
- if (tid < desc_len)\r
+ smem[threadIdx.x] = mySum;\r
+ \r
+ if (threadIdx.x < 8) \r
{\r
- smem[tid] = (float)descs[tid];\r
- }\r
- __syncthreads();\r
-\r
- #pragma unroll\r
- for (int i = threadIdx.x; i < MAX_DESCRIPTORS_LEN; i += BLOCK_DIM_X)\r
- {\r
- *queryVals = smem[i];\r
- ++queryVals;\r
+ smem[threadIdx.x] = mySum += smem[threadIdx.x + 8]; \r
+ smem[threadIdx.x] = mySum += smem[threadIdx.x + 4]; \r
+ smem[threadIdx.x] = mySum += smem[threadIdx.x + 2];\r
+ smem[threadIdx.x] = mySum += smem[threadIdx.x + 1]; \r
}\r
}\r
\r
///////////////////////////////////////////////////////////////////////////////\r
// Distance\r
\r
- template <int BLOCK_DIM_X>\r
class L1Dist\r
{\r
public:\r
- __device__ L1Dist() : mySum(0) {}\r
+ __device__ L1Dist() : mySum(0.0f) {}\r
\r
__device__ void reduceIter(float val1, float val2)\r
{\r
mySum += fabs(val1 - val2);\r
}\r
\r
- __device__ void reduceAll(float* sdiff, int tid)\r
+ template <int BLOCK_DIM_X>\r
+ __device__ void reduceAll(float* sdiff_row)\r
{\r
- reduceSum<BLOCK_DIM_X>(sdiff, mySum, tid);\r
+ reduceSum<BLOCK_DIM_X>(sdiff_row, mySum);\r
}\r
\r
- static __device__ float finalResult(float res)\r
+ __device__ operator float() const\r
{\r
- return res;\r
+ return mySum;\r
}\r
+\r
private:\r
float mySum;\r
};\r
\r
- template <int BLOCK_DIM_X>\r
class L2Dist\r
{\r
public:\r
- __device__ L2Dist() : mySum(0) {}\r
+ __device__ L2Dist() : mySum(0.0f) {}\r
\r
__device__ void reduceIter(float val1, float val2)\r
{\r
mySum += reg * reg;\r
}\r
\r
- __device__ void reduceAll(float* sdiff, int tid)\r
+ template <int BLOCK_DIM_X>\r
+ __device__ void reduceAll(float* sdiff_row)\r
{\r
- reduceSum<BLOCK_DIM_X>(sdiff, mySum, tid);\r
+ reduceSum<BLOCK_DIM_X>(sdiff_row, mySum);\r
}\r
\r
- static __device__ float finalResult(float res)\r
+ __device__ operator float() const\r
{\r
- return sqrtf(res);\r
+ return sqrtf(mySum);\r
}\r
+\r
private:\r
float mySum;\r
};\r
// reduceDescDiff\r
\r
template <int BLOCK_DIM_X, typename Dist, typename T> \r
- __device__ void reduceDescDiff(const T* queryDescs, const T* trainDescs, int desc_len, float* sdiff)\r
+ __device__ void reduceDescDiff(const T* queryDescs, const T* trainDescs, int desc_len, Dist& dist, \r
+ float* sdiff_row)\r
{\r
- const int tid = threadIdx.x;\r
+ for (int i = threadIdx.x; i < desc_len; i += BLOCK_DIM_X)\r
+ dist.reduceIter(queryDescs[i], trainDescs[i]);\r
\r
- Dist dist;\r
+ dist.reduceAll<BLOCK_DIM_X>(sdiff_row);\r
+ }\r
\r
- for (int i = tid; i < desc_len; i += BLOCK_DIM_X)\r
- dist.reduceIter(queryDescs[i], trainDescs[i]);\r
+///////////////////////////////////////////////////////////////////////////////////\r
+////////////////////////////////////// Match //////////////////////////////////////\r
+///////////////////////////////////////////////////////////////////////////////////\r
+ \r
+ ///////////////////////////////////////////////////////////////////////////////\r
+ // loadDescsVals\r
+\r
+ template <int BLOCK_DIM_X, int MAX_DESCRIPTORS_LEN, typename T> \r
+ __device__ void loadDescsVals(const T* descs, int desc_len, float* queryVals, float* smem)\r
+ {\r
+ const int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
+\r
+ if (tid < desc_len)\r
+ {\r
+ smem[tid] = (float)descs[tid];\r
+ }\r
+ __syncthreads();\r
\r
- dist.reduceAll(sdiff, tid);\r
+ #pragma unroll\r
+ for (int i = threadIdx.x; i < MAX_DESCRIPTORS_LEN; i += BLOCK_DIM_X)\r
+ {\r
+ *queryVals = smem[i];\r
+ ++queryVals;\r
+ }\r
}\r
\r
///////////////////////////////////////////////////////////////////////////////\r
- // reduceDescDiff_smem\r
+ // reduceDescDiffCached\r
\r
template <int N> struct UnrollDescDiff\r
{\r
template <typename Dist, typename T>\r
- static __device__ void calcCheck(Dist& dist, const float* queryVals, const T* trainDescs, \r
- int ind, int desc_len)\r
+ static __device__ void calcCheck(const float* queryVals, const T* trainDescs, int desc_len, \r
+ Dist& dist, int ind)\r
{\r
if (ind < desc_len)\r
+ {\r
dist.reduceIter(*queryVals, trainDescs[ind]);\r
\r
- ++queryVals;\r
+ ++queryVals;\r
\r
- UnrollDescDiff<N - 1>::calcCheck(dist, queryVals, trainDescs, ind + blockDim.x, desc_len);\r
+ UnrollDescDiff<N - 1>::calcCheck(queryVals, trainDescs, desc_len, dist, ind + blockDim.x);\r
+ }\r
}\r
\r
template <typename Dist, typename T>\r
- static __device__ void calcWithoutCheck(Dist& dist, const float* queryVals, const T* trainDescs)\r
+ static __device__ void calcWithoutCheck(const float* queryVals, const T* trainDescs, Dist& dist)\r
{\r
dist.reduceIter(*queryVals, *trainDescs);\r
\r
++queryVals;\r
trainDescs += blockDim.x;\r
\r
- UnrollDescDiff<N - 1>::calcWithoutCheck(dist, queryVals, trainDescs);\r
+ UnrollDescDiff<N - 1>::calcWithoutCheck(queryVals, trainDescs, dist);\r
}\r
};\r
template <> struct UnrollDescDiff<0>\r
{\r
template <typename Dist, typename T>\r
- static __device__ void calcCheck(Dist& dist, const float* queryVals, const T* trainDescs, \r
- int ind, int desc_len)\r
+ static __device__ void calcCheck(const float* queryVals, const T* trainDescs, int desc_len, \r
+ Dist& dist, int ind)\r
{\r
}\r
\r
template <typename Dist, typename T>\r
- static __device__ void calcWithoutCheck(Dist& dist, const float* queryVals, const T* trainDescs)\r
+ static __device__ void calcWithoutCheck(const float* queryVals, const T* trainDescs, Dist& dist)\r
{\r
}\r
};\r
struct DescDiffCalculator<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN, false>\r
{\r
template <typename Dist, typename T>\r
- static __device__ void calc(Dist& dist, const float* queryVals, const T* trainDescs, int desc_len)\r
+ static __device__ void calc(const float* queryVals, const T* trainDescs, int desc_len, Dist& dist)\r
{\r
- UnrollDescDiff<MAX_DESCRIPTORS_LEN / BLOCK_DIM_X>::calcCheck(dist, queryVals, trainDescs, \r
- threadIdx.x, desc_len);\r
+ UnrollDescDiff<MAX_DESCRIPTORS_LEN / BLOCK_DIM_X>::calcCheck(queryVals, trainDescs, desc_len, \r
+ dist, threadIdx.x);\r
}\r
};\r
template <int BLOCK_DIM_X, int MAX_DESCRIPTORS_LEN> \r
struct DescDiffCalculator<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN, true>\r
{\r
template <typename Dist, typename T>\r
- static __device__ void calc(Dist& dist, const float* queryVals, const T* trainDescs, int desc_len)\r
+ static __device__ void calc(const float* queryVals, const T* trainDescs, int desc_len, Dist& dist)\r
{\r
- UnrollDescDiff<MAX_DESCRIPTORS_LEN / BLOCK_DIM_X>::calcWithoutCheck(dist, queryVals, \r
- trainDescs + threadIdx.x);\r
+ UnrollDescDiff<MAX_DESCRIPTORS_LEN / BLOCK_DIM_X>::calcWithoutCheck(queryVals, \r
+ trainDescs + threadIdx.x, dist);\r
}\r
};\r
\r
template <int BLOCK_DIM_X, int MAX_DESCRIPTORS_LEN, bool DESC_LEN_EQ_MAX_LEN, typename Dist, typename T>\r
- __device__ void reduceDescDiff_smem(const float* queryVals, const T* trainDescs, int desc_len, float* sdiff)\r
- {\r
- const int tid = threadIdx.x;\r
+ __device__ void reduceDescDiffCached(const float* queryVals, const T* trainDescs, int desc_len, Dist& dist, \r
+ float* sdiff_row)\r
+ { \r
+ DescDiffCalculator<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN, DESC_LEN_EQ_MAX_LEN>::calc(queryVals, \r
+ trainDescs, desc_len, dist);\r
\r
- Dist dist;\r
-\r
- DescDiffCalculator<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN, DESC_LEN_EQ_MAX_LEN>::calc(dist, queryVals, \r
- trainDescs, desc_len);\r
- \r
- dist.reduceAll(sdiff, tid);\r
+ dist.reduceAll<BLOCK_DIM_X>(sdiff_row);\r
}\r
\r
-///////////////////////////////////////////////////////////////////////////////////\r
-////////////////////////////////////// Match //////////////////////////////////////\r
-///////////////////////////////////////////////////////////////////////////////////\r
-\r
///////////////////////////////////////////////////////////////////////////////\r
- // warpReduceMin\r
+ // warpReduceMinIdxIdx\r
\r
template <int BLOCK_DIM_Y> \r
- __device__ void warpReduceMin(int tid, volatile float* sdata, volatile int* strainIdx, volatile int* simgIdx)\r
+ __device__ void warpReduceMinIdxIdx(float& myMin, int& myBestTrainIdx, int& myBestImgIdx, \r
+ volatile float* sdata, volatile int* strainIdx, volatile int* simgIdx);\r
+\r
+ template <> \r
+ __device__ void warpReduceMinIdxIdx<16>(float& myMin, int& myBestTrainIdx, int& myBestImgIdx, \r
+ volatile float* smin, volatile int* strainIdx, volatile int* simgIdx)\r
{\r
- float minSum = sdata[tid];\r
+ const int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
\r
- if (BLOCK_DIM_Y >= 64) \r
- {\r
- float reg = sdata[tid + 32];\r
- if (reg < minSum)\r
- {\r
- sdata[tid] = minSum = reg;\r
- strainIdx[tid] = strainIdx[tid + 32];\r
- simgIdx[tid] = simgIdx[tid + 32];\r
- }\r
- }\r
- if (BLOCK_DIM_Y >= 32) \r
+ if (tid < 8)\r
{\r
- float reg = sdata[tid + 16];\r
- if (reg < minSum)\r
- {\r
- sdata[tid] = minSum = reg;\r
- strainIdx[tid] = strainIdx[tid + 16];\r
- simgIdx[tid] = simgIdx[tid + 16];\r
- }\r
- }\r
- if (BLOCK_DIM_Y >= 16) \r
- {\r
- float reg = sdata[tid + 8];\r
- if (reg < minSum)\r
+ myMin = smin[tid];\r
+ myBestTrainIdx = strainIdx[tid];\r
+ myBestImgIdx = simgIdx[tid];\r
+\r
+ float reg = smin[tid + 8];\r
+ if (reg < myMin)\r
{\r
- sdata[tid] = minSum = reg;\r
- strainIdx[tid] = strainIdx[tid + 8];\r
- simgIdx[tid] = simgIdx[tid + 8];\r
+ smin[tid] = myMin = reg;\r
+ strainIdx[tid] = myBestTrainIdx = strainIdx[tid + 8];\r
+ simgIdx[tid] = myBestImgIdx = simgIdx[tid + 8];\r
}\r
- }\r
- if (BLOCK_DIM_Y >= 8) \r
- { \r
- float reg = sdata[tid + 4];\r
- if (reg < minSum)\r
+\r
+ reg = smin[tid + 4];\r
+ if (reg < myMin)\r
{\r
- sdata[tid] = minSum = reg;\r
- strainIdx[tid] = strainIdx[tid + 4];\r
- simgIdx[tid] = simgIdx[tid + 4];\r
+ smin[tid] = myMin = reg;\r
+ strainIdx[tid] = myBestTrainIdx = strainIdx[tid + 4];\r
+ simgIdx[tid] = myBestImgIdx = simgIdx[tid + 4];\r
}\r
- }\r
- if (BLOCK_DIM_Y >= 4) \r
- { \r
- float reg = sdata[tid + 2];\r
- if (reg < minSum)\r
+ \r
+ reg = smin[tid + 2];\r
+ if (reg < myMin)\r
{\r
- sdata[tid] = minSum = reg;\r
- strainIdx[tid] = strainIdx[tid + 2];\r
- simgIdx[tid] = simgIdx[tid + 2];\r
+ smin[tid] = myMin = reg;\r
+ strainIdx[tid] = myBestTrainIdx = strainIdx[tid + 2];\r
+ simgIdx[tid] = myBestImgIdx = simgIdx[tid + 2];\r
}\r
- }\r
- if (BLOCK_DIM_Y >= 2) \r
- {\r
- float reg = sdata[tid + 1];\r
- if (reg < minSum)\r
+ \r
+ reg = smin[tid + 1];\r
+ if (reg < myMin)\r
{\r
- sdata[tid] = minSum = reg;\r
- strainIdx[tid] = strainIdx[tid + 1];\r
- simgIdx[tid] = simgIdx[tid + 1];\r
+ smin[tid] = myMin = reg;\r
+ strainIdx[tid] = myBestTrainIdx = strainIdx[tid + 1];\r
+ simgIdx[tid] = myBestImgIdx = simgIdx[tid + 1];\r
}\r
}\r
}\r
///////////////////////////////////////////////////////////////////////////////\r
// findBestMatch\r
\r
- template <int BLOCK_DIM_Y, typename Dist>\r
- __device__ void findBestMatch(int queryIdx, float myMin, int myBestTrainIdx, int myBestImgIdx, \r
- float* smin, int* strainIdx, int* simgIdx, int* trainIdx, int* imgIdx, float* distance)\r
+ template <int BLOCK_DIM_Y>\r
+ __device__ void findBestMatch(float& myMin, int& myBestTrainIdx, int& myBestImgIdx, \r
+ float* smin, int* strainIdx, int* simgIdx)\r
{\r
if (threadIdx.x == 0)\r
{\r
}\r
__syncthreads();\r
\r
- const int tid = threadIdx.y * blockDim.x + threadIdx.x;\r
-\r
- if (tid < 32)\r
- warpReduceMin<BLOCK_DIM_Y>(tid, smin, strainIdx, simgIdx);\r
-\r
- if (threadIdx.x == 0 && threadIdx.y == 0)\r
- {\r
- float minSum = smin[0];\r
- int bestTrainIdx = strainIdx[0];\r
- int bestImgIdx = simgIdx[0];\r
-\r
- imgIdx[queryIdx] = bestImgIdx;\r
- trainIdx[queryIdx] = bestTrainIdx;\r
- distance[queryIdx] = Dist::finalResult(minSum);\r
- }\r
+ warpReduceMinIdxIdx<BLOCK_DIM_Y>(myMin, myBestTrainIdx, myBestImgIdx, smin, strainIdx, simgIdx);\r
}\r
\r
///////////////////////////////////////////////////////////////////////////////\r
// ReduceDescCalculator\r
\r
- template <int BLOCK_DIM_X, typename Dist, typename T>\r
+ template <int BLOCK_DIM_X, typename T>\r
class ReduceDescCalculatorSimple\r
{\r
public:\r
queryDescs = queryDescs_;\r
}\r
\r
- __device__ void calc(const T* trainDescs, int desc_len, float* sdiff_row) const\r
+ template <typename Dist>\r
+ __device__ void calc(const T* trainDescs, int desc_len, Dist& dist, float* sdiff_row) const\r
{\r
- reduceDescDiff<BLOCK_DIM_X, Dist>(queryDescs, trainDescs, desc_len, sdiff_row);\r
+ reduceDescDiff<BLOCK_DIM_X>(queryDescs, trainDescs, desc_len, dist, sdiff_row);\r
}\r
\r
private:\r
const T* queryDescs;\r
};\r
\r
- template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int MAX_DESCRIPTORS_LEN, bool DESC_LEN_EQ_MAX_LEN, \r
- typename Dist, typename T>\r
- class ReduceDescCalculatorSmem\r
+ template <int BLOCK_DIM_X, int MAX_DESCRIPTORS_LEN, bool DESC_LEN_EQ_MAX_LEN, typename T>\r
+ class ReduceDescCalculatorCached\r
{\r
public:\r
__device__ void prepare(const T* queryDescs, int desc_len, float* smem)\r
{\r
- loadDescsVals<BLOCK_DIM_X, BLOCK_DIM_Y, MAX_DESCRIPTORS_LEN>(queryDescs, desc_len, smem, queryVals);\r
+ loadDescsVals<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN>(queryDescs, desc_len, queryVals, smem);\r
}\r
\r
- __device__ void calc(const T* trainDescs, int desc_len, float* sdiff_row) const\r
+ template <typename Dist>\r
+ __device__ void calc(const T* trainDescs, int desc_len, Dist& dist, float* sdiff_row) const\r
{\r
- reduceDescDiff_smem<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN, DESC_LEN_EQ_MAX_LEN, Dist>(queryVals, trainDescs, \r
- desc_len, sdiff_row);\r
+ reduceDescDiffCached<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN, DESC_LEN_EQ_MAX_LEN>(queryVals, trainDescs, \r
+ desc_len, dist, sdiff_row);\r
}\r
\r
private:\r
///////////////////////////////////////////////////////////////////////////////\r
// matchDescs loop\r
\r
- template <typename ReduceDescCalculator, typename T, typename Mask>\r
- __device__ void matchDescs(int queryIdx, const int imgIdx, const DevMem2D_<T>& trainDescs_, \r
+ template <typename Dist, typename ReduceDescCalculator, typename T, typename Mask>\r
+ __device__ void matchDescs(int queryIdx, int imgIdx, const DevMem2D_<T>& trainDescs_, \r
const Mask& m, const ReduceDescCalculator& reduceDescCalc,\r
- float* sdiff_row, float& myMin, int& myBestTrainIdx, int& myBestImgIdx)\r
+ float& myMin, int& myBestTrainIdx, int& myBestImgIdx, float* sdiff_row)\r
{\r
- const T* trainDescs = trainDescs_.ptr(threadIdx.y);\r
- const int trainDescsStep = blockDim.y * trainDescs_.step / sizeof(T);\r
- for (int trainIdx = threadIdx.y; trainIdx < trainDescs_.rows; \r
- trainIdx += blockDim.y, trainDescs += trainDescsStep)\r
+ for (int trainIdx = threadIdx.y; trainIdx < trainDescs_.rows; trainIdx += blockDim.y)\r
{\r
if (m(queryIdx, trainIdx))\r
{\r
- reduceDescCalc.calc(trainDescs, trainDescs_.cols, sdiff_row);\r
+ const T* trainDescs = trainDescs_.ptr(trainIdx);\r
+\r
+ Dist dist;\r
+\r
+ reduceDescCalc.calc(trainDescs, trainDescs_.cols, dist, sdiff_row);\r
\r
if (threadIdx.x == 0)\r
{\r
- float reg = sdiff_row[0];\r
- if (reg < myMin)\r
+ if (dist < myMin)\r
{\r
- myMin = reg;\r
+ myMin = dist;\r
myBestTrainIdx = trainIdx;\r
myBestImgIdx = imgIdx;\r
}\r
{\r
}\r
\r
- template <typename ReduceDescCalculator, typename Mask>\r
+ template <typename Dist, typename ReduceDescCalculator, typename Mask>\r
__device__ void loop(int queryIdx, Mask& m, const ReduceDescCalculator& reduceDescCalc, \r
- float* sdiff_row, float& myMin, int& myBestTrainIdx, int& myBestImgIdx) const\r
+ float& myMin, int& myBestTrainIdx, int& myBestImgIdx, float* sdiff_row) const\r
{\r
- matchDescs(queryIdx, 0, trainDescs, m, reduceDescCalc, \r
- sdiff_row, myMin, myBestTrainIdx, myBestImgIdx);\r
+ matchDescs<Dist>(queryIdx, 0, trainDescs, m, reduceDescCalc, \r
+ myMin, myBestTrainIdx, myBestImgIdx, sdiff_row);\r
}\r
\r
__device__ int desc_len() const\r
{\r
return trainDescs.cols;\r
}\r
+\r
private:\r
DevMem2D_<T> trainDescs;\r
};\r
{\r
}\r
\r
- template <typename ReduceDescCalculator, typename Mask>\r
+ template <typename Dist, typename ReduceDescCalculator, typename Mask>\r
__device__ void loop(int queryIdx, Mask& m, const ReduceDescCalculator& reduceDescCalc, \r
- float* sdiff_row, float& myMin, int& myBestTrainIdx, int& myBestImgIdx) const\r
+ float& myMin, int& myBestTrainIdx, int& myBestImgIdx, float* sdiff_row) const\r
{\r
for (int imgIdx = 0; imgIdx < nImg; ++imgIdx)\r
{\r
DevMem2D_<T> trainDescs = trainCollection[imgIdx];\r
m.nextMask();\r
- matchDescs(queryIdx, imgIdx, trainDescs, m, reduceDescCalc, \r
- sdiff_row, myMin, myBestTrainIdx, myBestImgIdx);\r
+ matchDescs<Dist>(queryIdx, imgIdx, trainDescs, m, reduceDescCalc, \r
+ myMin, myBestTrainIdx, myBestImgIdx, sdiff_row);\r
}\r
}\r
\r
{\r
return desclen;\r
}\r
+\r
private:\r
const DevMem2D_<T>* trainCollection;\r
int nImg;\r
\r
template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename ReduceDescCalculator, typename Dist, typename T, \r
typename Train, typename Mask>\r
- __global__ void match(PtrStep_<T> queryDescs_, Train train, Mask mask, int* trainIdx, int* imgIdx, float* distance)\r
+ __global__ void match(const PtrStep_<T> queryDescs_, const Train train, const Mask mask, \r
+ int* trainIdx, int* imgIdx, float* distance)\r
{\r
- __shared__ float sdiff[BLOCK_DIM_X * BLOCK_DIM_Y];\r
- __shared__ float smin[64];\r
- __shared__ int strainIdx[64];\r
- __shared__ int simgIdx[64];\r
+ __shared__ float smem[BLOCK_DIM_X * BLOCK_DIM_Y]; \r
\r
const int queryIdx = blockIdx.x;\r
\r
float myMin = numeric_limits_gpu<float>::max();\r
\r
{\r
- float* sdiff_row = sdiff + BLOCK_DIM_X * threadIdx.y;\r
+ float* sdiff_row = smem + BLOCK_DIM_X * threadIdx.y;\r
+\r
Mask m = mask;\r
+\r
ReduceDescCalculator reduceDescCalc;\r
- reduceDescCalc.prepare(queryDescs_.ptr(queryIdx), train.desc_len(), sdiff);\r
+\r
+ reduceDescCalc.prepare(queryDescs_.ptr(queryIdx), train.desc_len(), smem);\r
\r
- train.loop(queryIdx, m, reduceDescCalc, sdiff_row, myMin, myBestTrainIdx, myBestImgIdx);\r
+ train.template loop<Dist>(queryIdx, m, reduceDescCalc, myMin, myBestTrainIdx, myBestImgIdx, sdiff_row);\r
}\r
+ __syncthreads();\r
+\r
+ float* smin = smem;\r
+ int* strainIdx = (int*)(smin + BLOCK_DIM_Y);\r
+ int* simgIdx = strainIdx + BLOCK_DIM_Y;\r
\r
- findBestMatch<BLOCK_DIM_Y, Dist>(queryIdx, myMin, myBestTrainIdx, myBestImgIdx, \r
- smin, strainIdx, simgIdx, trainIdx, imgIdx, distance);\r
+ findBestMatch<BLOCK_DIM_Y>(myMin, myBestTrainIdx, myBestImgIdx, \r
+ smin, strainIdx, simgIdx);\r
+\r
+ if (threadIdx.x == 0 && threadIdx.y == 0)\r
+ {\r
+ imgIdx[queryIdx] = myBestImgIdx;\r
+ trainIdx[queryIdx] = myBestTrainIdx;\r
+ distance[queryIdx] = myMin;\r
+ }\r
}\r
\r
///////////////////////////////////////////////////////////////////////////////\r
// Match kernel callers\r
\r
- template <int BLOCK_DIM_X, int BLOCK_DIM_Y, template <int> class Dist, typename T, \r
+ template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename Dist, typename T, \r
typename Train, typename Mask>\r
- void match_caller(const DevMem2D_<T>& queryDescs, const Train& train, \r
+ void matchSimple_caller(const DevMem2D_<T>& queryDescs, const Train& train, \r
const Mask& mask, const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance)\r
{\r
StaticAssert<BLOCK_DIM_Y <= 64>::check(); // blockDimY vals must reduce by warp\r
dim3 grid(queryDescs.rows, 1, 1);\r
dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);\r
\r
- match<BLOCK_DIM_X, BLOCK_DIM_Y, ReduceDescCalculatorSimple<BLOCK_DIM_X, Dist<BLOCK_DIM_X>, T>, \r
- Dist<BLOCK_DIM_X>, T><<<grid, threads>>>(queryDescs, train, mask, trainIdx.data, \r
+ match<BLOCK_DIM_X, BLOCK_DIM_Y, ReduceDescCalculatorSimple<BLOCK_DIM_X, T>, Dist, T>\r
+ <<<grid, threads>>>(queryDescs, train, mask, trainIdx.data, \r
imgIdx.data, distance.data);\r
\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
template <int BLOCK_DIM_X, int BLOCK_DIM_Y, int MAX_DESCRIPTORS_LEN, bool DESC_LEN_EQ_MAX_LEN, \r
- template <int> class Dist, typename T, typename Train, typename Mask>\r
- void match_smem_caller(const DevMem2D_<T>& queryDescs, const Train& train, \r
+ typename Dist, typename T, typename Train, typename Mask>\r
+ void matchCached_caller(const DevMem2D_<T>& queryDescs, const Train& train, \r
const Mask& mask, const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance)\r
{\r
StaticAssert<BLOCK_DIM_Y <= 64>::check(); // blockDimY vals must reduce by warp\r
dim3 grid(queryDescs.rows, 1, 1);\r
dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);\r
\r
- match<BLOCK_DIM_X, BLOCK_DIM_Y, ReduceDescCalculatorSmem<BLOCK_DIM_X, BLOCK_DIM_Y, \r
- MAX_DESCRIPTORS_LEN, DESC_LEN_EQ_MAX_LEN, Dist<BLOCK_DIM_X>, T>, \r
- Dist<BLOCK_DIM_X>, T><<<grid, threads>>>(queryDescs, train, mask, trainIdx.data, \r
+ match<BLOCK_DIM_X, BLOCK_DIM_Y, \r
+ ReduceDescCalculatorCached<BLOCK_DIM_X, MAX_DESCRIPTORS_LEN, DESC_LEN_EQ_MAX_LEN, T>, \r
+ Dist, T>\r
+ <<<grid, threads>>>(queryDescs, train, mask, trainIdx.data, \r
imgIdx.data, distance.data);\r
\r
cudaSafeCall( cudaThreadSynchronize() );\r
///////////////////////////////////////////////////////////////////////////////\r
// Match kernel chooser\r
\r
- template <template <int> class Dist, typename T, typename Train, typename Mask>\r
+ template <typename Dist, typename T, typename Train, typename Mask>\r
void match_chooser(const DevMem2D_<T>& queryDescs, const Train& train, \r
const Mask& mask, const DevMem2Di& trainIdx, const DevMem2Di& imgIdx, const DevMem2Df& distance)\r
{\r
if (queryDescs.cols < 64)\r
- match_smem_caller<16, 16, 64, false, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
+ matchCached_caller<16, 16, 64, false, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
else if (queryDescs.cols == 64)\r
- match_smem_caller<16, 16, 64, true, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
+ matchCached_caller<16, 16, 64, true, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
else if (queryDescs.cols < 128)\r
- match_smem_caller<16, 16, 128, false, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
+ matchCached_caller<16, 16, 128, false, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
else if (queryDescs.cols == 128)\r
- match_smem_caller<16, 16, 128, true, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
+ matchCached_caller<16, 16, 128, true, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
else if (queryDescs.cols < 256)\r
- match_smem_caller<16, 16, 256, false, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
+ matchCached_caller<16, 16, 256, false, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
else if (queryDescs.cols == 256)\r
- match_smem_caller<16, 16, 256, true, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
+ matchCached_caller<16, 16, 256, true, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
else\r
- match_caller<16, 16, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
+ matchSimple_caller<16, 16, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance);\r
\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
{\r
const T* trainDescs = trainDescs_.ptr(trainIdx);\r
\r
- float dist = numeric_limits_gpu<float>::max();\r
+ float myDist = numeric_limits_gpu<float>::max();\r
\r
if (mask(queryIdx, trainIdx))\r
{\r
- reduceDescDiff<BLOCK_DIM_X, Dist>(queryDescs, trainDescs, trainDescs_.cols, sdiff_row);\r
+ Dist dist;\r
+\r
+ reduceDescDiff<BLOCK_DIM_X>(queryDescs, trainDescs, trainDescs_.cols, dist, sdiff_row);\r
\r
if (threadIdx.x == 0)\r
- {\r
- dist = Dist::finalResult(sdiff_row[0]);\r
- }\r
+ myDist = dist;\r
}\r
\r
if (threadIdx.x == 0)\r
- distance.ptr(queryIdx)[trainIdx] = dist;\r
+ distance.ptr(queryIdx)[trainIdx] = myDist;\r
}\r
}\r
\r
///////////////////////////////////////////////////////////////////////////////\r
// Calc distance kernel caller\r
\r
- template <int BLOCK_DIM_X, int BLOCK_DIM_Y, template <int> class Dist, typename T, typename Mask>\r
+ template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename Dist, typename T, typename Mask>\r
void calcDistance_caller(const DevMem2D_<T>& queryDescs, const DevMem2D_<T>& trainDescs, \r
const Mask& mask, const DevMem2Df& distance)\r
{\r
dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);\r
dim3 grid(queryDescs.rows, divUp(trainDescs.rows, BLOCK_DIM_Y), 1);\r
\r
- calcDistance<BLOCK_DIM_X, BLOCK_DIM_Y, Dist<BLOCK_DIM_X>, T><<<grid, threads>>>(\r
+ calcDistance<BLOCK_DIM_X, BLOCK_DIM_Y, Dist, T><<<grid, threads>>>(\r
queryDescs, trainDescs, mask, distance);\r
\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
\r
///////////////////////////////////////////////////////////////////////////////\r
- // reduceMin\r
+ // warpReduceMinIdx\r
\r
template <int BLOCK_SIZE> \r
__device__ void warpReduceMinIdx(volatile float* sdist, volatile int* strainIdx, float& myMin, int tid)\r
{\r
#if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 110\r
\r
- __shared__ float sdiff[BLOCK_DIM_X * BLOCK_DIM_Y];\r
+ __shared__ float smem[BLOCK_DIM_X * BLOCK_DIM_Y];\r
\r
- float* sdiff_row = sdiff + BLOCK_DIM_X * threadIdx.y;\r
+ float* sdiff_row = smem + BLOCK_DIM_X * threadIdx.y;\r
\r
const int queryIdx = blockIdx.x;\r
const T* queryDescs = queryDescs_.ptr(queryIdx);\r
\r
const int trainIdx = blockIdx.y * BLOCK_DIM_Y + threadIdx.y;\r
+\r
if (trainIdx < trainDescs_.rows)\r
{\r
const T* trainDescs = trainDescs_.ptr(trainIdx);\r
\r
if (mask(queryIdx, trainIdx))\r
{\r
- reduceDescDiff<BLOCK_DIM_X, Dist>(queryDescs, trainDescs, trainDescs_.cols, sdiff_row);\r
+ Dist dist;\r
+\r
+ reduceDescDiff<BLOCK_DIM_X>(queryDescs, trainDescs, trainDescs_.cols, dist, sdiff_row);\r
\r
if (threadIdx.x == 0)\r
{\r
- float dist = Dist::finalResult(sdiff_row[0]);\r
if (dist < maxDistance)\r
{\r
unsigned int i = atomicInc(nMatches + queryIdx, (unsigned int) -1);\r
///////////////////////////////////////////////////////////////////////////////\r
// Radius Match kernel caller\r
\r
- template <int BLOCK_DIM_X, int BLOCK_DIM_Y, template <int> class Dist, typename T, typename Mask>\r
+ template <int BLOCK_DIM_X, int BLOCK_DIM_Y, typename Dist, typename T, typename Mask>\r
void radiusMatch_caller(const DevMem2D_<T>& queryDescs, const DevMem2D_<T>& trainDescs, \r
float maxDistance, const Mask& mask, const DevMem2Di& trainIdx, unsigned int* nMatches, \r
const DevMem2Df& distance)\r
dim3 threads(BLOCK_DIM_X, BLOCK_DIM_Y, 1);\r
dim3 grid(queryDescs.rows, divUp(trainDescs.rows, BLOCK_DIM_Y), 1);\r
\r
- radiusMatch<BLOCK_DIM_X, BLOCK_DIM_Y, Dist<BLOCK_DIM_X>, T><<<grid, threads>>>(\r
+ radiusMatch<BLOCK_DIM_X, BLOCK_DIM_Y, Dist, T><<<grid, threads>>>(\r
queryDescs, trainDescs, maxDistance, mask, trainIdx, nMatches, distance);\r
\r
cudaSafeCall( cudaThreadSynchronize() );\r