Add kernel function but it's not working

boolli
1 parent 4e59b2b5
Showing 5 changed files with 159 additions and 2 deletions
openbr/plugins/cuda/copyfrom.cpp
openbr/plugins/cuda/copyto.cpp
openbr/plugins/cuda/cudacvt.cpp
openbr/plugins/cuda/cudacvt.cu
openbr/plugins/cuda/cudadefines.hpp
@@ -36,6 +36,9 @@ private:
       case CV_8UC1:
         br::cuda::cudacopyfrom::wrapper(dataPtr[0], dstMat.ptr<unsigned char>(), rows, cols);
         break;
+      case CV_8UC3:
+        br::cuda::cudacopyfrom::wrapper(dataPtr[0], dstMat.ptr<unsigned char>(), rows, cols * 3);
+        break;
       default:
         cout << "ERR: Invalid image format" << endl;
         break;
@@ -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<void*>::type);
@@ -44,6 +47,9 @@ private:
       case CV_8UC1:
         br::cuda::cudacopyto::wrapper(srcMat.ptr<unsigned char>(), &dstMatData[0], rows, cols);
         break;
+      case CV_8UC3:
+        br::cuda::cudacopyto::wrapper(srcMat.ptr<unsigned char>(), &dstMatData[0], rows, 3*cols);
+        break;
       default:
         cout << "ERR: Invalid image type! " << type2str(srcMat.type()) << endl;
         return;
+/* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * *
+ * 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 <iostream>
+#include <opencv2/imgproc/imgproc.hpp>
+
+#include <openbr/plugins/openbr_internal.h>
+#include <openbr/core/opencvutils.h>
+#include <thrust/host_vector.h>
+#include <thrust/device_vector.h>
+
+
+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<void*>();
+        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<void*>();
+        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<Mat> mv;
+            split(dst, mv);
+            dst = mv[channel % (int)mv.size()];
+        } */
+    }
+};
+
+BR_REGISTER(Transform, CUDACvtTransform)
+
+} // namespace br
+
+#include "imgproc/cudacvt.moc"
+#include <stdio.h>
+#include <iostream>
+#include <opencv2/opencv.hpp>
+#include <opencv2/gpu/gpu.hpp>
+#include "cudadefines.hpp"
+#include <opencv2/imgproc/imgproc.hpp>
+#include <opencv2/imgproc/imgproc_c.h>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/highgui/highgui_c.h>
+
+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<<<numBlocks, threadsPerBlock>>>((uint8_t*)srcPtr, (uint8_t*) (*dstPtr), rows, cols);
+    CUDA_KERNEL_ERR_CHK(&err);
+    CUDA_SAFE_FREE(srcPtr, &err);
+  } 
+
+}}
@@ -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); \