From e5a544f0fb25f32f0baadc80036cf4739c089c36 Mon Sep 17 00:00:00 2001 From: boolli Date: Sat, 20 Feb 2016 11:38:34 -0500 Subject: [PATCH] Finish CUDARGB2GrayScale plugin --- openbr/plugins/cuda/copyto.cpp | 3 --- openbr/plugins/cuda/cudacvt.cpp | 99 --------------------------------------------------------------------------------------------------- openbr/plugins/cuda/cudacvt.cu | 49 ------------------------------------------------- openbr/plugins/cuda/cudadefines.hpp | 2 +- openbr/plugins/cuda/cudargb2grayscale.cpp | 80 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ openbr/plugins/cuda/cudargb2grayscale.cu | 47 +++++++++++++++++++++++++++++++++++++++++++++++ 6 files changed, 128 insertions(+), 152 deletions(-) delete mode 100644 openbr/plugins/cuda/cudacvt.cpp delete mode 100644 openbr/plugins/cuda/cudacvt.cu create mode 100644 openbr/plugins/cuda/cudargb2grayscale.cpp create mode 100644 openbr/plugins/cuda/cudargb2grayscale.cu diff --git a/openbr/plugins/cuda/copyto.cpp b/openbr/plugins/cuda/copyto.cpp index 86b4800..cc63e58 100644 --- a/openbr/plugins/cuda/copyto.cpp +++ b/openbr/plugins/cuda/copyto.cpp @@ -26,9 +26,6 @@ private: const Mat& srcMat = src.m(); const int rows = srcMat.rows; const int cols = srcMat.cols; -std::cout << "CopyTo" << std::endl; - std::cout << "rows: " << rows << std::endl; - std::cout << "cols: " << cols << std::endl; // output will be a single pointer to graphics card memory Mat dstMat = Mat(4, 1, DataType::type); diff --git a/openbr/plugins/cuda/cudacvt.cpp b/openbr/plugins/cuda/cudacvt.cpp deleted file mode 100644 index 1f3297d..0000000 --- a/openbr/plugins/cuda/cudacvt.cpp +++ /dev/null @@ -1,99 +0,0 @@ -/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * - * Copyright 2012 The MITRE Corporation * - * * - * Licensed under the Apache License, Version 2.0 (the "License"); * - * you may not use this file except in compliance with the License. * - * You may obtain a copy of the License at * - * * - * http://www.apache.org/licenses/LICENSE-2.0 * - * * - * Unless required by applicable law or agreed to in writing, software * - * distributed under the License is distributed on an "AS IS" BASIS, * - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. * - * See the License for the specific language governing permissions and * - * limitations under the License. * - * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */ - -#include -#include - -#include -#include -#include -#include - - -using namespace cv; - -namespace br { namespace cuda{ - void cudacvt_wrapper(void* srcPtr, void**dstPtr, int rows, int cols); -}} - -namespace br -{ - -/*! - * \ingroup transforms - * \brief Colorspace conversion. - * \author Li Li \cite Josh Klontz \cite jklontz - */ -class CUDACvtTransform : public UntrainableTransform -{ - Q_OBJECT - Q_ENUMS(ColorSpace) - Q_PROPERTY(ColorSpace colorSpace READ get_colorSpace WRITE set_colorSpace RESET reset_colorSpace STORED false) - Q_PROPERTY(int channel READ get_channel WRITE set_channel RESET reset_channel STORED false) - -public: - enum ColorSpace { Gray = CV_BGR2GRAY, - RGBGray = CV_RGB2GRAY, - HLS = CV_BGR2HLS, - HSV = CV_BGR2HSV, - Lab = CV_BGR2Lab, - Luv = CV_BGR2Luv, - RGB = CV_BGR2RGB, - XYZ = CV_BGR2XYZ, - YCrCb = CV_BGR2YCrCb, - Color = CV_GRAY2BGR }; - -private: - BR_PROPERTY(ColorSpace, colorSpace, Gray) - BR_PROPERTY(int, channel, -1) - - void project(const Template &src, Template &dst) const - { - void* const* srcDataPtr = src.m().ptr(); - int rows = *((int*) srcDataPtr[1]); - int cols = *((int*) srcDataPtr[2]); - int type = *((int*) srcDataPtr[3]); - std::cout << "CVT" << std::endl; - std::cout << "rows: " << rows << std::endl; - std::cout << "cols: " << cols << std::endl; - - Mat dstMat = Mat(src.m().rows, src.m().cols, src.m().type()); - void** dstDataPtr = dstMat.ptr(); - dstDataPtr[1] = srcDataPtr[1]; - dstDataPtr[2] = srcDataPtr[2]; - dstDataPtr[3] = srcDataPtr[3]; - *((int*)dstDataPtr[3]) = CV_8UC1; // not sure if the type of the new mat is the same - - br::cuda::cudacvt_wrapper(srcDataPtr[0], &dstDataPtr[0], rows, cols); - dst = dstMat; - - /* - if (src.m().channels() > 1 || colorSpace == Color) cvtColor(src, dst, colorSpace); - else dst = src; - - if (channel != -1) { - std::vector mv; - split(dst, mv); - dst = mv[channel % (int)mv.size()]; - } */ - } -}; - -BR_REGISTER(Transform, CUDACvtTransform) - -} // namespace br - -#include "imgproc/cudacvt.moc" diff --git a/openbr/plugins/cuda/cudacvt.cu b/openbr/plugins/cuda/cudacvt.cu deleted file mode 100644 index 9c07c6b..0000000 --- a/openbr/plugins/cuda/cudacvt.cu +++ /dev/null @@ -1,49 +0,0 @@ -#include -#include -#include -#include -#include "cudadefines.hpp" -#include -#include -#include -#include - -using namespace std; -using namespace cv; -using namespace cv::gpu; - -namespace br{ namespace cuda { - - __global__ void cudacvt_kernel(uint8_t* srcPtr, uint8_t* dstPtr, int rows, int cols) - { - int rowInd = blockIdx.y*blockDim.y+threadIdx.y; - int colInd = blockIdx.x*blockDim.x+threadIdx.x; - int index = rowInd*cols + colInd; - if (rowInd < 1 || rowInd >= rows-1 || colInd < 1 || colInd >= cols-1) { - if (rowInd >= rows || colInd >= cols) { - return; - } else { - return; - } - } - - dstPtr[index] = 0.299f * srcPtr[3*index] + 0.587f * srcPtr[3*index+1] + 0.114f * srcPtr[3*index+2]; - return; - } - - void cudacvt_wrapper(void* srcPtr, void** dstPtr, int rows, int cols) - { - cudaError_t err; - dim3 threadsPerBlock(8, 8); - dim3 numBlocks(cols/threadsPerBlock.x + 1, - rows/threadsPerBlock.y + 1); - std::cout << "Before malloc" << std::endl; - CUDA_SAFE_MALLOC(dstPtr, rows*cols*sizeof(uint8_t), &err); - std::cout << "After malloc" << std::endl; - - cudacvt_kernel<<>>((uint8_t*)srcPtr, (uint8_t*) (*dstPtr), rows, cols); - CUDA_KERNEL_ERR_CHK(&err); - CUDA_SAFE_FREE(srcPtr, &err); - } - -}} diff --git a/openbr/plugins/cuda/cudadefines.hpp b/openbr/plugins/cuda/cudadefines.hpp index 0930077..a5981e4 100644 --- a/openbr/plugins/cuda/cudadefines.hpp +++ b/openbr/plugins/cuda/cudadefines.hpp @@ -16,7 +16,7 @@ using namespace std; cout << pthread_self() << ": CUDA Malloc Error(" << *errPtr << "): " << cudaGetErrorString(*errPtr) << endl; \ throw 0; \ } \ - cout << pthread_self() << ": CUDA Malloc: " << (void*)*(int**)cudaPtrPtr << endl; + //cout << pthread_self() << ": CUDA Malloc: " << (void*)*(int**)cudaPtrPtr << endl; #define CUDA_SAFE_MEMCPY(dstPtr, srcPtr, count, kind, errPtr) \ *errPtr = cudaMemcpy(dstPtr, srcPtr, count, kind); \ diff --git a/openbr/plugins/cuda/cudargb2grayscale.cpp b/openbr/plugins/cuda/cudargb2grayscale.cpp new file mode 100644 index 0000000..6f72728 --- /dev/null +++ b/openbr/plugins/cuda/cudargb2grayscale.cpp @@ -0,0 +1,80 @@ +/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * + * Copyright 2012 The MITRE Corporation * + * * + * Licensed under the Apache License, Version 2.0 (the "License"); * + * you may not use this file except in compliance with the License. * + * You may obtain a copy of the License at * + * * + * http://www.apache.org/licenses/LICENSE-2.0 * + * * + * Unless required by applicable law or agreed to in writing, software * + * distributed under the License is distributed on an "AS IS" BASIS, * + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. * + * See the License for the specific language governing permissions and * + * limitations under the License. * + * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * */ + +#include +#include + +#include +#include +#include +#include + + +using namespace cv; + +namespace br { namespace cuda{ + void cudargb2grayscale_wrapper(void* srcPtr, void**dstPtr, int rows, int cols); +}} + +namespace br +{ + +/*! + * \ingroup transforms + * \brief Colorspace conversion. + * \author Li Li \cite Josh Klontz \cite jklontz + */ +class CUDARGB2GrayScaleTransform : public UntrainableTransform +{ + Q_OBJECT + +public: + +private: + void project(const Template &src, Template &dst) const + { + void* const* srcDataPtr = src.m().ptr(); + int rows = *((int*) srcDataPtr[1]); + int cols = *((int*) srcDataPtr[2]); + int type = *((int*) srcDataPtr[3]); + + Mat dstMat = Mat(src.m().rows, src.m().cols, src.m().type()); + void** dstDataPtr = dstMat.ptr(); + dstDataPtr[1] = srcDataPtr[1]; + dstDataPtr[2] = srcDataPtr[2]; + dstDataPtr[3] = srcDataPtr[3]; + *((int*)dstDataPtr[3]) = CV_8UC1; // not sure if the type of the new mat is the same + + br::cuda::cudargb2grayscale_wrapper(srcDataPtr[0], &dstDataPtr[0], rows, cols); + dst = dstMat; + + /* + if (src.m().channels() > 1 || colorSpace == Color) cvtColor(src, dst, colorSpace); + else dst = src; + + if (channel != -1) { + std::vector mv; + split(dst, mv); + dst = mv[channel % (int)mv.size()]; + } */ + } +}; + +BR_REGISTER(Transform, CUDARGB2GrayScaleTransform) + +} // namespace br + +#include "imgproc/cudargb2grayscale.moc" diff --git a/openbr/plugins/cuda/cudargb2grayscale.cu b/openbr/plugins/cuda/cudargb2grayscale.cu new file mode 100644 index 0000000..ba10b80 --- /dev/null +++ b/openbr/plugins/cuda/cudargb2grayscale.cu @@ -0,0 +1,47 @@ +#include +#include +#include +#include +#include "cudadefines.hpp" +#include +#include +#include +#include + +using namespace std; +using namespace cv; +using namespace cv::gpu; + +namespace br{ namespace cuda { + + __global__ void cudargb2grayscale_kernel(uint8_t* srcPtr, uint8_t* dstPtr, int rows, int cols) + { + int rowInd = blockIdx.y*blockDim.y+threadIdx.y; + int colInd = blockIdx.x*blockDim.x+threadIdx.x; + int index = rowInd*cols + colInd; + if (rowInd < 1 || rowInd >= rows-1 || colInd < 1 || colInd >= cols-1) { + if (rowInd >= rows || colInd >= cols) { + return; + } else { + return; + } + } + + dstPtr[index] = 0.299f * srcPtr[3*index] + 0.587f * srcPtr[3*index+1] + 0.114f * srcPtr[3*index+2]; + return; + } + + void cudargb2grayscale_wrapper(void* srcPtr, void** dstPtr, int rows, int cols) + { + cudaError_t err; + dim3 threadsPerBlock(8, 8); + dim3 numBlocks(cols/threadsPerBlock.x + 1, + rows/threadsPerBlock.y + 1); + CUDA_SAFE_MALLOC(dstPtr, rows*cols*sizeof(uint8_t), &err); + + cudargb2grayscale_kernel<<>>((uint8_t*)srcPtr, (uint8_t*) (*dstPtr), rows, cols); + CUDA_KERNEL_ERR_CHK(&err); + CUDA_SAFE_FREE(srcPtr, &err); + } + +}} -- libgit2 0.21.4