关于CUDA/C++/OpenCV混合编程的笔记

使用CUDA C、C++和OpenCV混合编程的小笔记，随时添加新内容，纯属编程过程中得到的经验，不保证完全正确。

1.内核函数中可以用printf函数，但是不能用cout << endl;

__global__ void print_i(int* d_i)   // /home/psdz/test/cuda_c_test_9/main.cu(16):-1: error: error:
                                    //calling a __host__ function("std::basic_ostream<char, std::char_traits<char> >::operator <<")
                                    //from a __global__ function("print_i") is not allowed
{
    printf("%d", d_i);
    //cout << d_i << endl;
}

__global__ void print_j(char* d_j)
{
    printf(d_j);
}

2.device前缀定义的函数只能在GPU上执行，所以device修饰的函数里面不能调用一般常见的函数；global前缀，CUDA允许能够在CPU，GPU两个设备上运行，但是也不能运行CPU里常见的函数；host前缀修饰的是普通函数，默认缺省，可以调用普通函数。
3.第1条可能是因为在CUDA内核函数中只能用C中最基础的数据类型和语句。
4.关于内核参数的参数，可以用指针的形式传入已经为device上分配好的数组，不能用host上的变量以引用的方式传入内核函数，但可以以复制的方式传入，比如：

width, height是host的变量，d_in, d_out是device上的数组（指针），用cudaMalloc()分配好了空间的。

//复制传参是可以的
__global__ void deal_image1(const uchar* d_in, uchar* d_out, int width, int height);
//引用传参就不行
__global__ void deal_image1(const uchar* d_in, uchar* d_out, int& width, int& height);

5.精确计算二维grid和二维block的thread_id的计算公式。

int tid = (gridDim.x*blockIdx.y+blockIdx.x)*blockDim.x*blockDim.y+threadIdx.y*blockDim.x+threadIdx.x;

6.处理二维数组（如二维矩阵、图像）的专用双层for循环：

M行N列
for (int row = blockDim.y * blockIdx.y + threadIdx.y; row < M; row = row + gridDim.y * blockDim.y) {//x, y are src, u, v are dst. Search in dst. (x, y) = T^(-1)(u, v)
        for (int col = blockDim.x * blockIdx.x + threadIdx.x; col < N; col = col + gridDim.x * blockDim.x) {
            //do_something_here;
            __syncthreads();//block中线程对齐
    }
}

7.传入GPU的大数组一定要确保在内存上是连续的，否则不能传入，比如能用C的二维数组int arr[n][m]就不用

vector<vector<int>> vec(n, vector<int>(m, 0));