// since grids are 2D, need to compute xBlock and yBlock indices\r
const int xBlock = (blockIdx.y & 3); // blockIdx.y % 4\r
const int yBlock = (blockIdx.y >> 2); // floor(blockIdx.y/4)\r
- const int xIndex = xBlock * blockDim.x + threadIdx.x;\r
- const int yIndex = yBlock * blockDim.y + threadIdx.y;\r
+ const int xIndex = xBlock * 5 + threadIdx.x;\r
+ const int yIndex = yBlock * 5 + threadIdx.y;\r
\r
s_PATCH[threadIdx.y][threadIdx.x] = calcPATCH(yIndex, xIndex, centerX, centerY, win_offset, cos_dir, sin_dir, win_size);\r
- if (threadIdx.x == 0)\r
- s_PATCH[threadIdx.y][5] = calcPATCH(yIndex, xBlock * blockDim.x + 5, centerX, centerY, win_offset, cos_dir, sin_dir, win_size);\r
- if (threadIdx.y == 0)\r
- s_PATCH[5][threadIdx.x] = calcPATCH(yBlock * blockDim.y + 5, xIndex, centerX, centerY, win_offset, cos_dir, sin_dir, win_size);\r
- if (threadIdx.x == 0 && threadIdx.y == 0)\r
- s_PATCH[5][5] = calcPATCH(yBlock * blockDim.y + 5, xBlock * blockDim.x + 5, centerX, centerY, win_offset, cos_dir, sin_dir, win_size);\r
__syncthreads();\r
\r
- const float dw = c_DW[yIndex * PATCH_SZ + xIndex];\r
+ if (threadIdx.x < 5 && threadIdx.y < 5)\r
+ {\r
+ tid = threadIdx.y * 5 + threadIdx.x;\r
+\r
+ const float dw = c_DW[yIndex * PATCH_SZ + xIndex];\r
\r
- const float vx = (s_PATCH[threadIdx.y ][threadIdx.x + 1] - s_PATCH[threadIdx.y][threadIdx.x] + s_PATCH[threadIdx.y + 1][threadIdx.x + 1] - s_PATCH[threadIdx.y + 1][threadIdx.x ]) * dw;\r
- const float vy = (s_PATCH[threadIdx.y + 1][threadIdx.x ] - s_PATCH[threadIdx.y][threadIdx.x] + s_PATCH[threadIdx.y + 1][threadIdx.x + 1] - s_PATCH[threadIdx.y ][threadIdx.x + 1]) * dw;\r
+ const float vx = (s_PATCH[threadIdx.y ][threadIdx.x + 1] - s_PATCH[threadIdx.y][threadIdx.x] + s_PATCH[threadIdx.y + 1][threadIdx.x + 1] - s_PATCH[threadIdx.y + 1][threadIdx.x ]) * dw;\r
+ const float vy = (s_PATCH[threadIdx.y + 1][threadIdx.x ] - s_PATCH[threadIdx.y][threadIdx.x] + s_PATCH[threadIdx.y + 1][threadIdx.x + 1] - s_PATCH[threadIdx.y ][threadIdx.x + 1]) * dw;\r
\r
- s_dx_bin[tid] = vx;\r
- s_dy_bin[tid] = vy;\r
+ s_dx_bin[tid] = vx;\r
+ s_dy_bin[tid] = vy;\r
+ }\r
}\r
\r
__device__ void reduce_sum25(volatile float* sdata1, volatile float* sdata2, volatile float* sdata3, volatile float* sdata4, int tid)\r
\r
if (descriptors.cols == 64)\r
{\r
- compute_descriptors64<<<dim3(nFeatures, 16, 1), dim3(5, 5, 1)>>>(descriptors, featureX, featureY, featureSize, featureDir);\r
+ compute_descriptors64<<<dim3(nFeatures, 16, 1), dim3(6, 6, 1)>>>(descriptors, featureX, featureY, featureSize, featureDir);\r
cudaSafeCall( cudaGetLastError() );\r
\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
else\r
{\r
- compute_descriptors128<<<dim3(nFeatures, 16, 1), dim3(5, 5, 1)>>>(descriptors, featureX, featureY, featureSize, featureDir); \r
+ compute_descriptors128<<<dim3(nFeatures, 16, 1), dim3(6, 6, 1)>>>(descriptors, featureX, featureY, featureSize, featureDir); \r
cudaSafeCall( cudaGetLastError() );\r
\r
cudaSafeCall( cudaThreadSynchronize() );\r