#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <opencv2\opencv.hpp>
#include <iostream>
#include <string>

using namespace cv;
using namespace std;
  
__global__ void addKernel(uchar **pSrcImg,  uchar* pDstImg, int imgW, int imgH)  
{  
    int
 
 tidx = threadIdx.x + blockDim.x * blockIdx.x;  
    int tidy = threadIdx.y + blockDim.y * blockIdx.y;  
    if (tidx<imgW && tidy<imgH)
    {
        int idx=tidy*imgW+tidx;
        uchar lenaValue=pSrcImg[0][idx];
        uchar moonValue=pSrcImg[1][idx];
        pDstImg[idx]= uchar(0.5
 
*lenaValue+0.5*moonValue);
    }
}  

int main()  
{  
    //OpenCV讀取兩幅影象
    Mat img[2];
    img[0]=imread("data/lena.jpg", 0);
    img[1]=imread("data/moon.jpg", 0);
    int imgH=img[0].rows;
    int imgW=img[0].cols;
    //輸出影象
    Mat dstImg=Mat::zeros(imgH, imgW, CV_8UC1);
    //主機指標
    uchar **pImg=(uchar**)malloc
 
(sizeof(uchar*)*2); //輸入 二級指標

    //裝置指標
    uchar **pDevice;//輸入 二級指標
    uchar *pDeviceData;//輸入 一級指標
    uchar *pDstImgData;//輸出影象對應裝置指標

    //分配GPU記憶體
    cudaError err;
    //目標輸出影象分配GPU記憶體
    err=cudaMalloc(&pDstImgData, imgW*imgH*sizeof(uchar));
    //裝置二級指標分配GPU記憶體
    err=cudaMalloc(&pDevice, sizeof(uchar*)*2);
    //裝置一級指標分配GPU記憶體
    err=cudaMalloc(&pDeviceData, sizeof(uchar)*imgH*imgW*2);
    
    //關鍵：主機二級指標指向裝置一級指標位置，這樣才能使裝置的二級指標指向裝置的一級指標位置
    for (int i=0; i<2; i++)
    {
        pImg[i]=pDeviceData+i*imgW*imgH;
    }

    //拷貝資料到GPU
    //拷貝主機二級指標中的元素到裝置二級指標指向的GPU位置 （這個二級指標中的元素是裝置中一級指標的地址）
    err=cudaMemcpy(pDevice, pImg, sizeof(uchar*)*2, cudaMemcpyHostToDevice);
    //拷貝影象資料(主機一級指標指向主機記憶體) 到 裝置一級指標指向的GPU記憶體中
    err=cudaMemcpy(pDeviceData, img[0].data, sizeof(uchar)*imgH*imgW, cudaMemcpyHostToDevice);
    err=cudaMemcpy(pDeviceData+imgH*imgW, img[1].data, sizeof(uchar)*imgH*imgW, cudaMemcpyHostToDevice);

    //核函式實現lena圖和moon圖的簡單加權和
    dim3 block(8, 8);
    dim3 grid( (imgW+block.x-1)/block.x, (imgH+block.y-1)/block.y);
    addKernel<<<grid, block>>>(pDevice, pDstImgData, imgW, imgH);
    cudaThreadSynchronize();

    //拷貝輸出影象資料至主機，並寫入到本地
    err=cudaMemcpy(dstImg.data, pDstImgData, imgW*imgH*sizeof(uchar), cudaMemcpyDeviceToHost);
    imwrite("data/synThsis.jpg", dstImg);
}