potok/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2026-07-31 00:03:03 +04:00
Code Packages Releases Wiki Activity

*applied patch from NVidia (nppstTraspose bug)

Browse Source 
*fixed some warnings
*finished gpu test port to gtest framework
This commit is contained in:
Anatoly Baksheev
2011-02-18 12:23:18 +00:00
parent 916690a674
commit 047c7e0fd6
12 changed files with 69 additions and 47 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/gpu/src/nvidia/NPP_staging/NPP_staging.cu
+14 -8
View File
@@ -1198,26 +1198,32 @@ __global__ void transpose(T *d_src, Ncv32u srcStride,
Ncv32u xIndex = blockIdx_x * TRANSPOSE_TILE_DIM + threadIdx.x;
Ncv32u yIndex = blockIdx_y * TRANSPOSE_TILE_DIM + threadIdx.y;
Ncv32u index_in = xIndex + yIndex * srcStride;
Ncv32u index_gmem = xIndex + yIndex * srcStride;
xIndex = blockIdx_y * TRANSPOSE_TILE_DIM + threadIdx.x;
yIndex = blockIdx_x * TRANSPOSE_TILE_DIM + threadIdx.y;
Ncv32u index_out = xIndex + yIndex * dstStride;
for (Ncv32u i=0; i<TRANSPOSE_TILE_DIM; i+=TRANSPOSE_BLOCK_ROWS)
if (xIndex < srcRoi.width)
{
tile[threadIdx.y+i][threadIdx.x] = d_src[index_in+i*srcStride];
for (Ncv32u i=0; i<TRANSPOSE_TILE_DIM; i+=TRANSPOSE_BLOCK_ROWS)
{
if (yIndex + i < srcRoi.height)
{
tile[threadIdx.y+i][threadIdx.x] = d_src[index_gmem+i*srcStride];
}
}
}
__syncthreads();
xIndex = blockIdx_y * TRANSPOSE_TILE_DIM + threadIdx.x;
yIndex = blockIdx_x * TRANSPOSE_TILE_DIM + threadIdx.y;
index_gmem = xIndex + yIndex * dstStride;
if (xIndex < srcRoi.height)
{
for (Ncv32u i=0; i<TRANSPOSE_TILE_DIM; i+=TRANSPOSE_BLOCK_ROWS)
{
if (yIndex + i < srcRoi.width)
{
d_dst[index_out+i*dstStride] = tile[threadIdx.x][threadIdx.y+i];
d_dst[index_gmem+i*dstStride] = tile[threadIdx.x][threadIdx.y+i];
}
}
}
modules/gpu/src/opencv2/gpu/device/transform.hpp
+8 -8
View File
@@ -71,21 +71,21 @@ namespace cv { namespace gpu { namespace device
template <size_t src_elem_size, size_t dst_elem_size>
struct UnReadWriteTraits_
{
enum {shift=1};
enum { shift = 1 };
};
template <size_t src_elem_size>
struct UnReadWriteTraits_<src_elem_size, 1>
{
enum {shift=4};
enum { shift = 4 };
};
template <size_t src_elem_size>
struct UnReadWriteTraits_<src_elem_size, 2>
{
enum {shift=2};
enum { shift = 2 };
};
template <typename T, typename D> struct UnReadWriteTraits
{
enum {shift=UnReadWriteTraits_<sizeof(T), sizeof(D)>::shift};
enum { shift = UnReadWriteTraits_<sizeof(T), sizeof(D)>::shift };
typedef typename TypeVec<T, shift>::vec_t read_type;
typedef typename TypeVec<D, shift>::vec_t write_type;
@@ -94,21 +94,21 @@ namespace cv { namespace gpu { namespace device
template <size_t src_elem_size1, size_t src_elem_size2, size_t dst_elem_size>
struct BinReadWriteTraits_
{
enum {shift=1};
enum { shift = 1 };
};
template <size_t src_elem_size1, size_t src_elem_size2>
struct BinReadWriteTraits_<src_elem_size1, src_elem_size2, 1>
{
enum {shift=4};
enum { shift = 4 };
};
template <size_t src_elem_size1, size_t src_elem_size2>
struct BinReadWriteTraits_<src_elem_size1, src_elem_size2, 2>
{
enum {shift=2};
enum { shift = 2 };
};
template <typename T1, typename T2, typename D> struct BinReadWriteTraits
{
enum {shift=BinReadWriteTraits_<sizeof(T1), sizeof(T2), sizeof(D)>::shift};
enum {shift = BinReadWriteTraits_<sizeof(T1), sizeof(T2), sizeof(D)>::shift};
typedef typename TypeVec<T1, shift>::vec_t read_type1;
typedef typename TypeVec<T2, shift>::vec_t read_type2;