CUDA瞎琢磨笔记（updated on 20180201）

CUDA瞎琢磨笔记

host, gloable, device

CUDA函数在函数定义以前需要标明前缀，host, gloable, device，用来描述这个函数是在哪里调用的。下面一段说明来自网络资料：

（1）__host__ int foo(int a){}与C或者C++中的foo(int a){}相同，是由CPU调用，由CPU执行的函数 
（2）__global__ int foo(int a){}表示一个内核函数，是一组由GPU执行的并行计算任务，以foo<<>>(a)的形式或者driver API的形式调用。目前__global__函数必须由CPU调用，并将并行计算任务发射到GPU的任务调用单元。随着GPU可编程能力的进一步提高，未来可能可以由GPU调用。
（3）__device__ int foo(int a){}则表示一个由GPU中一个线程调用的函数。由于Tesla架构的GPU允许线程调用函数，因此实际上是将__device__ 函数以__inline形式展开后直接编译到二进制代码中实现的，并不是真正的函数。

在.cpp函数中调用.cu函数

一个很有意思的应用是：如果要写一个入口在cpp文件中，但内核用CUDA加速的函数，该怎么写。（下面的只是本人在写CUDA函数的时候试出来的结论，很有可能不对）
1. 首先，gloabl函数应该是记录实际GPU并行任务的函数，使用它需要用<<< block_size, thread_num>>>来修饰的，例如func<<< block_size, NUM_THREADS>>>(parameters)。问题是不能在cpp文件中<<<>>>有其他的语义，显然不能直接在.cpp文件中直接写gloabl函数。
2. 而host函数正好是由cpp函数调用的，且不需要<<<…, …>>>来修饰，因此，我们可以在.cu文件中先定义一个host的接口函数，这个函数可以在.cpp文件中被调用，调用方法和普通的.cpp函数一样。然后在.cu文件中写一个用于完成实际用途的gloabl函数，然后在host函数中调用这个gloabl函数即可。

template泛型编程

CUDA函数在template的写法上也有一个需要注意的地方，就是在写完带template的CUDA函数定义后，需要把实际这个函数在之后被调用时用到的类型写出来。举例来说：

template <typename D>
__global__ void func_kernal(D data, int size) {
    ....
}
template <typename D>
// 默认前缀是__host__
void func(D* data, int size) {
    func_kernal<D><<<..., ...>>>(data, size);
}
template void func<float>(float*, int);
template void func<double>(double*, int);