这篇其实是前文 CUDA版Grabcut的实现 的后续，和上文一样，先放视频。
(博客园好像不支持视频，gif文件太大，视频链接）
在上文用CUDA实现opencv下的grabcut后，当时问题主要是最后需要mincut需要上千次push-relabel才能得到满意结果，后改为种子点方式，不到100次就可以得到满意结果，但是种子点需要自己来画，不是很方便，因此，引入深度神经网络先用单桢计算种子点，然后根据这些确认的种子点来计算GMM，如视频这样，以很小代价成功处理1080P下的数据（上面开了录制视频与Unity编辑器占用了大量CPU）。
和前文一样，先说其中大体的实现过程。
1 对于实现yolov3的框架darknet的改造，现在darknet只能传入CPU数据进行计算，但是我们框架内部在读取摄像头后就直接丢到GPU上了，原始的YUV转RGBA都算好在GPU上，如果丢CPU数据，就要先download,然后再upload，这个操作太浪费时间并且是没必要的，我们需要改造这些接口。
2 根据一桢计算一个满意的分割图，然后根据分割图算GMM，然后根据这个GMM算后续图的mask(前景与背景变化不大的情况下), 其中grabcut的mincut里需要迭代多次push-relabel，因此这里给的分辨率需要先降采样，那么算的结果mask并不是太清晰，需要整合 CUDA加opencv复现导向滤波算法 我们才能得到满意的清晰边缘。
3 为了更高效率，我们只考虑显存直接对接方案，相比于CUDA提供的与DX11交互的例子，直接把引擎UE4/Unity3D里的Dx11上下文拿过来，然后把我们处理完后的数据给对应的DX11中，有个很大问题，摄像头30FPS就够了，游戏这块我们需要60FPS向上，在这过程中，游戏引擎的渲染线程被CUDA处理数据线程占用导致游戏会卡，就想CUDA处理线程和游戏引擎的渲染线程都去上厕所，CUDA处理线程占着厕所，游戏引擎的渲染线程就只能在外面等着，而我们要的是，渲染线程看厕所被占了，就马上去处理事情，等过会再来看看厕所是否空着了。
4 cuda端引用的dll过多，原来一些基本的蓝绿幕/深度扣像在DX11都有实现，现在的CUDA端要做成一个能方便接入的系统，提供一些高端功能，如神经网络检查是否有人在镜头内，镜头内的方位等，并且没有也不影响原来项目，类似插件模块。
相对上文的一些改变，中间因为显卡需要适配2070等20系最新显卡的原因，原来的openc4.0alpha(含cuda模块)+cuda8不能使用，于是重新编译了下opencv4.1(含cuda模块)+cuda10.1，中间有些坑，以后如果有时间就记录下，换了后也有个不好的地方，如果不是比较新的N卡驱动，就加载不了现在的cuda模块。考虑性价比，用的是yolov3-tiny框架，每桢中占比不到5ms,检测效果还行。
首先针对第一点，我们需要改进的是能直接传入的GPU方法，首先达到类似如下效果。
cv::cuda::resize(mat, netFrame, cv::Size(netWidth, netHeight), 0, 0, cv::INTER_LINEAR); image2netData_gpu(netFrame, netInput); network_predict_gpudata(net, netInput); detection *dets = get_network_boxes(net, netWidth, netHeight, cthresh, 0, 0, 1, &nboxes);
我们传入的GpuMat,直接在显卡完成缩放，转化成特征图，然后把数据丢给神经网络去计算，详细实现代码如下。
inline __global__ void image2netData(PtrStepSz<uchar4> source, float* outData, int size) { //rgbargbargba->rrrgggbbb const int idx = blockDim.x * blockIdx.x + threadIdx.x; const int idy = blockDim.y * blockIdx.y + threadIdx.y; if (idx < source.cols && idy < source.rows) { float4 color = rgbauchar42float4(source(idy, idx)); int nindex = idy * source.cols + idx; outData[nindex] = color.x; outData[size + nindex] = color.y; outData[2 * size + nindex] = color.z; } } void network_predict_gpudata(network * net, float *input) { network orig = *net; cudaMemcpy(net->input_gpu, input, net->w*net->h * sizeof(float) * 3, cudaMemcpyDeviceToDevice); net->truth = 0; net->train = 0; net->delta = 0; forward_network_gpudata(net); *net = orig; } void forward_network_gpudata(network *netp) { network net = *netp; //cuda_set_device(net.gpu_index); int i; for (i = 0; i < net.n; ++i) { net.index = i; layer l = net.layers[i]; if (l.delta_gpu) { fill_gpu(l.outputs * l.batch, 0, l.delta_gpu, 1); } l.forward_gpu(l, net); net.input_gpu = l.output_gpu; net.input = l.output; if (l.truth) { net.truth_gpu = l.output_gpu; net.truth = l.output; } } pull_network_output(netp); calc_network_cost(netp); }
这个改造比我想像中简单的多，毕竟C框架，实现都在明面上。
然后我们要根据我们设想，先根据单桢计算GMM，就是视频中点击计算GMM的button过程，如下是实现过程。
inline __global__ void setMask(PtrStepSz<uchar> source, PtrStepSz<uchar> mask, int radius, cv::Rect rect) { const int idx = blockDim.x * blockIdx.x + threadIdx.x; const int idy = blockDim.y * blockIdx.y + threadIdx.y; if (idx < source.cols && idy < source.rows) { int mmask = 0; //前景为255 if (idx > rect.x && idx < rect.x + rect.width && idy > rect.y && idy < rect.y + rect.height) { mmask = mask(idy, idx); int smask = source(idy, idx); if (smask == 3) mmask = 1; } mask(idy, idx) = mmask; } } inline __global__ void drawRect(PtrStepSz<uchar4> source, int xmin, int xmax, int ymin, int ymax, int radius, uchar4 drawColor) { const int idx = blockDim.x * blockIdx.x + threadIdx.x; const int idy = blockDim.y * blockIdx.y + threadIdx.y; if (idx < source.cols && idy < source.rows) { int4 xx = make_int4(idx, xmax, idy, ymax); int4 yy = make_int4(xmin, idx, ymin, idy); int4 xy = abs(xx - yy); //只要有一个条件满足就行（分别代表左边，右边，上边，下边） int sum = (xy.x < radius) + (xy.y < radius) + (xy.z < radius) + (xy.w < radius); float2 lr = make_float2(xy.x + xy.y, xy.z + xy.w); float2 rl = make_float2(xmax - xmin, ymax - ymin); if (sum > 0 && length(lr - rl) < radius) { source(idy, idx) = drawColor; } } } #if MRDNN //深度学习下自动计算一桢 vector<PersonBox> boxs; int num = detect->checkImage(grabCutTempFrame, 0.3f, boxs); if (num > 0) { GpuMat halfCutTemp = GpuMat(grabCutHeight, grabCutWidth, CV_8UC4); GpuMat halfTemp = GpuMat(grabCutHeight, grabCutWidth, CV_8UC3); GpuMat resultMask = GpuMat(grabCutHeight, grabCutWidth, CV_8UC1); cv::Mat cpuResult, cpuHalfCutTemp; cv::cuda::resize(grabCutTempFrame, halfCutTemp, cv::Size(grabCutWidth, grabCutHeight), 0, 0, cv::INTER_LINEAR, cvCudaStream); //自己实现，避免引入opencv_cudafilters410/opencv_cudaarithm410，加起来快七百M rgba2bgr_gpu(halfCutTemp, halfTemp, cudaStream); halfTemp.download(cpuHalfCutTemp, cvCudaStream); cudaStreamSynchronize(cudaStream); //选择一个边框 cv::Rect rectangle; rectangle.width = boxs[0].width*grabCutWidth; rectangle.height = boxs[0].height*grabCutHeight; rectangle.x = boxs[0].centerX*grabCutWidth - rectangle.width / 2; rectangle.y = boxs[0].centerY*grabCutHeight - rectangle.height / 2; //用opencv的方式来计算mask,因为opencv本身用八向，切割的边缘更清晰 cv::Mat bgModel, fgModel; cv::grabCut(cpuHalfCutTemp, cpuResult, rectangle, bgModel, fgModel, 3, cv::GC_INIT_WITH_RECT); //显示为3的就是GC_PR_FGD， resultMask.upload(cpuResult, cvCudaStream); //根据resultMask与rectangle构建grabCut->mask setMask_gpu(resultMask, grabCut->mask, 1, rectangle, cudaStream); //显示边框 uchar4 drawColor = make_uchar4(128, 0, 0, 255); cv::Rect rectangle2; rectangle2.width = boxs[0].width*width; rectangle2.height = boxs[0].height*height; rectangle2.x = boxs[0].centerX*width - rectangle2.width / 2; rectangle2.y = boxs[0].centerY*height - rectangle2.height / 2; drawRect_gpu(grabCutRectFrame, rectangle2, 3, drawColor, cudaStream); //GpuMat的析构函数会调用，可以不用手动调用 release halfCutTemp.release(); halfTemp.release(); resultMask.release(); } #endif
里面用的opencv本身提供的grabcut来算mask,因为单桢他用的八方向的，边缘清晰度更好，分割效果更好，然后用们上一篇的实现根据他的mask来计算GMM的值，然后每桢用这个GMM与push-relabel算的一张低分辨率下的mask,正好用来做导向滤波的引导图，经过CUDA加opencv复现导向滤波算法 的处理，就能得到边缘清晰的扣图效果。
其中算k-means/GMM的这个过程可以看上文，和前文不一样的是，原来k-means计算整张图所有点，这个k-means只计算其中确定的前景与背景，需要针对上文的updateCluster修改如下，让计算过程可以针对种子点模式。
//把source所有收集到一块gridDim.x*gridDim.y块数据上。 template<int blockx, int blocky, bool bSeed> __global__ void updateCluster(PtrStepSz<uchar4> source, PtrStepSz<uchar> clusterIndex, PtrStepSz<uchar> mask, float4* kencter, int* kindexs) { __shared__ float3 centers[blockx*blocky][CUDA_GRABCUT_K2]; __shared__ int indexs[blockx*blocky][CUDA_GRABCUT_K2]; const int idx = blockDim.x * blockIdx.x + threadIdx.x; const int idy = blockDim.y * blockIdx.y + threadIdx.y; const int threadId = threadIdx.x + threadIdx.y * blockDim.x; #pragma unroll CUDA_GRABCUT_K2 for (int i = 0; i < CUDA_GRABCUT_K2; i++) { centers[threadId][i] = make_float3(0.f); indexs[threadId][i] = 0; } __syncthreads(); if (idx < source.cols && idy < source.rows) { //所有值都放入共享centers int index = clusterIndex(idy, idx); uchar umask = mask(idy, idx); if (!bSeed || (bSeed && checkObvious(umask))) { bool bFg = checkFg(umask); int kindex = bFg ? index : (index + CUDA_GRABCUT_K); float4 color = rgbauchar42float4(source(idy, idx)); centers[threadId][kindex] = make_float3(color); indexs[threadId][kindex] = 1; } __syncthreads(); //每个线程块进行二分聚合,每次操作都保存到前一半数组里，直到最后保存在线程块里第一个线程上(这块比较费时,0.1ms) for (uint stride = blockDim.x*blockDim.y / 2; stride > 0; stride >>= 1) { //int tid = (threadId&(stride - 1)); if (threadId < stride)//stride 2^n { #pragma unroll CUDA_GRABCUT_K2 for (int i = 0; i < CUDA_GRABCUT_K2; i++) { centers[threadId][i] += centers[threadId + stride][i]; indexs[threadId][i] += indexs[threadId + stride][i]; } } //if (stride > 32) __syncthreads(); } //每块的第一个线程集合，把共享centers存入临时显存块上kencter if (threadIdx.x == 0 && threadIdx.y == 0) { int blockId = blockIdx.x + blockIdx.y * gridDim.x; #pragma unroll CUDA_GRABCUT_K2 for (int i = 0; i < CUDA_GRABCUT_K2; i++) { int id = blockId * 2 * CUDA_GRABCUT_K + i; kencter[id] = make_float4(centers[0][i], 0.f); kindexs[id] = indexs[0][i]; } } } }
其中bSeed为true表明是计算种子点模式，类似还需要修改updateGMM方法实现，种子点模式下，可以增加手动选择前景与背景，如graphcut类似。
这样，主要实现就算完成了，接下来，我们整合进引擎，公司主要使用UE4，我一般来说测试功能更喜欢用Unity3D,对于我测试功能来说，UI与逻辑编写方便与编译快是我需要的，所以主要以Unity3D讲解，后面会简单说下与UE4的对接，其实对接部分，unity3d要编写原生插件才能拿到Unity3D的DX11上下文，以及要写C++/C#的转接层，对于对接部分来说，UE4明显编写更舒服。
就如前面所说，要让CUDA处理数据的线程不占着引擎渲染线程，需要的是一个能在CUDA与DX11不同线程安全访问的显存设计，遗憾的是我没找到(有知道CUDA与DX11有类似DX11之间共享显存能有锁访问类似设计的告诉一声)，但是在我前文 UE4/Unity3D中同时捕获多高清摄像头的高效插件 里我们知道DX11本身之间有共享显存锁访问机制，根据这个，我重新整理了下思路，先声明一个独立的DX11上下文，在这个上下文中有共享纹理，这个DX11的共享纹理与CUDA本身绑定，并且CUDA在处理数据线程中占着这个DX11上下文，等CUDA处理完线程就丢给这个DX11的共享纹理，然后再让游戏引擎UE4/Unity3D的渲染线程的DX11上下文去访问这个DX11的共享纹理，这样用一个独立的DX11做中转，也能达到我们需要的效果。
如下是CUDA与DX11交互的主要代码，其中DX11之间用锁共享访问的逻辑请看 UE4/Unity3D中同时捕获多高清摄像头的高效插件这个链接里，相关代码在这不贴了，这样就避免在CUDA中处理完后还download到内存中，然后在游戏中把内存数据upload在对应引擎的显存上，这样太浪费并且低效。
//一块CUDA与DX11绑定的资源 struct Dx11CudaResource { //cuda资源 cudaGraphicsResource *cudaResource = nullptr; //cuda具体资源 cudaArray *cuArray = nullptr; //对应DX11纹理 ID3D11Texture2D *texture = nullptr; //对应CPU数据 uint8_t *cpuData = nullptr; }; //表明一个DX11上下文的共享显存 struct MGDLL_EXPORT GpuSharedResource { public: ID3D11Texture2D* texture = nullptr; HANDLE sharedHandle = nullptr; bool bGpuUpdate = false; public: //申请一块共享显存,mc表示一个Dx11上下文章 bool restart(class MainCompute* mc, int width, int height, DXGI_FORMAT format = DXGI_FORMAT_R8G8B8A8_UNORM); void release(); }; bool GpuSharedResource::restart(MainCompute* mc, int width, int height, DXGI_FORMAT format) { SafeRelease(texture); bool result = mc->CreateTextureBuffer(nullptr, width, height, format, &texture, false, true); sharedHandle = MainCompute::GetSharedHandle(texture); bGpuUpdate = false; return result; } void GpuSharedResource::release() { SafeRelease(texture); } //Dx11CudaResource 表示Cuda与Dx11联接二方的资源 //GpuSharedResource 表示共享纹理与句柄，MainCompute为对应DX11上下文包装 inline bool registerCudaResource(Dx11CudaResource& cudaDx11, GpuSharedResource& sharedResource, MainCompute* mc, int width, int height) { if (cudaDx11.cudaResource != nullptr && cudaDx11.texture != nullptr) { cudaGraphicsUnregisterResource(cudaDx11.cudaResource); cudaDx11.cudaResource = nullptr; } bool bInit = sharedResource.restart(mc, width, height); if (bInit) { cudaDx11.texture = sharedResource.texture; auto result = cudaGraphicsD3D11RegisterResource(&cudaDx11.cudaResource, cudaDx11.texture, cudaGraphicsRegisterFlagsNone); if (result != cudaSuccess) { LogMessage(warn, "cudaGraphicsD3D11RegisterResource fails."); } } return bInit; } //把CUDA资源复制给DX11纹理 inline void gpuMat2D3dTexture(cv::cuda::GpuMat frame, Dx11CudaResource& cudaResource, cudaStream_t stream) { if (cudaResource.texture != nullptr) { //cuda map dx11,资源数组间map cudaGraphicsMapResources(1, &cudaResource.cudaResource, stream); //map单个资源 cudaGraphicsSubResourceGetMappedArray(&cudaResource.cuArray, cudaResource.cudaResource, 0, 0); //从cuda显存里把数据复制给dx11注册的cuda显存里 cudaMemcpy2DToArray(cudaResource.cuArray, 0, 0, frame.ptr(), frame.step, frame.cols * sizeof(int32_t), frame.rows, cudaMemcpyDeviceToDevice); //cuda unmap dx11 cudaGraphicsUnmapResources(1, &cudaResource.cudaResource, stream); } } //复制GPU处理后的扣像线程给共享显存，通过锁访问，避免和引擎的游戏线程一起访问这块显存 if (dataType & KeyingGPU) { CComPtr<IDXGIKeyedMutex> pDX11Mutex = nullptr; auto hResult = keyingRc.texture->QueryInterface(__uuidof(IDXGIKeyedMutex), (LPVOID*)&pDX11Mutex); DWORD dresult = pDX11Mutex->AcquireSync(0, 0); if (dresult == WAIT_OBJECT_0) { gpuMat2Texture(keying, resultOpFrame); } dresult = pDX11Mutex->ReleaseSync(1); keyingRc.bGpuUpdate = true; } //dx11共享显存把结果丢给游戏引擎的纹理 //d3ddevice 游戏引擎的dx11上下文 //texture 游戏引擎的纹理 //sharedHandle CUDA那边的游戏引擎dx11上下文里共享纹理的句柄 void copySharedToTexture(ID3D11Device * d3ddevice, HANDLE & sharedHandle, ID3D11Texture2D * texture) { if (!d3ddevice) return; CComPtr<ID3D11DeviceContext> d3dcontext = nullptr; //ID3D11DeviceContext* d3dcontext = nullptr; d3ddevice->GetImmediateContext(&d3dcontext); if (!d3dcontext) return; if (sharedHandle && texture) { CComPtr<ID3D11Texture2D> pBuffer = nullptr; HRESULT hr = d3ddevice->OpenSharedResource(sharedHandle, __uuidof(ID3D11Texture2D), (void**)(&pBuffer)); if (FAILED(hr) || pBuffer == nullptr) { LogMessage(error, "open shared texture error."); return; } CComPtr<IDXGIKeyedMutex> pDX11Mutex = nullptr; auto hResult = pBuffer->QueryInterface(__uuidof(IDXGIKeyedMutex), (LPVOID*)&pDX11Mutex); if (FAILED(hResult) || pDX11Mutex == nullptr) { LogMessage(error, "get IDXGIKeyedMutex failed."); return; } DWORD result = pDX11Mutex->AcquireSync(1, 0); if (result == WAIT_OBJECT_0 && pBuffer) { d3dcontext->CopyResource(texture, pBuffer); } result = pDX11Mutex->ReleaseSync(0); } }
让照上文这个顺序，先声明一个独立的DX11上下文，里面生成对应共享纹理，CUDA先绑定这边资源与这个DX11的共享纹理，等CUDA处理完数据，然后把数据复制到这个绑定的DX11共享资源上，这一步，通过锁知道游戏引擎的渲染线程是否在访问，如果在访问，就放弃，继续向下执行，而游戏引擎也通过锁安全访问这个CUDA的计算结果，然后带给UE4呈现，不会给游戏引擎本身带来延迟。注意显存泄漏的问题，Gpumat我们不需要多管，但是这里我们自己声明的，需要注意释放，注册CUDA与DX11资源后面再次注册是要先取消，否则显存泄漏。
现在越来越多数据密集计算，大部分适合GPU算，现在显卡的显存也越来越大，完全可把数据加载到显存，不同的流计算不同的部分，最后统一计算呈现效果，把CPU从数据密集处理中脱离出来，专注逻辑处理。
整个完整过程就差不多了，最后就是CUDA模块与相关引用的dll确实过大，可能有些情况原来的DX11模块就完全能处理，并且A卡暂时还用不了CUDA，所以这块设计成一个插件模块，所有的DLL放入一个文件夹中，如果不要，直接去掉这个文件夹，相关功能自动使用DX11处理，如果检测到有这个文件夹，就先检测是否能成功加载，然后确认是否使用CUDA模块来处理这个功能。
整个实现非常简单，对于这个需求也完全足够了。如下，先定义插件层与工厂类。
template<typename T> class ImageFactory { public: virtual T *create(int type) = 0; }; template<typename T> class PluginManager { public: static PluginManager<T>& getInstance() { static PluginManager m_instance; return m_instance; }; ~PluginManager() { //Release(); }; private: //static PluginManager<T>* instance; PluginManager() {}; private: std::vector<ImageFactory<T>*> factorys; std::vector<T*> models; public: //注册生产类 void registerFactory(ImageFactory<T>* vpf, int type) { if (type >= factorys.size()) factorys.resize(type + 1); factorys[type] = vpf; }; //产生一个实体 T* createModel(int type) { if (type >= factorys.size()) return nullptr; ImageFactory<T>* factory = factorys[type]; T* model = factory->create(type); models.push_back(model); return model; }; void Release() { for (T* model : models) { SafeDelete(model); } models.empty(); for (ImageFactory<T>* factory : factorys) { SafeDelete(factory); } factorys.empty(); } }; void registerFactory(ImageFactory<VideoProcess>* factory, int type) { PluginManager<VideoProcess>::getInstance().registerFactory(factory, type); }
然后在CUDA层定义相应的实现接口与加载是否满足加载的API。
class CudaVideoProcessFactory :public ImageFactory<VideoProcess> { public: CudaVideoProcessFactory() {}; ~CudaVideoProcessFactory() {}; public: virtual VideoProcess* create(int type) override; }; extern "C" GUDLL_EXPORT void registerFactory(); extern "C" GUDLL_EXPORT bool bCanLoad(); VideoProcess * CudaVideoProcessFactory::create(int type) { VideoProcess * vp = new VideoProcessCuda(); return vp; } void registerFactory() { registerFactory(new CudaVideoProcessFactory(), Cuda); registerFactory(new CudaComputePipeFactory(), Cuda); } bool bCanLoad() { int count = cv::cuda::getCudaEnabledDeviceCount(); if (count > 0) { int deviceId = 0; auto device = cv::cuda::DeviceInfo::DeviceInfo(deviceId); bool compation = device.isCompatible(); if (compation) { //cv::cuda::setDevice(deviceId); } return true; } return false; }
最后是在插件层根据定义的目录结构加载CUDA层。
inline HINSTANCE LoadGPUCompuate() { //HINSTANCE hdll = nullptr; typedef bool(*bCanLoad)(); typedef void(*registerfactory)(); hdll = LoadLibraryEx(L"GPUCompute.dll", nullptr, LOAD_WITH_ALTERED_SEARCH_PATH); if (hdll == nullptr) { //DWORD error_id = GetLastError(); //LogMessage(error, to_string(error_id).c_str()); wchar_t sz[512] = { 0 }; HMODULE ihdll = GetModuleHandle(L"ImageUtility.dll"); ::GetModuleFileName(ihdll, sz, 512); ::PathRemoveFileSpec(sz); ::PathAppend(sz, L"GPUCompute"); SetDllDirectory(sz);//SetCurrentDirectory(sz); hdll = LoadLibraryEx(L"GPUCompute.dll", nullptr, LOAD_WITH_ALTERED_SEARCH_PATH); //error_id = GetLastError(); //LogMessage(error, to_string(error_id).c_str()); } if (hdll) { bCanLoad bcd = (bCanLoad)GetProcAddress(hdll, "bCanLoad"); if (bcd()) { registerfactory rf = (registerfactory)GetProcAddress(hdll, "registerFactory"); if (rf) rf(); } else { FreeLibrary(hdll); hdll = nullptr; } } return hdll; }
这样就能根据需求确定是否要加载CUDA与深度神经网络模块，后面改成caffe这样也是适用的。
在相应的游戏引擎Unity3D/UE4目录下，只需要如上定义DLL的目录就行，如果不需要有深度神经网络与CUDA相关功能，直接去掉这个GPUCompute目录也不会对原来的DX11计算模块有影响。
最后记，现在这个Grabcut效果只能说特定环境下还可以，比如人和背景差别比较大的情况下，如果相近，效果不是太好，如上视频中，衣服的左边和右边边缘扣的差很多，就是因为左边受光照多导致与背景差别小，这个功能暂时只能做个备用项，还需要继续改进，但是从这里面，我们底层框架适配了引擎与深度神经网络，这为后面扩展更多功能打下基础。