cleaned up PCA

DepthDeluxe
1 parent fd23ffcd
Showing 3 changed files with 26 additions and 43 deletions
openbr/plugins/cuda/copyfrom.cpp
openbr/plugins/cuda/cudapca.cpp
openbr/plugins/cuda/cudapca.cu
@@ -2,6 +2,8 @@
 #include <opencv2/opencv.hpp>
+//#include <gperftools/profiler.h>
+
 #include <openbr/plugins/openbr_internal.h>
 using namespace std;
@@ -22,6 +24,8 @@ namespace br
 private:
     void project(const Template &src, Template &dst) const
     {
+//      ProfilerStart("PROFILEME.log");
+
       // pull the data back out of the Mat
       void* const* dataPtr = src.m().ptr<void*>();
       int rows = *((int*)dataPtr[1]);
@@ -44,6 +48,8 @@ private:
         break;
       }
       dst = dstMat;
+
+//      ProfilerStop();
     }
   };
@@ -29,11 +29,10 @@ using namespace cv;
 #include <openbr/core/eigenutils.h>
 #include <openbr/core/opencvutils.h>
-namespace br { namespace cuda {
-  void cudapca_loadwrapper(float* evPtr, int evRows, int evCols, float* meanPtr, int meanElems);
-  void cudapca_trainwrapper(void* cudaDataPtr, float* dataPtr, int rows, int cols);
-  void cudapca_projectwrapper(void* src, void** dst);
-}}
+namespace br { namespace cuda { namespace pca {
+  void loadwrapper(float* evPtr, int evRows, int evCols, float* meanPtr, int meanElems);
+  void wrapper(void* src, void** dst);
+}}}
 namespace br
 {
@@ -95,7 +94,6 @@ private:
           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));
         }
@@ -136,7 +134,7 @@ private:
       dstDataPtr[2] = srcDataPtr[2];  *((int*)dstDataPtr[2]) = keep;
       dstDataPtr[3] = srcDataPtr[3];
-      br::cuda::cudapca_projectwrapper(srcDataPtr[0], &dstDataPtr[0]);
+      br::cuda::pca::wrapper(srcDataPtr[0], &dstDataPtr[0]);
       dst = dstMat;
@@ -161,7 +159,11 @@ private:
     void load(QDataStream &stream)
     {
-        stream >> keep >> drop >> whiten >> originalRows >> mean >> eVals >> eVecs;
+        Eigen::MatrixXf originalEVecs;
+        stream >> keep >> drop >> whiten >> originalRows >> mean >> eVals >> originalEVecs;
+
+        // perform transpose before copying over
+        eVecs = originalEVecs; //originalEVecs.transpose();
         cout << "Mean Dimensions" << endl;
         cout << "\tRows: " << mean.rows() << " Cols: " << mean.cols() << endl;
@@ -173,6 +175,7 @@ private:
         cout << "Mean first value: " << mean(0, 0) << endl;
+
         // TODO(colin): use Eigen Map class to generate map files so we don't have to copy the data
         // serialize the eigenvectors
         float* evBuffer = new float[eVecs.rows() * eVecs.cols()];
@@ -191,7 +194,7 @@ private:
         }
         // call the wrapper function
-        cuda::cudapca_loadwrapper(evBuffer, eVecs.rows(), eVecs.cols(), meanBuffer, mean.rows()*mean.cols());
+        br::cuda::pca::loadwrapper(evBuffer, eVecs.rows(), eVecs.cols(), meanBuffer, mean.rows()*mean.cols());
         delete evBuffer;
         delete meanBuffer;
@@ -9,28 +9,8 @@ using namespace std;
 using namespace cv;
 using namespace cv::gpu;
-namespace br { namespace cuda {
-  __global__ void calculateCovariance_kernel(float* trainingSet, float* cov, int numRows, int numCols) {
-    int rowInd = blockIdx.y*blockDim.y + threadIdx.y;
-    int colInd = blockIdx.x*blockDim.x + threadIdx.x;
-
-    // this calculates trainingSet' * trainingSet
-    if (rowInd >= numRows || colInd >= numCols) {
-      return;
-    }
-
-    // get a reference the value we wish to write
-    float& out = cov[rowInd*numRows + colInd];
-
-    // calculate the value of this position
-    out = 0;
-    for (int i=0; i<numRows; i++) {
-      out += trainingSet[rowInd*numCols + colInd] * trainingSet[rowInd*numCols + numRows]; // XXX(colin): not sure if this is correct
-    }
-    out = out / (numRows-1);
-  }
-
-  __global__ void cudapca_project_multiply_kernel(float* src, float* dst, float* evPtr, int evRows, int evCols) {
+namespace br { namespace cuda { namespace pca {
+  __global__ void multiplyKernel(float* src, float* dst, float* evPtr, int evRows, int evCols) {
     int colInd = blockIdx.x*blockDim.x+threadIdx.x;
     // check dimensions
@@ -44,7 +24,7 @@ namespace br { namespace cuda {
     }
   }
-  __global__ void cudapca_project_subtractmean_kernel(float* out, float* mean, int numCols) {
+  __global__ void subtractMeanKernel(float* out, float* mean, int numCols) {
     int colInd = blockIdx.x*blockDim.x+threadIdx.x;
     // perform bound checking
@@ -61,7 +41,7 @@ namespace br { namespace cuda {
   float* _cudaSrcPtr;
   float* _cudaDstPtr;
-  void cudapca_loadwrapper(float* evPtr, int evRows, int evCols, float* meanPtr, int meanElems) {
+  void loadwrapper(float* evPtr, int evRows, int evCols, float* meanPtr, int meanElems) {
     _evRows = evRows; _evCols = evCols;
     _meanElems = meanElems;
@@ -79,13 +59,7 @@ namespace br { namespace cuda {
     CUDA_SAFE_MALLOC(&_cudaDstPtr, _evCols*sizeof(float), &err);
   }
-  void cudapca_trainwrapper(void* cudaDataPtr, float* dataPtr, int rows, int cols) {
-    cudaError_t err;
-    CUDA_SAFE_MEMCPY(dataPtr, cudaDataPtr, rows*cols*sizeof(float), cudaMemcpyDeviceToHost, &err);
-    CUDA_SAFE_FREE(cudaDataPtr, &err);
-  }
-
-  void cudapca_projectwrapper(void* src, void** dst) {
+  void wrapper(void* src, void** dst) {
     // copy the image to the GPU
     //cudaMemcpy(_cudaSrcPtr, src, _meanElems*sizeof(float), cudaMemcpyHostToDevice);
     cudaError_t err;
@@ -94,13 +68,13 @@ namespace br { namespace cuda {
     // 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>>>((float*)src, cudaMeanPtr, _meanElems);
+    subtractMeanKernel<<<numBlocks, threadsPerBlock>>>((float*)src, cudaMeanPtr, _meanElems);
     CUDA_KERNEL_ERR_CHK(&err);
     // perform the multiplication
     threadsPerBlock = 64;
     numBlocks = _evCols / threadsPerBlock + 1;
-    cudapca_project_multiply_kernel<<<numBlocks, threadsPerBlock>>>((float*)src, (float*)(*dst), cudaEvPtr, _evRows, _evCols);
+    multiplyKernel<<<numBlocks, threadsPerBlock>>>((float*)src, (float*)(*dst), cudaEvPtr, _evRows, _evCols);
     CUDA_KERNEL_ERR_CHK(&err);
     CUDA_SAFE_FREE(src, &err);    // TODO(colin): figure out why adding this free causes memory corruption...
@@ -108,4 +82,4 @@ namespace br { namespace cuda {
     // copy the data back to the CPU
     //cudaMemcpy(dst, _cudaDstPtr, _evCols*sizeof(float), cudaMemcpyDeviceToHost);
   }
-}}
+}}}