车道线识别之——增强黄色车道线

为了更好的识别车道线，需要给黄色车道线增强，方法是将RGB图片转换为HSV图片，增强HSV图片的第三位(V)的数值，再转换成黑白图片。这个方法只适用于白天并且车道线清晰的场景。
如下是普通代码和CUDA代码

#include <iostream>
#include <opencv2/opencv.hpp>
#include <opencv2/core.hpp>
#include <cuda.h>
#include <cuda_runtime.h>

using namespace std;
using namespace cv;

void RGB_to_HSV(const cv::Mat img_RGB, cv::Mat &img_HSV){
    int row = img_RGB.rows;
    int col = img_RGB.cols;
    int B, G, R;//Clearly what are they mean
    int V;    //V = max{R, G, B}
    int m;      //m = min{R, G, B}
    int S;    //S = (V - min{R, G, B}) / V, if V > min{R, G, B}
                //S = 0, if R = G = B
                //delta = max{R, G, B} - min{R, G, B}
    int H;      //H = 0, if max{R, G, B} = min{R, G, B}
                //H = (60 * (G - B)) / delta, if max{R, G, B} = R
                //H = 120 + (60 * (B - R)) / delta, if max{R, G, B} = G
                //H = 240 + (60 * (R - G)) / delta, if max{R, G, B} = B

    for (int i = 0; i < row; i = i + 1) {
        for (int j = 0; j < col; j = j + 1) {
            B = img_RGB.at<cv::Vec3b>(i, j)[0];
            G = img_RGB.at<cv::Vec3b>(i, j)[1];
            R = img_RGB.at<cv::Vec3b>(i, j)[2];
            //Confirm value V and m
            if ((B >= G) && (G >= R)) {
                V = B;
                m = R;
            } else if ((B >= G) && (G < R) && (B >= R)) {
                V = B;
                m = G;
            } else if ((B >= G) && (G < R) && (B < R)) {
                V = R;
                m = G;
            } else if ((B < G) && (G < R)) {
                V = R;
                m = B;
            } else if ((B < G) && (G >= R) && (B >= R)) {
                V = G;
                m = R;
            } else {
                V = G;
                m = B;
            }
            //Confirm value S
            if (V > m) {
                S = (int)((V - m) / V);
            } else {
                S = 0;
            }
            //Confirm value H
            if (V == m) {
                H = 0;
            } else if (V == R) {
                H = (int)(60 * (G - B) / (V - m));
            } else if (V == G) {
                H = (int)(120 + 60 * (B - R) / (V - m));
            } else {
                H = (int)(240 + 60 * (R - G) / (V - m));
            }
            //if H < 0, H should +360
            if (H < 0) {
                H = H + 360;
            }

            img_HSV.at<cv::Vec3b>(i, j)[0] = H;
            img_HSV.at<cv::Vec3b>(i, j)[1] = S;
            img_HSV.at<cv::Vec3b>(i, j)[2] = V;

            //cout << i << "  " << j << endl;   //testcode
        }

    }//end double for loop

}

//kernel code of void RGB_to_HSV(const cv::Mat img_RGB, cv::Mat&img_HSV)
__global__ void global_RGB_to_HSV(uchar3* d_image_RGB, uchar3* d_image_HSV ,int height, int width){
    int R, G, B;
    int V;    //V = max{R, G, B}
    int m;      //m = min{R, G, B}
    int S;    //S = (V - min{R, G, B}) / V, if V > min{R, G, B}
                //S = 0, if R = G = B
                //delta = max{R, G, B} - min{R, G, B}
    int H;      //H = 0, if max{R, G, B} = min{R, G, B}
                //H = (60 * (G - B)) / delta, if max{R, G, B} = R
                //H = 120 + (60 * (B - R)) / delta, if max{R, G, B} = G
                //H = 240 + (60 * (R - G)) / delta, if max{R, G, B} = B
    for (int row = blockDim.y * blockIdx.y + threadIdx.y; row < height; row = row + gridDim.y * blockDim.y) {
        for (int col = blockDim.x * blockIdx.x + threadIdx.x; col < width; col = col + gridDim.x * blockDim.x) {
            B = d_image_RGB[row * width + col].x;
            G = d_image_RGB[row * width + col].y;
            R = d_image_RGB[row * width + col].z;
            //Confirm value V and m
            if ((B >= G) && (G >= R)) {
                V = B;
                m = R;
            } else if ((B >= G) && (G < R) && (B >= R)) {
                V = B;
                m = G;
            } else if ((B >= G) && (G < R) && (B < R)) {
                V = R;
                m = G;
            } else if ((B < G) && (G < R)) {
                V = R;
                m = B;
            } else if ((B < G) && (G >= R) && (B >= R)) {
                V = G;
                m = R;
            } else {
                V = G;
                m = B;
            }
            //Confirm value S
            if (V > m) {
                S = (int)((V - m) / V);
            } else {
                S = 0;
            }
            //Confirm value H
            if (V == m) {
                H = 0;
            } else if (V == R) {
                H = (int)(60 * (G - B) / (V - m));
            } else if (V == G) {
                H = (int)(120 + 60 * (B - R) / (V - m));
            } else {
                H = (int)(240 + 60 * (R - G) / (V - m));
            }
            //if H < 0, H should +360
            if (H < 0) {
                H = H + 360;
            }
            d_image_HSV[row * width + col].x = H;
            d_image_HSV[row * width + col].y = S;
            d_image_HSV[row * width + col].z = V;
        }
    }
}

//change(enhance) the V
void enhance_HSV(cv::Mat &image_hsv) {
    //int V;
    int row = image_hsv.rows;
    int col = image_hsv.cols;
    for (int i = 0; i < row; i = i + 1) {
        for (int j = 0; j < col; j = j + 1) {
            if (image_hsv.at<cv::Vec3b>(i, j)[2] >= 205) {
                image_hsv.at<cv::Vec3b>(i, j)[2] = 255;
            } else {
                image_hsv.at<cv::Vec3b>(i, j)[2] = 0;
            }
        }
    }
}

//kernel code of void enhance_HSV(cv::Mat &image_hsv)
//CAUTION: __device__ 函数实际上是以__inline形式展开后直接编译到二进制代码中实现的，并不是真正的函数。
__device__ void device_enhance_HSV(uchar3* d_image_HSV, int height, int width) {
    for (int row = blockDim.y * blockIdx.y + threadIdx.y; row < height; row = row + gridDim.y * blockDim.y) {
        for (int col = blockDim.x * blockIdx.x + threadIdx.x; col < width; col = col + gridDim.x * blockDim.x) {
            if (d_image_HSV[row * width + col].z >= 205) {
                d_image_HSV[row * width + col].z = 255;
            } else {
                d_image_HSV[row * width + col].z = 0;
            }
        }
    }
}

void HSV_to_gray(const cv::Mat image_hsv, cv::Mat &image_gray, cv::Mat &image_gray_) {
    int row = image_hsv.rows;
    int col = image_hsv.cols;
    for (int i = 0; i < row; i = i + 1) {
        for (int j = 0; j < col; j = j + 1) {
            if (image_hsv.at<cv::Vec3b>(i, j)[2] == 255) {
                image_gray.at<uchar>(i, j) = 255;
            } else {
                image_gray.at<uchar>(i, j) = 0;
            }
        }
    }
    for (int i = 1; i < row - 1; i = i + 1) {
        for (int j = 1; j < col - 1; j = j + 1) {
            if (image_gray.at<uchar>(i - 1, j - 1) +
                image_gray.at<uchar>(i - 1, j) +
                image_gray.at<uchar>(i - 1, j + 1) +
                image_gray.at<uchar>(i, j - 1) +
                image_gray.at<uchar>(i, j + 1) +
                image_gray.at<uchar>(i + 1, j - 1) +
                image_gray.at<uchar>(i + 1, j) +
                image_gray.at<uchar>(i + 1, j + 1) > 1000) {
                image_gray_.at<uchar>(i, j) = 255;
            } else {
                image_gray_.at<uchar>(i, j) = 0;
            }
        }
    }
}

//kernel code of void HSV_to_gray(const cv::Mat image_hsv, cv::Mat &image_gray, cv::Mat &image_gray_)
__device__ void device_HSV_to_gray(const uchar3* d_image_HSV, uchar* d_image_gray, uchar* d_image_gray_, int height, int width) {
    for (int row = blockDim.y * blockIdx.y + threadIdx.y; row < height; row = row + gridDim.y * blockDim.y) {
        for (int col = blockDim.x * blockIdx.x + threadIdx.x; col < width; col = col + gridDim.x * blockDim.x) {
            if (d_image_HSV[row * width + col].z == 255) {
                d_image_gray[row * width + col] = 255;
            } else {
                d_image_gray[row * width + col] = 0;
            }
        }
    }
    __syncthreads();
    for (int row = blockDim.y * blockIdx.y + threadIdx.y; row < height; row = row + gridDim.y * blockDim.y) {
        for (int col = blockDim.x * blockIdx.x + threadIdx.x; col < width; col = col + gridDim.x * blockDim.x) {
            if (d_image_gray[(row - 1) * width + col - 1] +
                d_image_gray[(row - 1) * width + col] +
                d_image_gray[(row - 1) * width + col + 1] +
                d_image_gray[row * width + col - 1] +
                d_image_gray[row * width + col + 1] +
                d_image_gray[(row + 1) * width + col - 1] +
                d_image_gray[(row + 1) * width + col] +
                d_image_gray[(row + 1) * width + col + 1]  > 1000) {
                d_image_gray_[row * width + col] = 255;
            } else {
                d_image_gray_[row * width + col] = 0;
            }
        }
    }
}

//Execute device functions, because device cannot be executed in main(), even cannot be configured(source: threads).
__global__ void global_do(uchar3* d_image_RGB, uchar3* d_image_HSV, uchar* d_image_gray, uchar* d_image_gray_, int height, int width) {
    //global_RGB_to_HSV(d_image_RGB, d_image_HSV, height, width);
    device_enhance_HSV(d_image_HSV, height, width);
    device_HSV_to_gray(d_image_HSV, d_image_gray, d_image_gray_, height, width);
}

//a test of  __host__ __device__  together. Build SUCCEED!
//__host__ __device__ void foo(){
//    printf("Hello world!\n");
//}

int main()
{
    cv::Mat image_RGB = cv::imread("../../bgrtohsv/front.jpg");
    int height = image_RGB.rows;
    int width = image_RGB.cols;
    cv::Mat image_HSV(height, width, CV_8UC3);
    cv::Mat image_gray(height, width, CV_8UC1);
    cv::Mat image_gray_(height, width, CV_8UC1);

    size_t size_image_RGB = sizeof(uchar3) * height * width;
    size_t size_image_HSV = sizeof(uchar3) * height * width;
    size_t size_image_gray = sizeof(uchar) * height * width;
    size_t size_image_gray_ = sizeof(uchar) * height * width;

    uchar3* d_image_RGB = NULL;
    uchar3* d_image_HSV = NULL;
    uchar* d_image_gray = NULL;
    uchar* d_image_gray_ = NULL;

    cudaMalloc((void**)&d_image_RGB, size_image_RGB);
    cudaMalloc((void**)&d_image_HSV, size_image_HSV);
    cudaMalloc((void**)&d_image_gray, size_image_gray);
    cudaMalloc((void**)&d_image_gray_, size_image_gray_);
    cudaMemcpy(d_image_RGB, image_RGB.data, size_image_RGB, cudaMemcpyHostToDevice);
    cudaMemcpy(d_image_HSV, image_HSV.data, size_image_HSV, cudaMemcpyHostToDevice);
    cudaMemcpy(d_image_gray, image_gray.data, size_image_gray, cudaMemcpyHostToDevice);
    cudaMemcpy(d_image_gray_, image_gray_.data, size_image_gray_, cudaMemcpyHostToDevice);

    dim3 dimGrid(16, 16, 1);
    dim3 dimBlock(32, 32, 1);

    global_RGB_to_HSV<< <dimGrid, dimBlock>> >(d_image_RGB, d_image_HSV, height, width);
    cudaDeviceSynchronize();
    global_do<< <dimGrid, dimBlock>> >(d_image_RGB, d_image_HSV, d_image_gray, d_image_gray_, height, width);
    cudaDeviceSynchronize();    //wait for ALL

    cudaMemcpy(image_RGB.data, d_image_RGB, size_image_RGB, cudaMemcpyDeviceToHost);
    cudaMemcpy(image_HSV.data, d_image_HSV, size_image_HSV, cudaMemcpyDeviceToHost);
    cudaMemcpy(image_gray.data, d_image_gray, size_image_gray, cudaMemcpyDeviceToHost);
    cudaMemcpy(image_gray_.data, d_image_gray_, size_image_gray_, cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();

//    RGB_to_HSV(image_RGB, image_HSV);
//    enhance_HSV(image_HSV);
//    HSV_to_gray(image_HSV, image_gray, image_gray_);

    cv::imshow("image_HSV", image_HSV);
    cv::imshow("image_gray", image_gray);
    cv::imshow("image_gray_", image_gray_);

    cv::waitKey(0);

    cv::imwrite("../../bgrtohsv/image_HSV.jpg", image_HSV);
    cv::imwrite("../../bgrtohsv/image_gray.jpg", image_gray);
    cv::imwrite("../../bgrtohsv/image_gray_.jpg", image_gray_);

    cudaFree(d_image_RGB);
    cudaFree(d_image_HSV);
    cudaFree(d_image_gray);
    cudaFree(d_image_gray_);

    //foo();
    return 0;
}

原图(image_RGB)：image_HSV：image_gray:image_gray_:最后两张图是有差别的。

总结：
1.使用CUDA处理二维图像的万能公式：

for (int row = blockDim.y * blockIdx.y + threadIdx.y; row < height; row = row + gridDim.y * blockDim.y) {
        for (int col = blockDim.x * blockIdx.x + threadIdx.x; col < width; col = col + gridDim.x * blockDim.x) {
        	do_something();
        }
    }

2.__device__定义的函数不能在main()函数里执行，甚至不能分配计算资源，需要放到__global__中执行，再在main()中执行。这条不对！！！

//Execute device functions, because device cannot be executed in main(), even cannot be configured(source: threads).
__global__ void global_abc(uchar3* d_image_HSV, uchar* d_image_gray, uchar* d_image_gray_, int height, int width) {
    device_enhance_HSV(d_image_HSV, height, width);
    device_HSV_to_gray(d_image_HSV, d_image_gray, d_image_gray_, height, width);
}

CMakeLists.txt

cmake_minimum_required(VERSION 2.8)

project(bgrtohsv)


set(CMAKE_CXX_STANDARD 11)
set(CMAKE_CXX_STANDARD_REQUIRED ON)

include_directories(include
${CUDA_INCLUDE_DIRS}
${OpenCV_INCLUDE_DIRS}
)

link_directories(${OpenCV_LIBRARY_DIRS})

find_package(CUDA REQUIRED)
find_package(OpenCV REQUIRED)

#INCLUDE(/home/psdz/cmake-3.9.0/Modules/FindCUDA.cmake)

FILE(GLOB SOURCES "*.cu" "*.cpp" "*.c" "*.h")

set(CUDA_NVCC_FLAGS "-g -G")

CUDA_ADD_EXECUTABLE(${PROJECT_NAME} main.cu)

target_link_libraries(${PROJECT_NAME} ${OpenCV_LIBS})

#add_executable(${PROJECT_NAME} "main.cpp")