西樵网站制作公司,高端的响应式网站建设公司,黑马程序员就业情况,上海做网站的公司名称自定义设备函数
核函数#xff1a;__global__修饰#xff1b;在设备中执行#xff1b;设备函数#xff1a;__device__修饰#xff1b;在设备中执行#xff1b;只能被核函数或其他设备函数调用#xff1b;主机函数#xff1a;__host__修饰#xff08;可省略#xff0…自定义设备函数
核函数__global__修饰在设备中执行设备函数__device__修饰在设备中执行只能被核函数或其他设备函数调用主机函数__host__修饰可省略在主机中执行
#include stdio.h
#include cuda_runtime.h
#include device_launch_parameters.h
#includemath.h
#include malloc.h
#include opencv2/opencv.hpp#include stdlib.h#define BLOCK_SIZE 1void __device__ thread_gpu(unsigned char in, unsigned char* out, int thread)
{in thread ? *out 255 : *out 0;
}//图像卷积 GPU
__global__ void sobel_gpu(unsigned char* in, unsigned char* out, const int Height, const int Width)
{int x blockDim.x * blockIdx.x threadIdx.x;int y blockDim.y blockIdx.y threadIdx.y;int index y * Width x;int Gx 0;int Gy 0;unsigned char x0, x1, x2, x3, x4, x5, x6, x7, x8;if (x0 x(Width-1) y0 y(Height-1)){x0 in[(y - 1)*Width (x - 1)];x1 in[(y - 1)*Width (x)];x2 in[(y - 1)*Width (x 1)];x3 in[(y)*Width (x - 1)];x5 in[(y)*Width (x 1)];x6 in[(y 1)*Width (x - 1)];x7 in[(y 1)*Width (x)];x8 in[(y 1)*Width (x 1)];Gx (x0 2 * x3 x6) - (x2 2 * x5 x8);Gy (x0 2 * x1 x2) - (x6 2 * x7 x8);out[index] (abs(Gx) abs(Gy)) / 2;thread_gpu(out[index], out[index], 80);}
}int main()
{cv::Mat src;src cv::imread(photo16.jpg);cv::Mat grayImg,gaussImg;cv::cvtColor(src, grayImg, cv::COLOR_BGR2GRAY);cv::GaussianBlur(grayImg, gaussImg, cv::Size(3,3), 0, 0, cv::BORDER_DEFAULT);int height src.rows;int width src.cols;//输出图像cv::Mat dst_gpu(height, width, CV_8UC1, cv::Scalar(0));//GPU存储空间int memsize height * width * sizeof(unsigned char);//输入 输出unsigned char* in_gpu;unsigned char* out_gpu;cudaMalloc((void**)in_gpu, memsize);cudaMalloc((void**)out_gpu, memsize);dim3 threadsPreBlock(BLOCK_SIZE, BLOCK_SIZE);dim3 blocksPreGrid((width threadsPreBlock.x - 1)/threadsPreBlock.x, (height threadsPreBlock.y - 1)/threadsPreBlock.y);cudaMemcpy(in_gpu, gaussImg.data, memsize, cudaMemcpyHostToDevice);sobel_gpu blocksPreGrid, threadsPreBlock (in_gpu, out_gpu, height, width);cudaMemcpy(dst_gpu.data, out_gpu, memsize, cudaMemcpyDeviceToHost);cv::imwrite(dst_gpu_save.png, dst_gpu);//cv::namedWindow(src, cv::WINDOW_NORMAL);cv::imshow(src, src);cv::imshow(dst_gpu, dst_gpu);cv::waitKey();cudaFree(in_gpu);cudaFree(out_gpu);return 0;
}
