深度学习的兴起，使得多线程以及GPU编程逐渐成为算法工程师无法规避的问题。这里主要记录自己的GPU自学历程。
在之前的小节中，我们已经遇到了 __global__ 和 __shared__这两种类型限定符。 前者属于函数类型限定符，后者则属于变量类型限定符。 接下来，我们来来了解一下这两类限定符。
函数类型限定符用来标识函数运行在主机还是设备上，函数由主机还是设备调用。
__global__
__global__修饰的函数为 核函数。运行在设备上；可以由主机调用；可以由计算能力大于3.2的设备调用；必须有void返回类型；调用时必须制定运行参数(<<< >>>)该函数的调用时异步的，即可以不必等候该函数全部完成，便可以在CPU上继续工作；__device__
运行在设备上；只能由设备调用；编译器会内联所有认为合适的__device__修饰的函数；__host__
运行在主机上；只能由主机调用；效果等同于函数不加任何限定符；不能与__global__共同使用， 但可以和__device__联合使用；__noinline__
声明不允许内联__forceinline__
强制编译器内联该函数变量类型限定符用来标识变量在设备上的内存位置。
__device__ (单独使用时)
位于 global memory space生命周期为整个应用期间(即与application同生死)可以被grid内的所有threads读取可以在主机中由以下函数读取 cudaGetSymbolAddress() cudaGetSymbolSize() cudaMemcpyToSymbol()cudaMemcpyFromSymbol()__constant__
可以和 __device__ 联合使用位于 constant memory space生命周期为整个应用期间可以被grid内的所有threads读取可以在主机中由以下函数读取 cudaGetSymbolAddress() cudaGetSymbolSize() cudaMemcpyToSymbol()cudaMemcpyFromSymbol()__shared__
可以和 __device__ 联合使用位于一个Block的shared memory space生命周期为整个Block只能被同一block内的threads读写__managed__
可以和 __device__ 联合使用可以被主机和设备引用，主机或者设备函数可以获取其地址或者读写其值生命周期为整个应用期间__restrict__
该关键字用来对指针进行限制性说明，目的是为了减少指针别名带来的问题。
C99标准中引入了restricted指针，用以缓解C语言中指针二义性的问题。缓解指针二义性问题可用于编译器的代码优化。下面是一个指针二义性的例子：
void foo(const float* a, const float* b, float* c) { c[0] = a[0] * b[0]; c[1] = a[0] * b[0]; c[2] = a[0] * b[0] * a[1]; c[3] = a[0] * a[1]; c[4] = a[0] * b[0]; c[5] = b[0]; ... }在C语言中，指针a, b, 和c可能有二义性(别名)，因而对数组c的写入可能会更改数组a和b的元素的值。这就意味着，为了保证程序的正确性，编译器不能把a[0]和b[0]装载入寄存器，对它们做乘法，然后把结果写入c[0]和c[1]，这是因为有这种可能a[0]和c[0]是同一个地址。故而编译器无法对相同的表达式进行优化。
通过把a, b, c声明为restricted指针，程序员可以断言这些指针实际上没有二义性（这里，所有的指针参数都要被设为restrict）
void foo(const float* __restrict__a, const float* __restrict__ b, float* __restrict__ c)在增加了restrict关键字以后，编译器可以根据需要对代码进行优化：
void foo(const float* __restrict__ a, const float* __restrict__ b, float* __restrict__ c) { float t0 = a[0]; float t1 = b[0]; float t2 = t0 * t2; float t3 = a[1]; c[0] = t2; c[1] = t2; c[4] = t2; c[2] = t2 * t3; c[3] = t0 * t3; c[5] = t1; ... }这样便可以减少访存次数和计算量，而代价是增加寄存器的使用量。考虑到额外的寄存器使用可能会降低occupancy，因此这种优化也可能会带来负面效果。
参考资料
“CUDA Toolkit Documentation B.C ”http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#c-language-extensions