From 699ab0b0d58a813f4625059ca1ee1cc5aedb6b57 Mon Sep 17 00:00:00 2001 From: boolli Date: Fri, 19 Feb 2016 10:31:39 -0500 Subject: [PATCH] Add kernel function but it's not working --- openbr/plugins/cuda/copyfrom.cpp | 3 +++ openbr/plugins/cuda/copyto.cpp | 8 +++++++- openbr/plugins/cuda/cudacvt.cpp | 99 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ openbr/plugins/cuda/cudacvt.cu | 49 +++++++++++++++++++++++++++++++++++++++++++++++++ openbr/plugins/cuda/cudadefines.hpp | 2 +- 5 files changed, 159 insertions(+), 2 deletions(-) create mode 100644 openbr/plugins/cuda/cudacvt.cpp create mode 100644 openbr/plugins/cuda/cudacvt.cu diff --git a/openbr/plugins/cuda/copyfrom.cpp b/openbr/plugins/cuda/copyfrom.cpp index 8b0a88b..44b6343 100644 --- a/openbr/plugins/cuda/copyfrom.cpp +++ b/openbr/plugins/cuda/copyfrom.cpp @@ -36,6 +36,9 @@ private: case CV_8UC1: br::cuda::cudacopyfrom::wrapper(dataPtr[0], dstMat.ptr(), rows, cols); break; + case CV_8UC3: + br::cuda::cudacopyfrom::wrapper(dataPtr[0], dstMat.ptr(), rows, cols * 3); + break; default: cout << "ERR: Invalid image format" << endl; break; diff --git a/openbr/plugins/cuda/copyto.cpp b/openbr/plugins/cuda/copyto.cpp index 902c335..86b4800 100644 --- a/openbr/plugins/cuda/copyto.cpp +++ b/openbr/plugins/cuda/copyto.cpp @@ -25,7 +25,10 @@ private: { const Mat& srcMat = src.m(); const int rows = srcMat.rows; - const int cols = srcMat.cols; + 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); @@ -44,6 +47,9 @@ private: case CV_8UC1: br::cuda::cudacopyto::wrapper(srcMat.ptr(), &dstMatData[0], rows, cols); break; + case CV_8UC3: + br::cuda::cudacopyto::wrapper(srcMat.ptr(), &dstMatData[0], rows, 3*cols); + break; default: cout << "ERR: Invalid image type! " << type2str(srcMat.type()) << endl; return; diff --git a/openbr/plugins/cuda/cudacvt.cpp b/openbr/plugins/cuda/cudacvt.cpp new file mode 100644 index 0000000..28e4999 --- /dev/null +++ b/openbr/plugins/cuda/cudacvt.cpp @@ -0,0 +1,99 @@ +/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * + * 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, CV_8UC1); + 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 new file mode 100644 index 0000000..372c36f --- /dev/null +++ b/openbr/plugins/cuda/cudacvt.cu @@ -0,0 +1,49 @@ +#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; + 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 354be2b..0930077 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); \ -- libgit2 0.21.4