minor bug in hog (unbind texture)
authorAnatoly Baksheev <no@email>
Thu, 31 Mar 2011 11:14:23 +0000 (11:14 +0000)
committerAnatoly Baksheev <no@email>
Thu, 31 Mar 2011 11:14:23 +0000 (11:14 +0000)
modules/gpu/src/cuda/hog.cu

index c358ef3..f42cbbc 100644 (file)
@@ -507,8 +507,7 @@ void extract_descrs_by_cols(int win_height, int win_width, int block_stride_y, i
     dim3 threads(nthreads, 1);\r
     dim3 grid(img_win_width, img_win_height);\r
 \r
-    int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) / \r
-                          block_stride_x;\r
+    int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) / block_stride_x;\r
     extract_descrs_by_cols_kernel<nthreads><<<grid, threads>>>(\r
         img_block_width, win_block_stride_x, win_block_stride_y, block_hists, descriptors);\r
     cudaSafeCall( cudaGetLastError() );\r
@@ -689,7 +688,7 @@ __global__ void compute_gradients_8UC1_kernel(int height, int width, const PtrEl
         hidx = (hidx + cnbins) % cnbins;\r
 \r
         ((uchar2*)qangle.ptr(blockIdx.y))[x] = make_uchar2(hidx, (hidx + 1) % cnbins);\r
-        ((float2*)grad.ptr(blockIdx.y))[x] = make_float2(mag * (1.f - ang), mag * ang);\r
+        ((float2*)  grad.ptr(blockIdx.y))[x] = make_float2(mag * (1.f - ang), mag * ang);\r
     }\r
 }\r
 \r
@@ -725,7 +724,7 @@ __global__ void resize_for_hog_kernel(float sx, float sy, DevMem2D_<uchar> dst,
     unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;\r
 \r
     if (x < dst.cols && y < dst.rows)\r
-        ((unsigned char*)dst.ptr(y))[x] = tex2D(resize8UC1_tex, x * sx + colOfs, y * sy) * 255;\r
+        dst.ptr(y)[x] = tex2D(resize8UC1_tex, x * sx + colOfs, y * sy) * 255;\r
 }\r
 \r
 __global__ void resize_for_hog_kernel(float sx, float sy, DevMem2D_<uchar4> dst, int colOfs)\r
@@ -760,12 +759,14 @@ static void resize_for_hog(const DevMem2D& src, DevMem2D dst, TEX& tex)
 \r
     dim3 threads(32, 8);\r
     dim3 grid(divUp(dst.cols, threads.x), divUp(dst.rows, threads.y));\r
-    float sx = static_cast<float>(src.cols) / dst.cols;\r
+    \r
+       float sx = static_cast<float>(src.cols) / dst.cols;\r
     float sy = static_cast<float>(src.rows) / dst.rows;\r
+\r
     resize_for_hog_kernel<<<grid, threads>>>(sx, sy, (DevMem2D_<T>)dst, colOfs);\r
     cudaSafeCall( cudaGetLastError() );\r
-    cudaSafeCall(cudaThreadSynchronize());\r
-    cudaSafeCall(cudaUnbindTexture(resize8UC1_tex));\r
+    cudaSafeCall( cudaThreadSynchronize() );\r
+    cudaSafeCall( cudaUnbindTexture(tex) );\r
 }\r
 \r
 void resize_8UC1(const DevMem2D& src, DevMem2D dst) { resize_for_hog<uchar> (src, dst, resize8UC1_tex); }\r