标签归档:cuda

CUDA + subPixel + EDGE 边缘检测

子像素的边缘检测算法很多,但使用CUDA进行的不是很多。github上可以找到一个带CUDA的子像素边缘检测算法,但经过运行发现有些小bug,其cpu版本正常,但gpu版本找轮廓时会偶发轮廓被切断的问题。一下就对其进行一些修正。当然也许还有其他bug~~

一.源项目位置

源项目位置https://github.com/CsCsongor/subPixelEdgeDetect 。算法基于论文https://www.ipol.im/pub/art/2017/216/

二.代码分析

通过分析,主要是因为使用CUDA后,导入了线程。而原有CPU算法是单线程的。在涉及到最优前向轮廓或最优后向轮廓的求解时,会因为多线程,导致没有找到最优轮廓,而保留了2个轮廓信息。
就是源算法中cu_chain_edge_points这个函数里的下面这段代码

if (fwd >= 0 && next[from] != fwd && ((alt = prev[fwd]) < 0 || chain(alt, fwd, Ex, Ey, Gx, Gy, rows, cols) < fwd_s)){
	if (next[from] >= 0){			// Remove previous from-x link if one */
		prev[next[from]] = -1;	// Only prev requires explicit reset  */
	}
	next[from] = fwd;					// Set next of from-fwd link          */
	if (alt >= 0){						// Remove alt-fwd link if one
		next[alt] = -1;					// Only next requires explicit reset
	}
	prev[fwd] = from;					// Set prev of from-fwd link
}
if (bck >= 0 && prev[from] != bck && ((alt = next[bck]) < 0 || chain(alt, bck, Ex, Ey, Gx, Gy, rows, cols) > bck_s)){
		if (alt >= 0){					// Remove bck-alt link if one
			prev[alt] = -1;				// Only prev requires explicit reset
		}
		next[bck] = from;				// Set next of bck-from link
		if (prev[from] >= 0){		// Remove previous x-from link if one
			next[prev[from]] = -1; // Only next requires explicit reset
		}
		prev[from] = bck;				// Set prev of bck-from link
}

而原作者好像也发现了这个问题,在拼接所有轮廓的时候,并没有使用cpu版本里的代码去先找每一个轮廓的起始点,然后再拼接轮廓,而是从任意点开始。这样就会导致轮廓被切断了。
以下list_chained_edge_points 函数中的注释掉的那行代码,就是找轮廓起始点的算法。

// Set k to the beginning of the chain, or to i if closed curve
// for (k = i; (n = prev[k]) >= 0 && n != i; k = n); //这句被注释了,替换成了下面这行
if ((n = prev[i]) >=0 && n != i){ k= n;}

三.算法修改

1.修改cu_chain_edge_points函数。 寻找最优5*5轮廓点的算法屏蔽,先记录下所有向前向后轮廓点

// Chain edge points
__global__
void cu_chain_edge_points(int * next, int * prev, double * Ex,	double * Ey,double * Gx, double * Gy, int rows, int cols){
	int x, y, i , j, alt;
	int dx, dy, to;

	x = blockIdx.x*blockDim.x+threadIdx.x+2;
	y = blockIdx.y*blockDim.y+threadIdx.y+2;

	// Try each point to make local chains
	// 2 pixel margin to include the tested neighbors
	if (x < (rows-2) && y < (cols-2)){
		// Must be an edge point
		if (Ex[x + y*rows] >= 0.0 && Ey[x + y*rows] >= 0.0){
			int from = x + y*rows;  // Edge point to be chained
			double fwd_s = 0.0;  	  // Score of best forward chaining
			double bck_s = 0.0;     // Score of best backward chaining
			int fwd = -1;           // Edge point of best forward chaining
			int bck = -1;           // Edge point of best backward chaining

			/* try all neighbors two pixels apart or less.
				looking for candidates for chaining two pixels apart, in most such cases,
				is enough to obtain good chains of edge points that	accurately describes the edge.
			*/
			for (i = -2; i <= 2; i++){
				for (j = -2; j <= 2; j++){
					to = x + i + (y + j)*rows; // Candidate edge point to be chained

					double s = chain(from, to, Ex, Ey, Gx, Gy, rows, cols);  //score from-to

					if (s > fwd_s){ // A better forward chaining found
						fwd_s = s;
						fwd = to;
					} else if (s < bck_s){ // A better backward chaining found
						bck_s = s;
						bck = to;
					}
				}
			}

            if (fwd >= 0){
                next[from] = fwd;					// Set next of from-fwd link
            }
            if (bck >= 0){
                prev[from] = bck;				// Set prev of bck-from link
            }

//			if (fwd >= 0 && next[from] != fwd && ((alt = prev[fwd]) < 0 || chain(alt, fwd, Ex, Ey, Gx, Gy, rows, cols) < fwd_s)){
//				if (next[from] >= 0){			// Remove previous from-x link if one */
//					prev[next[from]] = -1;	// Only prev requires explicit reset  */
//				}
//				next[from] = fwd;					// Set next of from-fwd link          */
//				if (alt >= 0){						// Remove alt-fwd link if one
//					next[alt] = -1;					// Only next requires explicit reset
//				}
//				prev[fwd] = from;					// Set prev of from-fwd link
//			}
//			if (bck >= 0 && prev[from] != bck && ((alt = next[bck]) < 0 || chain(alt, bck, Ex, Ey, Gx, Gy, rows, cols) > bck_s)){
//					if (alt >= 0){					// Remove bck-alt link if one
//						prev[alt] = -1;				// Only prev requires explicit reset
//					}
//					next[bck] = from;				// Set next of bck-from link
//					if (prev[from] >= 0){		// Remove previous x-from link if one
//						next[prev[from]] = -1; // Only next requires explicit reset
//					}
//					prev[from] = bck;				// Set prev of bck-from link
//			}
		}
	}
}

2.增加一个cu_cut_edge_points函数来完成单点最后向前向后轮廓的选择,切掉不是最优的轮廓

__global__
void cu_cut_edge_points(int * next, int * prev, double * Ex,	double * Ey,double * Gx, double * Gy, int rows, int cols){
    int x, y, i , j, alt;
    int dx, dy, to;

    x = blockIdx.x*blockDim.x+threadIdx.x+2;
    y = blockIdx.y*blockDim.y+threadIdx.y+2;

    if (x < (rows-2) && y < (cols-2)){
        if (Ex[x + y*rows] >= 0.0 && Ey[x + y*rows] >= 0.0){
            int center = x + y*rows;
            int temp_next = -1;
            int temp_prev = -1;

            if (x < (rows-2) && y < (cols-2)){
                for (i = -2; i <= 2; i++){
                    for (j = -2; j <= 2; j++){
                        to = x + i + (y + j)*rows;
                        if(next[to] == center){
                            if(temp_next >= 0){
                                if(chain(center, to, Ex, Ey, Gx, Gy, rows, cols) > chain(center, temp_next, Ex, Ey, Gx, Gy, rows, cols)){
                                    next[temp_next] = -1;
                                    temp_next = to;
                                } else {
                                    next[to] = -1;
                                }
                            } else {
                                temp_next = to;
                            }
                        }

                        if(prev[to] == center){
                            if(temp_prev >= 0){
                                if(chain(center, to, Ex, Ey, Gx, Gy, rows, cols) < chain(center, temp_prev, Ex, Ey, Gx, Gy, rows, cols)){
                                    prev[temp_prev] = -1;
                                    temp_prev = to;
                                } else {
                                    prev[to] = -1;
                                }
                            } else {
                                temp_prev = to;
                            }
                        }
                    }
                }
            }
        }
    }
}

3.修改cu_thresholds_remove函数,删掉一些错误的轮廓

__global__
void cu_thresholds_remove(int * next, int * prev,int rows, int cols, int * valid){
	int x, y;

	x = blockIdx.x*blockDim.x+threadIdx.x;
	y = blockIdx.y*blockDim.y+threadIdx.y;
	int idx = x+y*rows;

	if (idx < rows*cols){
		if ((prev[idx] >= 0 || next[idx] >= 0) && !valid[idx]){
			prev[idx] = next[idx] = -1;
		}

        if(prev[idx] != -1 && next[prev[idx]] != idx){
            prev[idx] = -1;
        }
        if(next[idx] != -1 && prev[next[idx]] != idx){
            next[idx] = -1;
        }
	}
}

4.修改list_chained_edge_points函数,重新启用找寻轮廓起始点的代码

// Set k to the beginning of the chain, or to i if closed curve
for (k = i; (n = prev[k]) >= 0 && n != i; k = n);
//if ((n = prev[i]) >=0 && n != i){ k= n;}

5.修改devernay函数,增加刚才添加的函数

cu_chain_edge_points<<<numberOfBlocks, threadsPerBlock, threadsPerBlock.x*threadsPerBlock.y*sizeof(uchar), stream>>>(next, prev, Ex, Ey, Gx, Gy, rows, cols);
cudaDeviceSynchronize();
cu_cut_edge_points<<<numberOfBlocks, threadsPerBlock, threadsPerBlock.x*threadsPerBlock.y*sizeof(uchar), stream>>>(next, prev, Ex, Ey, Gx, Gy, rows, cols);
cudaDeviceSynchronize();
cu_thresholds_with_hysteresis<<<numberOfBlocks, threadsPerBlock, threadsPerBlock.x*threadsPerBlock.y*sizeof(uchar), stream>>>(next, prev, modG, rows, cols, th_high, th_low, valid);
cudaDeviceSynchronize();

完整的代码请下载 subPixelGPU.zip

Opencv Inpaint + CUDA

opencv只有cpu接口的Inpaint函数,对于需要使用CUDA进行图片处理时,反复内存显存迁移数据会影响计算速度。在不考虑填充效果十分好的情况下,可以使用如下CUDA算法,简单的进行填充。

算法来源

算法主要参考了https://github.com/Po-Ting-lin/HairRemoval.git 中的填充代码。该项目主要是去除皮肤上的毛发。对于其如何寻找需要填充的区域就不讨论了,直接使用其分析出的mask图进行填充。

该填充算法对细线和小面积的填充效果还可以 大面积的 效果就很一般了

继续阅读

OPENCV + CUDA 实现小波滤波

opencv中有cpu 和 gpu版本的DFT函数,及傅里叶变换的函数,可以实现dft滤波。但opencv中没有DWT,及小波变换。下面将介绍一下实现的方法。

傅里叶变换与小波变换都能实现滤波,不好说那个更好。但傅里叶变换有个缺点,对于图像处理来说,其在处理图像锐利边缘时,很容易出现边缘抖动的情况。原因如下图:

傅里叶变换只能使用余弦波,很难拟合出突变。具体的解释和小波变换的特点。可以参考知乎上的文章 https://zhuanlan.zhihu.com/p/44215123

代码部分主要参考了GitHub上实现方法,经过一些简单的修改以适应vc++,和高频滤波的需要。原链接https://github.com/pierrepaleo/PDWT

继续阅读

Opencv中使用CUDA原函数

opencv中的cuda模块封装了大部分常用的图像处理函数。但一些函数只提供了8bit图片的接口,没有16bit图片的接口。如果需要处理10bit 12bit或更高big的图片就需要调用CUDA的原型函数了。下面就简单举例使用opencv中的GpuMat调用cuda原函数的方法。

opencv中存储GPU图片的类型为GpuMat,不需要考虑显存的分配和释放,使用起来比较方便。而CUDA使用的是Npp8u* Npp16u*等指针,这里就涉及指针的转换。以nppiLUTPalette_16u_C1R 函数举例。

Mat img = ….. //读取一张12bit图片
cuda::GpuMat src_img; //初始化Gpu
src_img.upload(img); //上传图片到Gpu
Mat lut(1, 4096, CV_16UC1, lineardata.data()); // 导入lut数据
cuda::GpuMat gpu_lut = new cv::cuda::GpuMat(1, 4096, CV_16UC1);
gpu_lut.upload(lut); //上传lut数据到Gpu,lut数据也要放入Gpu中
cuda::GpuMat dst_gpu(src_img.size(), CV_16UC1); //初始化转化后的GpuMat
NppiSize oSizeROI = { src_img.cols, src_img.rows }; //设置ROI 这里是整张图
//使用prt获取GpuMat的指针
NppStatus status1 = nppiLUTPalette_16u_C1R(src_img.ptr(), static_cast(src_img.step), dst_gpu.ptr(), static_cast(dst_gpu.step), oSizeROI, gpu_lut.ptr(),16);

得到的dst_gpu就是经过lut后的图片,后续的使用也不需要考虑显存的分配和释放。

Opencv中使用cuda进行 dft 与 idft滤波运算

opencv源代码中包含了dft的demo,但没有使用cuda的demo。下文会简单给出一个cuda例程,并进行简单的高频滤波。

为方便说明,将程序分成了若干部分
1.头文件

#include <iostream>

#include <opencv2/core.hpp>
#include <opencv2/imgcodecs.hpp>
#include <opencv2/imgproc.hpp>
#include <opencv2/core/cuda.hpp>
#include <opencv2/cudaarithm.hpp>

using namespace cv;

2.导入图片加载到GPU,为简单使用了单色图

继续阅读

tensorflowsharp中使用CUDA

tensorflowsharp使用C#封装了tensorflow的c-api接口,可以方便的在C#中使用tensorflow的模型。但作者只提供了cpu版本的tensorflow,如果要使用gpu版本的tensorflow,就需要自己编译出一个带CUDA的dll。好在网上资料比较多,编译起来并不是很困难

一.编译环境搭建

windows下搭建tensorflow编译环境可以参考官网上的教程,写得很详细,按照上面一步步操作就可以了。链接https://www.tensorflow.org/install/source_windows

继续阅读

OpencvSharp 中使用 cuda

opencvsharp 是 opencv的c#版本,近期有项目使用了opencvsharp来进行图像处理。这个github上星级很高的项目果然是不错的,运行起来比较稳定,没有出现大的问题。但opencvsharp中没有cuda的完整支持,只有最基本的类型支持,无任何算法支持,想用就只能靠自己添加了。作者的解释如下:

大概就是说cuda需要用户自己编译opencv ,没有一个统一版本的dll提供使用,所以就删除了cuda的支持
其实也对,cuda的使用涉及cuda版本,使用的显卡算力等。使用c++版本的opencv时,也是要自己编译的。
但c#上使用就没有办法了吗?还好作者已经打好了基础,提供了GpuMat的支持,并有大量cup版本的函数进行参考。添加起来还是比较容易的。
项目目录 https://github.com/shimat/opencvsharp

继续阅读