     int img_block_height = (height - CELLS_PER_BLOCK_Y * CELL_HEIGHT + block_stride_y) / 
                            block_stride_y;
 
-    dim3 grid(div_up(img_block_width, nblocks), img_block_height);
+    dim3 grid(divUp(img_block_width, nblocks), img_block_height);
     dim3 threads(32, 2, nblocks);
 
     cudaSafeCall(cudaFuncSetCacheConfig(compute_hists_kernel_many_blocks<nblocks>, 

    }

    if (threadIdx.x == 0)
        labels[blockIdx.y * img_win_width + blockIdx.x * blockDim.z + win_x] 
            = (product + free_coef >= threshold);

}



// Resize

texture<uchar4, 2, cudaReadModeNormalizedFloat> resize8UC4_tex;
texture<uchar,  2, cudaReadModeNormalizedFloat> resize8UC1_tex;

__global__ void resize_for_hog_kernel(float sx, float sy, DevMem2D_<uchar> dst, int colOfs)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < dst.cols && y < dst.rows)
        ((unsigned char*)dst.ptr(y))[x] = tex2D(resize8UC1_tex, x * sx + colOfs, y * sy) * 255;
}

__global__ void resize_for_hog_kernel(float sx, float sy, DevMem2D_<uchar4> dst, int colOfs)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < dst.cols && y < dst.rows)
	{        
		float4 val = tex2D(resize8UC4_tex, x * sx + colOfs, y * sy);
        dst.ptr(y)[x] = make_uchar4(val.x * 255, val.y * 255, val.z * 255, val.w * 255);
	}
}

template<class T, class TEX> 
static void resize_for_hog(const DevMem2D& src, DevMem2D dst, TEX& tex)
{
    tex.filterMode = cudaFilterModeLinear;

    size_t texOfs = 0;
    int colOfs = 0;

    cudaChannelFormatDesc desc = cudaCreateChannelDesc<T>();    
    cudaSafeCall( cudaBindTexture2D(&texOfs, tex, src.data, desc, src.cols, src.rows, src.step) );

    if (texOfs != 0) 
    {
        colOfs = static_cast<int>( texOfs/sizeof(T) );
        cudaSafeCall( cudaUnbindTexture(tex) );
        cudaSafeCall( cudaBindTexture2D(&texOfs, tex, src.data, desc, src.cols, src.rows, src.step) );
    }    

    dim3 threads(32, 8);
    dim3 grid(divUp(dst.cols, threads.x), divUp(dst.rows, threads.y));
    float sx = static_cast<float>(src.cols) / dst.cols;
    float sy = static_cast<float>(src.rows) / dst.rows;
    resize_for_hog_kernel<<<grid, threads>>>(sx, sy, (DevMem2D_<T>)dst, colOfs);
    cudaSafeCall( cudaGetLastError() );
    cudaSafeCall(cudaThreadSynchronize());
    cudaSafeCall(cudaUnbindTexture(resize8UC1_tex));
}

void resize_8UC1(const DevMem2D& src, DevMem2D dst) { resize_for_hog<uchar> (src, dst, resize8UC1_tex); }
void resize_8UC4(const DevMem2D& src, DevMem2D dst) { resize_for_hog<uchar4>(src, dst, resize8UC4_tex); }

}}}