CUDA OpenCV C++ фильтр Собеля

#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/imgproc/imgproc.hpp>
 
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
 
#include <cstdlib>
#include <cstdio>
#include <iostream>
#include <string>
#include <fstream>
#include <time.h>
#include <iomanip>
 
using namespace std;
using namespace cv;
 
string fileName = "tulip.jpg";
 
const char *getMatElementDepth(int depth);
const char *getMatElementType(int type);
void printImageInfo(const Mat &image);
 
// Variables declared globally.
__constant__ int ROWS;
__constant__ int COLS;
__constant__ size_t PITCH1;
__constant__ size_t PITCH2;
surface<void, cudaSurfaceType2D> inSurfRef;
surface<void, cudaSurfaceType2D> sobelPoziomy;
surface<void, cudaSurfaceType2D> sobelPionowy;
surface<void, cudaSurfaceType2D> operatorSobela;
surface<void, cudaSurfaceType2D> in_surf_ref;
surface<void, cudaSurfaceType2D> out_vertical_surf_ref;
surface<void, cudaSurfaceType2D> out_horizontal_surf_ref;
surface<void, cudaSurfaceType2D> out_average_surf_ref;
 
__device__ __forceinline__ int absValue(int value)
{
    return (abs(value) > 255) ? 255 : abs(value);
}
 
__global__ void shiftChannels()
{
    int x = blockIdx.x*blockDim.x + threadIdx.x;
    int y = blockIdx.y*blockDim.y + threadIdx.y;
 
    if ((x < COLS) && (y < ROWS))
    {
        uchar4 pixel;
        surf2Dread(&pixel, in_surf_ref, 4 * x, y);
 
        pixel.x = pixel.z;
        pixel.y = pixel.y;
        pixel.z = pixel.z;
 
        surf2Dwrite(pixel, out_vertical_surf_ref, 4 * x, y);
        surf2Dwrite(pixel, out_horizontal_surf_ref, 4 * x, y);
        surf2Dwrite(pixel, out_average_surf_ref, 4 * x, y);
    }
}
__global__ void sobelKernel(int *sobel1, int *sobel2, int size)
{
    int x = blockIdx.x*blockDim.x + threadIdx.x;
    int y = blockIdx.y*blockDim.y + threadIdx.y;
 
    int i, j;
    int growth = 0;
    uchar4 nextPixele[9]; 
    if ((x>(size / 2) - 1) && (y>(size / 2) - 1) && (x < COLS - (size / 2)) && (y < ROWS - (size / 2)))
    {
        uchar4 pixel1, pixel2, pixel3;
        surf2Dread(&pixel1, inSurfRef, 4 * x, y);
        pixel2 = pixel1;
        pixel3 = pixel1;
 
        int *row_sobel1 = (int *)((char *)sobel1 + y*PITCH1);
        int *row_sobel2 = (int *)((char *)sobel2 + y*PITCH2);
 
        
        for (i = -(size / 2); i <= (size / 2); i++)
        {
            for (j = -(size / 2); j <= (size / 2); j++)
            {
                surf2Dread(&nextPixele[growth], inSurfRef, 4 * (x + i), (y + j));
                growth++;
            }
        }
 
        for (i = 0; i < 9; i++)
        {
            pixel1.x += (nextPixele[i].x * sobel1[i]);
            pixel1.y += (nextPixele[i].y * sobel1[i]);
            pixel1.z += (nextPixele[i].z * sobel1[i]);
            pixel1.w += (nextPixele[i].w * sobel1[i]);
 
            pixel2.x += (nextPixele[i].x * sobel2[i]);
            pixel2.y += (nextPixele[i].y * sobel2[i]);
            pixel2.z += (nextPixele[i].z * sobel2[i]);
            pixel2.w += (nextPixele[i].w * sobel2[i]);
        }
 
        pixel1.x = absValue(pixel1.x);
        pixel1.y = absValue(pixel1.y);
        pixel1.z = absValue(pixel1.z);
        pixel1.w = absValue(pixel1.w);
 
        pixel2.x = absValue(pixel2.x);
        pixel2.y = absValue(pixel2.y);
        pixel2.z = absValue(pixel2.z);
        pixel2.w = absValue(pixel2.w);
 
        surf2Dwrite(pixel1, sobelPoziomy, 4 * x, y);
        surf2Dwrite(pixel2, sobelPionowy, 4 * x, y);
    }
}
Mat image, pic, greyscale; 
Mat picafterSobel1, picafterSobel2;
 
int horizmask[3][3] = {
    { -3,0,3 },
    { -10,0,10 },
    { -3,0,3 } };
 
int verticalmask[3][3] = {
    { -3,-10,-3 },
    { 0,0,0 },
    {3,10,3 } };
 
int widthFiltr = 3;
int main()
{
 
    srand(time(NULL));
    cout << "start" << endl;
 
    Mat image_original = imread(fileName, CV_LOAD_IMAGE_COLOR);
    Mat image_grey;
    Mat image_vertical, image_horizontal, image_average;
 
    if (image_original.empty())
    {
        std::cout << "Image " << fileName << " not found" << std::endl;
        system("pause");
        return -1;
    }
 
    cv::cvtColor(image_original, image_grey, CV_BGR2GRAY);
    cv::cvtColor(image_original, image_vertical, CV_BGR2GRAY);
    cv::cvtColor(image_original, image_horizontal, CV_BGR2GRAY);
    cv::cvtColor(image_original, image_average, CV_BGR2GRAY);
 
    size_t image_bytes = image_grey.total()*image_grey.elemSize();
 
    cout << "Info about the image converted to GREYSCALE:" << endl;
    printImageInfo(image_grey);
 
    cudaError_t err = cudaSuccess;
    dim3 block_dim(16, 16);
    dim3 grid_dim1((image_grey.cols + block_dim.x - 1) / block_dim.x, (image_grey.rows + block_dim.y - 1) / block_dim.y);
 
    cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc(8, 8, 8, 8,
        cudaChannelFormatKindUnsigned);
 
    cudaArray *d_in_array = NULL;
    cudaArray *d_out_horizontal_array = NULL;
    cudaArray *d_out_vertical_array = NULL;
    cudaArray *d_out_average_array = NULL;
 
    image = imread("tulip.jpg");
    cvtColor(image, pic, CV_BGR2RGBA);
    pic.copyTo(picafterSobel1);
    pic.copyTo(picafterSobel2);
    size_t image_bytes1 = pic.total()*pic.elemSize();
    //cvtColor(pic, greyscale, CV_BGR2GRAY);
 
    //cudaError_t err = cudaSuccess;
 
    dim3 grid_dim((pic.cols + block_dim.x - 1) / block_dim.x, (pic.rows + block_dim.y - 1) / block_dim.y);
 
 
    cudaArray *d_pic1 = NULL;
    cudaArray *d_pic2 = NULL;
    cudaArray *d_pic3 = NULL;
 
    int *d_filtr1 = NULL;
    int *d_filtr2 = NULL;
    size_t pitch1;
    size_t pitch2;
    size_t row_bytes = widthFiltr * sizeof(int);
 
    //alokacja pamieci na obrazki w kernelu
 
    err = cudaMallocArray(&d_in_array, &channel_desc, image_grey.cols,
        image_grey.rows, cudaArraySurfaceLoadStore);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaMallocArray(&d_out_horizontal_array, &channel_desc, image_grey.cols,
        image_grey.rows, cudaArraySurfaceLoadStore);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaMallocArray(&d_out_vertical_array, &channel_desc, image_grey.cols,
        image_grey.rows, cudaArraySurfaceLoadStore);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaMallocArray(&d_out_average_array, &channel_desc, image_grey.cols,
        image_grey.rows, cudaArraySurfaceLoadStore);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaMemcpyToArray(d_in_array, 0, 0, image_grey.data,
        image_bytes, cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaMemcpyToSymbol(ROWS, &image_grey.rows, sizeof(int));
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaMemcpyToSymbol(COLS, &image_grey.cols, sizeof(int));
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaBindSurfaceToArray(in_surf_ref, d_in_array);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaBindSurfaceToArray(out_vertical_surf_ref, d_out_vertical_array);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaBindSurfaceToArray(out_horizontal_surf_ref, d_out_horizontal_array);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaBindSurfaceToArray(out_average_surf_ref, d_out_average_array);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    shiftChannels << <grid_dim1, block_dim >> >();
 
    err = cudaGetLastError();
    if (err != cudaSuccess) {
        fprintf(stderr, "shiftChannels launch failed: %s\n", cudaGetErrorString(err));
    }
 
    err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaMemcpyFromArray(image_vertical.data, d_out_vertical_array, 0, 0,
        image_bytes, cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaMemcpyFromArray(image_horizontal.data, d_out_horizontal_array, 0, 0,
        image_bytes, cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
 
    err = cudaMemcpyFromArray(image_average.data, d_out_average_array, 0, 0,
        image_bytes, cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaGetSymbolAddress failed!");
    }
    err = cudaMallocArray(&d_in_array, &channel_desc, pic.cols, pic.rows, cudaArraySurfaceLoadStore);
    err = cudaMallocArray(&d_pic1, &channel_desc, pic.cols, pic.rows, cudaArraySurfaceLoadStore);
    err = cudaMallocArray(&d_pic2, &channel_desc, pic.cols, pic.rows, cudaArraySurfaceLoadStore);
    err = cudaMallocArray(&d_pic3, &channel_desc, pic.cols, pic.rows, cudaArraySurfaceLoadStore);
 
    err = cudaMallocPitch(&d_filtr1, &pitch1, row_bytes, widthFiltr);
    err = cudaMallocPitch(&d_filtr2, &pitch2, row_bytes, widthFiltr);
    //Copy to GPU
    err = cudaMemcpy2D(d_filtr1, pitch1, horizmask, row_bytes, row_bytes, widthFiltr, cudaMemcpyHostToDevice);
    err = cudaMemcpy2D(d_filtr2, pitch2, verticalmask, row_bytes, row_bytes, widthFiltr, cudaMemcpyHostToDevice);
 
    err = cudaMemcpyToArray(d_in_array, 0, 0, pic.data, image_bytes, cudaMemcpyHostToDevice);
 
    err = cudaMemcpyToSymbol(ROWS, &pic.rows, sizeof(int));
    err = cudaMemcpyToSymbol(COLS, &pic.cols, sizeof(int));
 
    err = cudaBindSurfaceToArray(inSurfRef, d_in_array);
    err = cudaBindSurfaceToArray(sobelPoziomy, d_pic1);
    err = cudaBindSurfaceToArray(sobelPionowy, d_pic2);
 
    
    sobelKernel << <grid_dim, block_dim >> >(d_filtr1, d_filtr2, widthFiltr);
    err = cudaDeviceSynchronize();
    err = cudaMemcpyFromArray(picafterSobel1.data, d_pic1, 0, 0, image_bytes, cudaMemcpyDeviceToHost);
    err = cudaMemcpyFromArray(picafterSobel2.data, d_pic2, 0, 0, image_bytes, cudaMemcpyDeviceToHost);
 
    //cvtColor(image_grey, image_original, CV_RGBA2BGR);
 
    
    //namedWindow("Original", CV_WINDOW_AUTOSIZE);
    //imshow("Original", image_original);
    /*cvtColor(picafterSobel1, picafterSobel1, CV_RGBA2BGR);
    cvtColor(picafterSobel2, picafterSobel2, CV_RGBA2BGR);
    imshow("FiltrSobel1", picafterSobel1);
    imshow("FiltrSobel2", picafterSobel2);*/
    
    imwrite("ORIGINAL.jpg", image);
    imwrite("FiltrSobel1.jpg", picafterSobel1);
    imwrite("FiltrSobel2.jpg", picafterSobel2);
 
    
 
    //namedWindow("Grey", CV_WINDOW_AUTOSIZE);
    //imshow("Grey", image_grey);
 
    //namedWindow("Destination", CV_WINDOW_AUTOSIZE);
    //imshow("Destination", image);
 
    
    imwrite("out_original.jpg", image_original);
    imwrite("out_grey.jpg", image_grey);
    imwrite("out_vertical.jpg", image_vertical);
    system("pause");
    pic.release();
    image.release();
    picafterSobel1.release();
    picafterSobel2.release();
    greyscale.release();
    //cvDestroyAllWindows();
    cudaFreeArray(d_pic1);
    cudaFreeArray(d_pic2);
    cudaFreeArray(d_pic3);
    cudaFree((void **)&d_filtr1);
    cudaFree((void **)&d_filtr2);
    cudaFreeArray(d_in_array);
    cudaFreeArray(d_out_vertical_array);
    cudaFreeArray(d_out_horizontal_array);
    cudaFreeArray(d_out_average_array);
 
 
 
    return 0;
}
 
 
const char *getMatElementDepth(int depth)
{
    switch (depth)
    {
    case CV_8U: return "CV_8U";
    case CV_8S: return "CV_8S";
    case CV_16U: return "CV_16U";
    case CV_16S: return "CV_16S";
    case CV_32S: return "CV_32S";
    case CV_32F: return "CV_32F";
    case CV_64F: return "CV_64F";
    default: return "Unknown OpenCV matrix element depth!";
    }
}
 
const char *getMatElementType(int type)
{
    switch (type)
    {
    case CV_8UC1: return "CV_8UC1";
    case CV_8SC1: return "CV_8SC1";
    case CV_16UC1: return "CV_16UC1";
    case CV_16SC1: return "CV_16SC1";
    case CV_32SC1: return "CV_32SC1";
    case CV_32FC1: return "CV_32FC1";
    case CV_64FC1: return "CV_64FC1";
    case CV_8UC2: return "CV_8UC2";
    case CV_8SC2: return "CV_8SC2";
    case CV_16UC2: return "CV_16UC2";
    case CV_16SC2: return "CV_16SC2";
    case CV_32SC2: return "CV_32SC2";
    case CV_32FC2: return "CV_32FC2";
    case CV_64FC2: return "CV_64FC2";
    case CV_8UC3: return "CV_8UC3";
    case CV_8SC3: return "CV_8SC3";
    case CV_16UC3: return "CV_16UC3";
    case CV_16SC3: return "CV_16SC3";
    case CV_32SC3: return "CV_32SC3";
    case CV_32FC3: return "CV_32FC3";
    case CV_64FC3: return "CV_64FC3";
    case CV_8UC4: return "CV_8UC4";
    case CV_8SC4: return "CV_8SC4";
    case CV_16UC4: return "CV_16UC4";
    case CV_16SC4: return "CV_16SC4";
    case CV_32SC4: return "CV_32SC4";
    case CV_32FC4: return "CV_32FC4";
    case CV_64FC4: return "CV_64FC4";
    default: return "Unknown OpenCV matrix element type!";
    }
}
 
void printImageInfo(const Mat &image)
{
    if ((NULL == image.data) || image.empty())
    {
        cout << "The image is empty!" << endl << endl;
        return;
    }
 
    cout << "Image dimensionality: " << image.dims << endl;
    cout << "Image cols x rows: " << image.cols << " x " << image.rows << endl;
    cout << "The number of image pixels: " << image.total() << endl;
    cout << "Type of each image pixel: " << getMatElementType(image.type()) << endl;
    cout << "Depth of each image pixel: " << getMatElementDepth(image.depth()) << endl;
    cout << "The number of image channels: " << image.channels() << endl;
    cout << "Size of each channel in bytes: " << image.elemSize1() << endl;
    cout << "Size of each image pixel in bytes: " << image.elemSize() << endl;
    cout << "Size of the image in bytes: " << image.total()*image.elemSize() << endl << endl;
}