added streamlined CUDA execution, fixed memory problems

DepthDeluxe
1 parent 7f5c5a5a
Showing 8 changed files with 69 additions and 130 deletions
openbr/plugins/cuda/copyfrom.cpp
openbr/plugins/cuda/cudacvtfloat.cpp
openbr/plugins/cuda/cudacvtfloat.cu
openbr/plugins/cuda/cudalbp.cpp
openbr/plugins/cuda/cudalbp.cu
openbr/plugins/cuda/cudapca.cpp
openbr/plugins/cuda/cudapca.cu
openbr/plugins/cuda/cudapca.hpp
@@ -22,8 +22,6 @@ namespace br
 private:
     void project(const Template &src, Template &dst) const
     {
-      cout << "CUDACopyFrom Start" << endl << endl << endl;
-
       // pull the data back out of the Mat
       void* const* dataPtr = src.m().ptr<void*>();
       void* cudaMemPtr = dataPtr[0];
@@ -44,8 +42,6 @@ private:
         break;
       }
       dst = dstMat;
-
-      cout << "CUDACopyFrom End" << endl;
     }
   };
  
@@ -27,8 +27,6 @@ class CUDACvtFloatTransform : public UntrainableTransform
   public:
     void project(const Template &src, Template &dst) const
     {
-      cout << "CUDACvtFloat Start" << endl;
-
       void* const* srcDataPtr = src.m().ptr<void*>();
       void* srcMemPtr = srcDataPtr[0];
       int rows = *((int*)srcDataPtr[1]);
@@ -51,8 +49,6 @@ class CUDACvtFloatTransform : public UntrainableTransform
  
       br::cuda::cudacvtfloat::wrapper((const unsigned char*)srcMemPtr, &dstDataPtr[0], rows, cols);
       dst = dstMat;
-
-      cout << "CUDACvtFloat End" << endl;
     }
 };
  
@@ -31,6 +31,9 @@ namespace br { namespace cuda { namespace cudacvtfloat {
     );
  
     kernel<<<threadsPerBlock, blocks>>>(src, (float*)(*dst), rows, cols);
+
+    // free the src memory since it is now in a newly allocated dst
+    cudaFree((void*)src);
   }
  
 }}}
@@ -166,8 +166,6 @@ class CUDALBPTransform : public UntrainableTransform
         //matManager->release(a);
         //matManager->release(b);
  
-        cout << "CUDALBP Start" << endl;
-
         void* const* srcDataPtr = src.m().ptr<void*>();
         void* cudaSrcPtr = srcDataPtr[0];
         int rows = *((int*)srcDataPtr[1]);
@@ -182,8 +180,6 @@ class CUDALBPTransform : public UntrainableTransform
  
         br::cuda::cudalbp_wrapper(cudaSrcPtr, &dstDataPtr[0], rows, cols);
         dst = dstMat;
-
-        cout << "CUDALBP End" << endl;
     }
 };
  
@@ -47,6 +47,8 @@ namespace br { namespace cuda {
  
     cudaMalloc(dstPtr, rows*cols*sizeof(uint8_t));
     cudalbp_kernel<<<numBlocks, threadsPerBlock>>>((uint8_t*)srcPtr, (uint8_t*)(*dstPtr), rows, cols, lut);
+
+    cudaFree(srcPtr);
   }
  
   void cudalbp_init_wrapper(uint8_t* cpuLut) {
@@ -16,14 +16,23 @@
 #include <iostream>
 using namespace std;
  
+#include <QList>
+
 #include <Eigen/Dense>
-#include <openbr/plugins/openbr_internal.h>
  
+#include <opencv2/opencv.hpp>
+using namespace cv;
+
+#include <openbr/plugins/openbr_internal.h>
 #include <openbr/core/common.h>
 #include <openbr/core/eigenutils.h>
 #include <openbr/core/opencvutils.h>
  
-#include "cudapca.hpp"
+namespace br { namespace cuda {
+  void cudapca_loadwrapper(float* evPtr, int evRows, int evCols, float* meanPtr, int meanElems);
+  void cudapca_trainwrapper(const void* cudaDataPtr, float* dataPtr, int rows, int cols);
+  void cudapca_projectwrapper(void* src, void** dst);
+}}
  
 namespace br
 {
@@ -71,14 +80,30 @@ private:
         return (srcMap - mean).squaredNorm() - projMap.squaredNorm();
     }
  
-    void train(const TemplateList &trainingSet)
+    void train(const TemplateList &cudaTrainingSet)
     {
+        const int instances = cudaTrainingSet.size();       // get the number of training set instances
+        QList<Template> trainingQlist;
+        for(int i=0; i<instances; i++) {
+          Template currentTemplate = cudaTrainingSet[i];
+          void* const* srcDataPtr = currentTemplate.m().ptr<void*>();
+          const void* cudaMemPtr = srcDataPtr[0];
+          int rows = *((int*)srcDataPtr[1]);
+          int cols = *((int*)srcDataPtr[2]);
+          int type = *((int*)srcDataPtr[3]);
+
+          Mat mat = Mat(rows, cols, type);
+          br::cuda::cudapca_trainwrapper(cudaMemPtr, mat.ptr<float>(), rows, cols);
+          trainingQlist.append(Template(mat));
+        TemplateList trainingSet;
+        }
+        TemplateList trainingSet(trainingQlist);
+
         if (trainingSet.first().m().type() != CV_32FC1)
             qFatal("Requires single channel 32-bit floating point matrices.");
  
         originalRows = trainingSet.first().m().rows;    // get number of rows of first image
         int dimsIn = trainingSet.first().m().rows * trainingSet.first().m().cols; // get the size of the first image
-        const int instances = trainingSet.size();       // get the number of training set instances
  
         // Map into 64-bit Eigen matrix
         Eigen::MatrixXd data(dimsIn, instances);        // create a mat
@@ -90,10 +115,32 @@ private:
  
     void project(const Template &src, Template &dst) const
     {
-        dst = cv::Mat(1, keep, CV_32FC1);
+
+      void* const* srcDataPtr = src.m().ptr<void*>();
+      void* cudaMemPtr = srcDataPtr[0];
+      int rows = *((int*)srcDataPtr[1]);
+      int cols = *((int*)srcDataPtr[2]);
+      int type = *((int*)srcDataPtr[3]);
+
+      if (type != CV_32FC1) {
+        cout << "ERR: Invalid image type" << endl;
+        return;
+      }
+
+      Mat dstMat = Mat(src.m().rows, src.m().cols, src.m().type());
+      void** dstDataPtr = dstMat.ptr<void*>();
+      dstDataPtr[1] = srcDataPtr[1];  *((int*)dstDataPtr[1]) = 1;
+      dstDataPtr[2] = srcDataPtr[2];  *((int*)dstDataPtr[2]) = keep;
+      dstDataPtr[3] = srcDataPtr[3];
+
+      br::cuda::cudapca_projectwrapper(cudaMemPtr, &dstDataPtr[0]);
+
+      dst = dstMat;
+
+        //dst = cv::Mat(1, keep, CV_32FC1);
  
         // perform the operation on the graphics card
-        cuda::cudapca_projectwrapper((float*)src.m().ptr<float>(), (float*)dst.m().ptr<float>());
+        //cuda::cudapca_projectwrapper((float*)src.m().ptr<float>(), (float*)dst.m().ptr<float>());
  
         // Map Eigen into OpenCV
         //Mat cpuDst = cv::Mat(1, keep, CV_32FC1);
@@ -7,8 +7,6 @@ using namespace std;
 using namespace cv;
 using namespace cv::gpu;
  
-#include "cudapca.hpp"
-
 namespace br { namespace cuda {
   __global__ void calculateCovariance_kernel(float* trainingSet, float* cov, int numRows, int numCols) {
     int rowInd = blockIdx.y*blockDim.y + threadIdx.y;
@@ -77,116 +75,29 @@ namespace br { namespace cuda {
     cudaMalloc(&_cudaDstPtr, _evCols*sizeof(float));
   }
  
-  void cudapca_trainwrapper() {
-    /*
-    if (trainingSet[0].type() != CV_32FC1) {
-      std::cout << "ERR: Requires single 32-bit floating point matrix!";
-      return;
-    }
-
-    cudaError_t status;
-
-    const int numExamples = trainingSetSize;
-    int numPixels = trainingSet[0].rows * trainingSet[0].cols;
-
-    // create a custom matrix
-    float* cudaDataPtr;
-    status = cudaMalloc(&cudaDataPtr, numPixels * numExamples * sizeof(float));
-    if (status != cudaSuccess) {
-      std::cout << "ERR: Memory allocation" << std::endl;
-      return;
-    }
-
-    // copy all the data to the graphics card
-    for (int i=0; i < numExamples; i++) {
-      status = cudaMemcpy(cudaDataPtr + i*numPixels, trainingSet[i].ptr<float>(), numPixels*sizeof(float), cudaMemcpyHostToDevice);
-      if (status != cudaSuccess) {
-        std::cout << "ERR: Memcpy at index " << i << std::endl;
-        return;
-      }
-    }
-
-    // start the core part of the algorithm
-    int numDimensions = numPixels;
-    const bool dominantEigenEstimation = (numDimensions > numExamples);
-
-    // malloc and init mean
-    mean = new float[numDimensions];
-    for (int i=0; i < numDimensions; i++) {
-      mean[i] = 0;
-    }
-    float* cudaMeanPtr;
-    status = cudaMalloc(&cudaMeanPtr, numDimensions*sizeof(float));
-    if (status != cudaSuccess) {
-      std::cout << " ERR: Malloc of mean" << std::endl;
-      return;
-    }
-
-    if (keep != 0) {
-      // compute the mean so we can subtract from data
-      for (int i=0; i < numExamples; i++) {
-        Mat& m = trainingSet[i];
-
-        for (int j=0; j < numDimensions; j++) {
-          mean[j] += m.ptr<float>()[i*numDimensions + j];
-        }
-      }
-      for (int i=0; i < numDimensions; i++) {
-        mean[i] = mean[i] / numExamples;
-      }
-
-      // copy mean over to graphics card
-      cudaMemcpy(cudaMeanPtr, mean, numExamples*sizeof(float), cudaMemcpyHostToDevice);
-      if (status != cudaSuccess) {
-        std::cout << " ERR: Cpy of mean" << std::endl;
-        return;
-      }
-
-      // set the thread dimensions and run the kernel
-      dim3 threadsPerBlock(64, 1);
-      dim3 numBlocks(numDimensions/threadsPerBlock.x + 1,
-                     numExamples/threadsPerBlock.y + 1);
-
-      subtractMean_kernel<<<numBlocks, threadsPerBlock>>>(cudaDataPtr, cudaMeanPtr, numExamples, numDimensions);
-
-      // calculate the covariance matrix using kernel
-      // malloc location for covariance matrix
-      float* cudaCovPtr;
-      status = cudaMalloc(&cudaCovPtr, numExamples*numExamples*sizeof(float));
-      if (status != cudaSuccess) h
-        std::cout << " ERR: Cpy of mean" << std::endl;
-        return;
-      }
-
-      // calculate the covariance matrix
-      threadsPerBlock = dim3(8, 8);
-      numBlocks = dim3(numExamples/threadsPerBlock.x + 1,
-                       numExamples/threadsPerBlock.y + 1);
-      calculateCovariance_kernel<<<numBlocks, threadsPerBlock>>>(cudaDataPtr, cudaCovPtr, numExamples, numDimensions);
-
-      // perform eigendecomposition
-      //std::cout << "Skipping eigendecomposition" << std::endl;
-      cusolverStatus_t cusolverStatus;
-      cusolverStatus = cusolverDnSgebrd(cusolverHandle,)
-    }
-    */
+  void cudapca_trainwrapper(const void* cudaDataPtr, float* dataPtr, int rows, int cols) {
+    cudaMemcpy(dataPtr, cudaDataPtr, rows*cols*sizeof(float), cudaMemcpyDeviceToHost);
   }
  
-  void cudapca_projectwrapper(float* src, float* dst) {
+  void cudapca_projectwrapper(void* src, void** dst) {
     // copy the image to the GPU
-    cudaMemcpy(_cudaSrcPtr, src, _meanElems*sizeof(float), cudaMemcpyHostToDevice);
+    //cudaMemcpy(_cudaSrcPtr, src, _meanElems*sizeof(float), cudaMemcpyHostToDevice);
+
+    cudaMalloc(dst, _evRows*_evCols*sizeof(float));
  
     // subtract out the mean of the image (mean is 1xpixels in size)
     int threadsPerBlock = 64;
     int numBlocks = _meanElems / threadsPerBlock + 1;
-    cudapca_project_subtractmean_kernel<<<numBlocks, threadsPerBlock>>>(_cudaSrcPtr, cudaMeanPtr, _meanElems);
+    cudapca_project_subtractmean_kernel<<<numBlocks, threadsPerBlock>>>((float*)src, cudaMeanPtr, _meanElems);
  
     // perform the multiplication
     threadsPerBlock = 64;
     numBlocks = _evCols / threadsPerBlock + 1;
-    cudapca_project_multiply_kernel<<<numBlocks, threadsPerBlock>>>(_cudaSrcPtr, _cudaDstPtr, cudaEvPtr, _evRows, _evCols);
+    cudapca_project_multiply_kernel<<<numBlocks, threadsPerBlock>>>((float*)src, (float*)(*dst), cudaEvPtr, _evRows, _evCols);
+
+    //cudaFree(src);    // TODO(colin): figure out why adding this free causes memory corruption...
  
     // copy the data back to the CPU
-    cudaMemcpy(dst, _cudaDstPtr, _evCols*sizeof(float), cudaMemcpyDeviceToHost);
+    //cudaMemcpy(dst, _cudaDstPtr, _evCols*sizeof(float), cudaMemcpyDeviceToHost);
   }
 }}
-#include <opencv2/opencv.hpp>
-#include <opencv2/gpu/gpu.hpp>
-
-using namespace cv;
-using namespace cv::gpu;
-
-namespace br { namespace cuda {
-  void cudapca_loadwrapper(float* evPtr, int evRows, int evCols, float* meanPtr, int meanElems);
-  void cudapca_trainwrapper();
-
-  void cudapca_projectwrapper(float* src, float* dst);
-}}