CUDA(2)

cudaMalloc
函数告诉CUDA运行时在设备上分配内存。
函数使用限制

 1. 可以将cudaMalloc()分配的指针传递给在设备上执行的函数。
 2. 可以在设备代码中使用cudaMalloc()分配的指针进行内存读/写操作。
 3. 可以将cudaMalloc()分配的指针传递给在主机上执行的函数。
 4. 不能在主机代码中使用cudaMalloc()分配的指针进行内存读/写操作。
 5. 不能使用标准C的free()函数来释放cudaMalloc()分配的内存；要释放cudaMalloc()分配的内存，需要调用cudaFree()。

设备指针的使用方式与标准C中指针的使用方式完全一样。主机指针只能访问主机代码中的内存，而设备指针也只能访问设备代码中的内存。 在主机代码中可以通过调用cudaMemcpy()来访问设备上的内存。

内置变量

CUDA运行时允许启动一个二维线程格，并且线程格中的每个线程块都是一个三维的线程数组。

 1. blockDim  线程块中的线程数量 。
 2. gridDim   线程格中线程块的数量。
 3. blockIdx  当前执行设备代码的线程块的索引。
 4. threadIdx 当前执行设备代码的线程索引。

对于使用多个线程块并且每个线程块中包含多个线程。计算线程索引为

int tid = threadIdx.x+blockIdx.x*blockDim.x

多CPU时，tid每次增加的是CPU的数量。在GPU实现中，将并行线程的数量作为处理器的数量。

tid += blockDim.x*gridDim.x;

CUDA将每个线程并行执行，需要计算每个并行线程的初始索引和递增的量值。递增的步长为线程格中正在运行的线程数量。

共享内存、常量

将__shared__添加到变量声明中，这将使这个变量驻留在共享内存中。CUDA C编译器对共享内存中的变量与普通变量将分别采取不同的处理方式。共享内存存在GPU上，因此访问共享内存的速度高于访问普通的系统内存。
__shared__ 共享变量是block中所有线程共享的变量。共享变量的声明可以在kernel内部，也可以作为全局变量。如果共享变量是在kernel中声明的，其作用域就是kernel内，否则对所有kernel有效。
如果共享变量的大小在声明时不指定 :

extern __shared__ int  title[];

常量内存用于保存在核函数执行期间不会发生变化的数据。CUDA对常量内存采取了不同于标准全局内存的处理方式。在某些情况下，用常量内存来替换全局内存能有效地减少内存带宽。要使用常量内存，需在变量前面加上__constant__关键字。

线程 、线程组

GPU上的线程调度方式与CPU有很大不同。CPU的调度算法包含优先级、时间片等。GPU上所有线程机会均等，线程状态只有等待资源和执行两种状态，如果资源未就绪，那么就等待；一旦就绪，立即执行。

当GPU资源很充裕时，所有线程都是并发执行的，这样加速效果很接近理论加速比；而GPU资源少于总线程个数时，有一部分线程就会等待前面执行的线程释放资源，从而变为串行化执行。

块并行相当于CPU多进程的情况。CUDA的线程块将一组线程组织到一起，共同分配一部分资源，然后内部调度执行。线程块与线程块之间没有关系。

线程并行和线程块并行的区别:

线程并行是细粒度并行，调度效率高；块并行是粗粒度并行，每次调度都要重新分配资源，有时资源只有一份，那么所有线程块都只能排成一队，串行执行。

设备信息

//获得CUDA设备个数
int count;
HANDLE_ERROR(cudaGetDeviceCount(&count));
//迭代查询每个设备的相关信息。cudaDeviceProp结构包含了设备的相关信息
cudaDeviceProp prop;
for(int i=0;i<count;i++)
{
    HANDLE_ERROR(cudaGetDeviceProperties(&prop,i));
}

// 选择满足条件的CUDA设备
cudaDeviceProp prop;
int dev;
memset(&prop,0,sizeof(cudaDeviceProp));
prop.major=1;
prop.minor=3;
HANDLE_ERROR(cudaChooseDevice(&dev,&prop));

kernel 函数

__device__ 修饰符表示代码将在GPU而不是host上运行。声明为__device__ 的函数只能被其他__device__ 或__global__ 修饰的函数调用。

线程同步

//对线程块中的线程进行同步
__syncthreads();

该函数确保线程块中每个线程都执行完__syncthreads()之前的语句后，才会执行下一条语句。
注意，将__syncthreads()放在if的结构中可能导致程序失败。这种情况下，可能会导致线程发散(某些线程需要执行一条指令，而其他线程不需要执行)。CUDA架构确保线程块中的每个线程都执行了

__syncthreads();

否则没有任何线程可以执行之后的语句。在线程发散的情况中，一些线程永远都无法执行__syncthreads() ,硬件将这些线程保持等待，一直等下去。