我最近开始学习 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 的基金会已经为 gemm 和函数转置实现了一个 OpenCL 内核。
所以我试图从一个直接来自 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);
}
协方差矩阵.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 的新手,我很想知道我做错了什么。
是否有人已经实现了 OpenCL 内核来使用 OpenCV 处理协方差矩阵?
提前感谢您的帮助。
我是一名优秀的程序员,十分优秀!