cuda
https://blog.csdn.net/qq_41554005/article/details/119765334
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions
https://blog.csdn.net/Rong_Toa/article/details/78655575
https://blog.csdn.net/nn1997729/article/details/118181359
cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);
cudaMalloc((void **)&d_a, size);
２．１ __device__，__device__表明声明的数据存放在显存中，所有的线程都可以访问，而且主机也可以通过运行时库访问；
２．２ __shared__，__shared__表示数据存放在共享存储器在，只有在所在的块内的线程可以访问，其它块内的线程不能访问；
２．３ __constant__，__constant__表明数据存放在常量存储器中，可以被所有的线程访问，也可以被主机通过运行时库访问；
对于函数部分前缀
1. __device__
使用 _device_ 限定符声明的函数具有以下特征：
n 在设备上执行；
n 仅可通过设备调用。
2. __global__
使用 _global_ 限定符可将函数声明为内核。此类函数：
n 在设备上执行；
n 仅可通过主机调用。
3. __host__
使用 _host_ 限定符声明的函数具有以下特征：
n 在主机上执行；
n 仅可通过主机调用。
仅使用 _host_ 限定符声明函数等同于不使用 _host_、_device_ 或 _global_ 限定符声明函数，这两种情况下，函数都将仅为主机进行编译。
函数前缀的一些限制：
_device_ 和 _global_ 函数不支持递归。
_device_ 和 _global_ 函数的函数体内无法声明静态变量。
_device_ 和 _global_ 函数不得有数量可变的参数。
_device_ 函数的地址无法获取，但支持 _global_ 函数的函数指针。
_global_ 和 _host_ 限定符无法一起使用。
_global_ 函数的返回类型必须为空。
对 _global_ 函数的任何调用都必须按规定指定其执行配置。
_global_ 函数的调用是异步的，也就是说它会在设备执行完成之前返回。
_global_ 函数参数将同时通过共享存储器传递给设备，且限制为 256 字节。
对于变量前缀：
1.__device__
_device_ 限定符声明位于设备上的变量。
在接下来的三节中介绍的其他类型限定符中，最多只能有一种可与 _device_ 限定符一起使用，以更具体地指定变量属于哪个存储器空间。如果未出现其他任何限定符，则变量具有以下特征：
n 位于全局存储器空间中；
n 与应用程序具有相同的生命周期；
可通过网格内的所有线程访问，也可通过运行时库从主机访问。
2.__constant__
_constant_ 限定符可选择与 _device_ 限定符一起使用，所声明的变量具有以下特征：
n 位于固定存储器空间中；
n 与应用程序具有相同的生命周期；
可通过网格内的所有线程访问，也可通过运行时库从主机访问。
3.__shared__
_shared_ 限定符可选择与 _device_ 限定符一起使用，所声明的变量具有以下特征：
n 位于线程块的共享存储器空间中；
n 与块具有相同的生命周期；
n 尽可通过块内的所有线程访问。
只有在 _syncthreads()_（参见第 4.4.2 节）的执行写入之后，才能保证共享变量对其他线程可见。除非变量被声明为瞬时变量，否则只要之前的语句完成，编译器即可随意优化共享存储器的读写操作。
限制：
不允许为在主机上执行的函数内的 struct 和 union 成员、形参和局部变量使用这些限定符。
_shared_ 和 _constant_ 变量具有隐含的静态存储。
_device_、_shared_ 和 _constant_ 变量无法使用 extern 关键字定义为外部变量。
_device_ 和 _constant_ 变量仅允许在文件作用域内使用。
不可为设备或从设备指派 _constant_ 变量，仅可通过主机运行时函数从主机指派（参见第 4.5.2.3 节和第 4.5.3.6 节）。
_shared_ 变量的声明中不可包含初始化。
下面是具体的一个应用：
将共享存储器中的变量声明为外部数组时，例如：
extern __shared__ float shared[];
数组的大小将在启动时确定（参见第 4.2.3 节）。所有变量均以这种形式声明，在存储器中的同一地址开始，因此数组中的变量布局必须通过偏移显式管理。例如，如果一名用户希望在动态分配的共享存储器内获得与以下代码对应的内容：
short array0[128];
float array1[64];
int array2[256];
则应通过以下方法声明和初始化数组：
extern __shared__ char array[];
__device__ void func() // __device__ or __global__ function
{
short* array0 = (short*)array;
float* array1 = (float*)&array0[128];
int* array2 = (int*)&array1[64];
}
在设备代码中声明、不带任何限定符的自动变量通常位于寄存器中。但在某些情况下，编译器可能选择将其置于本地存储器中。
只要编译器能够确定在设备上执行的代码中的指针指向的是共享存储器空间还是全局存储器空间，此类指针即受支持，否则将仅限于指向在全局存储器空间中分配或声明的存储器。
通过获取 _device_、_shared_ 或 _constant_ 变量的地址而获得的地址仅可在设备代码中使用。通过 cudaGetSymbolAddress()获取的 _device_ 或 _constant_ 变量的地址仅可在主机代码中使用。
对global函数进行配置
对 _global_ 函数的任何调用都必须指定该调用的执行配置。
执行配置定义将用于在该设备上执行函数的网格和块的维度，以及相关的流。可通过在函数名称和括号参数列表之间插入 <<<Dg, Db, Ns, s>>>形式的表达式来指定，其中：
Dg 的类型为 dim3，指定网格的维度和大小，Dg.x * Dg.y 等于所启动的块数量，Dg.z 无用；
Db 的类型为 dim3，指定各块的维度和大小，Db.x * Db.y * Db.z 等于各块的线程数量；
Ns 的类型为 size_t，指定各块为此调用动态分配的共享存储器（除静态分配的存储器之外），这些动态分配的存储器可供声明为外部数组的其他任何变量使用，Ns 是一个可选参数，默认值为 0；
S 的类型为 cudaStream_t，指定相关流；S 是一个可选参数，默认值为 0。
举例来说，一个函数的声明如下：
__global__ void Func(float* parameter);
必须通过如下方法来调用此函数：
Func<<< Dg, Db, Ns >>>(parameter);
执行配置的参数将在实际函数参数之前被评估，与函数参数相同，通过共享存储器同时传递给设备。
如果 Dg 或 Db 大于设备允许的最大大小，或 Ns 大于设备上可用的共享存储器最大值，或者小于静态分配、函数参数和执行配置所需的共享存储器数量，则函数将失败。
SIMT架构
NVIDIA GPU是典型的SIMT架构(Single-Instruction, Multiple-Thread Architecture，单指令多线程架构)。(半)线程束在同一时间内执行同样的指令(相同的PC)，但每个线程有自己的数据空间(寄存器)；可以同时做同样的事情，但是处理不同的数据。这样可以有效节省指令带宽(指令由线程束共享)，实现高效并行。
当线程束中的线程，必须要执行不同的条件分支时，满足分支条件的线程会被激活并执行分支内的内容；不满足分支条件的线程会接收同样的指令，但不会被激活，不会实际执行，但也不能跳过去执行其他指令。换言之，当线程束中的线程遇到分支时，不论线程是否需要执行分支，都会消耗执行该分支的时间，因为线程束中的线程执行同样的指令，是高度同步的。
举个例子，假如程序是按如下方式编写的：
...
if(条件)
{
操作1
}
else
{
操作2
}
...
如果是CPU，如果不满足条件，则会直接跳转去执行操作2，不会执行操作1。 但是对于GPU同一个线程束内的线程，即使不满足条件，依旧会去跟其他线程一起去执行操作1。等到操作1执行完毕后，该线程又会跟其他线程一起执行操作2。这也就意味着，不管线程实际上需不需要执行某一分支，它都要跟其他线程一起跑完这一分支。
另外，GPU不支持分支预测和推测执行，只会老老实实的一条条执行各个分支里的内容。因此，对于CUDA程序来说，分支实际上是一种低效的行为。
但是有两种情况是例外：
线程束中所有线程均只需要执行一个分支，如线程束中的所有线程满足条件，则所有线程执行完操作1后，不会去执行操作2。 线程束中的半个线程束(线程0~15、线程16~31)同时满足条件，则这半个线程束不会去执行操作2(因为硬件实际上是按照半个线程束调度的)。
但是从Volta架构开始，Independent Thread Scheduling被引入，线程束内的线程不再完全同步。每个线程都会有自己独立的PC。遇到分支时，不再像之前的架构一样，只有(半)线程束内的线程条件一致时，才会跳过分支；Volta架构的调度优化器会将线程束中的线程，按照分支条件是否满足，重新组合成SIMT单元，从而跳过分支。
Volta架构的Independent Thread Scheduling无疑是高效的，但是这是一个跟旧架构完全不同的特性。在编写旧架构的CUDA程序时，程序员会默认线程束内的线程一定会同步执行。Volta架构的新特性破坏了这一假设，无疑会给程序带来一些问题，需要注意。
