協方差矩陣與opencl和opencv

我最近開始學習opencl以及如何爲OpenCV創建內核。我仍在使用基礎知識。協方差矩陣與opencl和opencv

我需要實現一個opencl內核來計算協方差矩陣。這個函數沒有opencv基金已經實現的opencl內核。

技術上我要處理的是：

cv::RNG tutu(std::time(nullptr)); // This is only for check the improvement of the method over the executions 

int sz = tutu.uniform(1,20); 

cv::Mat_<float> a1(1,sz); 
cv::Mat_<float> a2(a1.size()); 
cv::Mat_<float> c2; 

for(std::size_t i=0;i<sz;i++) 
{ 
    a1(i) = tutu.uniform(0,300); 
    a2(i) = tutu.uniform(300,600); 
} 

cv::Mat_<float> f; 

cv::vconcat(a1,a2,f); 

// process the Covariance Matrix : 

cv::gemm(one,f,-0.5f,f,1.f,c2,0); 
cv::gemm(c2.t(),c2,1.f,cv::noArray(),0.f,c2);

我發現的OpenCV的基金會已經實現了在OpenCL內核都GEMM和功能轉。

所以我試圖直接從OpenCV的源代碼衍生出一個實例。

我寫了這個：

ocl.h：

void CovarMatrix(cv::Mat_<float>& src,cv::Mat_<float>& covar);

ocl.cpp：

#include <memory> 
    #include <fstream> 
    #include <sstream> 
    #include <iterator> 

    #include <opencv2/core.hpp> 
    #include <opencv2/core/ocl.hpp> 


    namespace test 
    { 

    namespace ocl 
    { 

    namespace 
    { 

    std::unique_ptr<cv::ocl::ProgramSource> cov_src; 

    void init_cov() 
    { 

     std::ifstream stream("../mahalanobis/covarianceMatrix.cl"); 
     std::ostringstream sstream; 

     sstream << stream.rdbuf(); 

     cv::String norm_file_content = sstream.str(); 

     stream.close(); 

     cov_src.reset(new cv::ocl::ProgramSource(norm_file_content)); 
    } 


    } 

static bool ocl_gemm(cv::Mat_<float>& matA, cv::Mat_<float>& matB, cv::Mat_<float>& CV_OUT matD) 
{ 

    cv::Mat_<float> tmp; 
    cv::Mat_<float> tmp2; 

    cv::Size sizeA = matA.size(), sizeB = matB.size(); 

    cv::Size sizeD(sizeB.width, sizeA.height); 

    const cv::ocl::Device & dev = cv::ocl::Device::getDefault(); 
    int max_wg_size = (int)dev.maxWorkGroupSize(); 
    int block_size = (max_wg_size/32 < 32) ? (max_wg_size/16 < 16) ? (max_wg_size/8 < 8) ? 1 : 8 : 16 : 32; 

// matD.create(sizeD); 
// tmp2.create(matD.t().size()); 

    tmp.create(sizeD); 
    tmp2.create(tmp.t().size()); 
    matD.create(sizeD.width,sizeD.width); 

    cv::UMat A = matA.getUMat(cv::ACCESS_READ,cv::USAGE_ALLOCATE_DEVICE_MEMORY); 
    cv::UMat B = matB.getUMat(cv::ACCESS_READ,cv::USAGE_ALLOCATE_DEVICE_MEMORY); 
// cv::UMat D = matD.getUMat(cv::ACCESS_WRITE,cv::USAGE_ALLOCATE_DEVICE_MEMORY); 
    cv::UMat D = tmp.getUMat(cv::ACCESS_WRITE,cv::USAGE_ALLOCATE_DEVICE_MEMORY); 
// cv::UMat E(sizeD.width,sizeD.height,CV_32FC1,cv::Scalar::all(0.),cv::USAGE_ALLOCATE_DEVICE_MEMORY); 

    cv::UMat E = tmp2.getUMat(cv::ACCESS_WRITE,cv::USAGE_ALLOCATE_DEVICE_MEMORY); 
    cv::UMat F = matD.getUMat(cv::ACCESS_WRITE,cv::USAGE_ALLOCATE_DEVICE_MEMORY); 

    matB.copyTo(D); 



    int vectorWidths[] = { 4, 4, 2, 2, 1, 4, 1, -1 }; 
    int kercn = cv::ocl::checkOptimalVectorWidth(vectorWidths, B, D); 

    cv::String opts = cv::format(
       "-I /home/administrateur/lib_dir/opencv_dir/opencv_304/opencv/modules/core/src/opencl/ -D T=float -D T1=float -D WT=%s -D cn=1 -D kercn=%d -D LOCAL_SIZE=%d %s -D HAVE_C -D TILE_DIM=32 -D BLOCK_ROWS=8 -D rowsPerWI=1 ", 

          cv::ocl::typeToStr(CV_32FC(kercn)), 
          kercn, block_size, 
          (sizeA.width % block_size !=0) ? "-D NO_MULT" : ""); 

    cv::ocl::Kernel k("covarianceMatrix", *cov_src, opts); 


    k.args(cv::ocl::KernelArg::ReadOnlyNoSize(A), 
      cv::ocl::KernelArg::ReadOnlyNoSize(B, 1, kercn), 
      cv::ocl::KernelArg::ReadWrite(D, 1, kercn), 
      sizeA.width, 
      cv::ocl::KernelArg::ReadWrite(E,kercn,1), 
      cv::ocl::KernelArg::ReadWrite(F,kercn,kercn) 
      ); 

    std::size_t globalsize[2] = { static_cast<std::size_t>(sizeD.width/kercn), static_cast<std::size_t>(sizeD.height)}; 
    std::size_t localsize[2] = { static_cast<std::size_t>(block_size), static_cast<std::size_t>(block_size)}; 
    return k.run(2, globalsize, block_size!=1 ? localsize : nullptr, false); 

} 





void CovarMatrix(cv::Mat_<float>& src,cv::Mat_<float>& covar) 
{ 

    if(!covar.empty()) 
     covar.release(); 


    cv::Mat_<float> o = cv::Mat_<float>::ones(src.rows,src.rows); 

    if(!cov_src) 
     init_cov(); 

    ocl_gemm(o,src,covar); 


}

covarianceMatrix.cl：

#include "gemm.cl" 
#include "transpose.cl" 



__kernel void covarianceMatrix 
(
        __global const uchar * A_ptr, int A_step, int A_offset, 
        __global const uchar * B_ptr, int B_step, int B_offset, 
        __global uchar * D_ptr, int D_step, int D_offset, int D_rows, int D_cols, 
        int n, 
        __global uchar * E_ptr, int E_step, int E_offset, int E_rows, int E_cols, 
        __global uchar * F_ptr, int F_step, int F_offset, int F_rows, int F_cols 
        ) 

{ 


// cv::gemm(src2,src1,-0.5,src1,1.f,src2); 
// cv::gemm(src2.t(),src2,1.f,cv::noArray(),0.f,dest); 



    gemm(A_ptr,A_step,A_offset, 
     B_ptr,B_step,B_offset, 
     D_ptr,D_step,D_offset,D_rows,D_cols, 
     n,-0.5f,1.f); 

    transpose(D_ptr,D_step,D_offset,D_rows,D_cols*sizeof(float), 
       E_ptr,E_step,E_offset); 


    gemm(E_ptr,E_step,E_offset, 
     D_ptr,D_step,D_offset, 
     F_ptr,F_step,F_offset,F_rows,F_cols, 
     n,1.f,0.f); 


}

如果的大小矩陣少於6是完美的工作:)。否則...不是真的。它可使用此代碼檢查：

cv::RNG tutu(std::time(nullptr)); 

int sz = tutu.uniform(1,20); 

cv::Mat_<float> a1(1,sz); 
cv::Mat_<float> a2(a1.size()); 

for(std::size_t i=0;i<sz;i++) 
{ 
    a1(i) = tutu.uniform(0,300); 
    a2(i) = tutu.uniform(300,600); 
} 

cv::Mat_<float> f; 

cv::vconcat(a1,a2,f); 

cv::Mat_<float> c1; 
cv::Mat_<float> c2; 

    cv::Mat_<float> mean; 
// reference 
    cv::calcCovarMatrix(f,c1,mean,cv::COVAR_ROWS | cv::COVAR_NORMAL,CV_32F); 
// check 
test::ocl::CovarMatrix(f,c2); 

std::size_t cnt(0.f); 

for(auto it = c1.begin(),it2 = c2.begin();it != c1.end();it++,it2++) 
    if(*it == *it2) 
     cnt++; 

std::cout<<"check "<<cnt<<" "<<c1.total()<<std::endl;

我仍然在新OpenCL和我interrested知道我做錯了什麼。

有人已經實現了OpenCL內核來處理OpenCV的協方差矩陣嗎？

提前感謝您的幫助。

來源

2015-10-08 Jonny_S

只是猜測：gemm（）函數被設計爲使用預定義的工作組大小，而其他大小隻會給出錯誤的結果。在OpenCV – DarkZeros

Hello中有如此多的函數和宏遵循流程是有點棘手的。我已經檢查並重新檢查參數。此外，我開始複製粘貼opencv函數，使其工作，我專門爲我的情況。在我有了轉置的元素之後。我看輸出。它看起來更像是輸出矩陣的某個部分的多重處理。 –

我辭職，我自己來寫這個代碼：

void ocl_CovarMatrix(cv::Mat_<float>& src,cv::Mat_<float>& covar) 
{ 

    cv::UMat usrc = src.getUMat(cv::ACCESS_READ,cv::USAGE_ALLOCATE_DEVICE_MEMORY); 
    cv::UMat ones = cv::UMat::ones(usrc.rows,usrc.rows,usrc.type()); 
    cv::UMat utmp; 

    double beta = 1.; 
    double alpha = -1./static_cast<double>(usrc.rows); 

    cv::gemm(ones,usrc,alpha,usrc,beta,utmp); 
    cv::gemm(utmp.t(),utmp,beta,cv::noArray(),0.,utmp); 

    utmp.copyTo(covar); 

    ones.release(); 
    utmp.release(); 
    usrc.release(); 
}

我懷疑GPU內存每次更新的功能就是調用，使代碼，而不是慢，如果它已經被寫在一個內核。但它工作效率高。

如果有人有想法，我仍然對另一種解決方案感興趣。

來源

2015-10-12 23:09:50

協方差矩陣與opencl和opencv

回答

相關問題