CUDA基础

https://blog.csdn.net/dcrmg/article/details/54446393
https://www.cnblogs.com/qingsunny/p/3384779.html

文章内容基本取自于别人的博客，仅在此记录笔记以作为自己对CUDA的入门学习。

1. CPU被称为host，GPU被称为device。Host和Device可以通过global memory、constant memory、texture memory进行通信。
2. grid->block->thread，一个grid分为多个block，一个block分为多个thread。
3. sm->warp->sp，一个sm分为多个warp，一个warp含有多个sp。
4. 一个kernel对应一个grid，每一个warp包含一些thread。sp包含两个ALU和一个FPU，没有cache，现代GPU就是一组sp的array，被称为SPA，每一个sp执行一个thread。block被加载到sm上运行。通过block实现了细粒度的通信手段，由于block被加载到SM上，因此可以利用SM提供的shared memory和_syncthreads()功能实现线程同步和通信。block之间不能有依赖。
5. block和thread都有各自的ID，记作blockIdx(1D,2D),threadIdx(1D,2D,3D)
6. block和thread还有Dim，即blockDim，threadDim，他们都有三个分量x，y，z
7. 每个thread都有自己的一份register和local memory的空间。
8. 同一个block中的thread共享一份shared memory。
9. 所有的thread都共享一份global memory、constant memory、texture memory。
10. 不同的grid有各自的global memory、constant memory、texture memory。

关于memory的效率操作，low-latency和high-bandwidth是高性能的理想情况。
register是最快的，但是需要合理分配使用，让自己的kernel使用较少的register就能够允许更多的block驻留在SM中。可以用
__launch_bounds__(maxThreadsPerBlock,minBlocksPerMultiprocessor)可以指定每个block可以包含的最大thread数目和每个处理器最少要包含的block数目。

__global__表示编译在GPU上，因此函数将被交给编译设备代码的NVCC编译器来处理。
__device__分配GPU上的global memory空间，grid中所有线程可访问。
__constant__分配GPU上的constant memory空间，grid中所有线程可访问。
__shared__分配GPU上的thread block空间，block中所有线程可访问。
local是thread的单独空间，仅该thread可访问。

线程硬件原理：
每个sm最多分配8个block，每个sm最多可接受768个thread，可以是1个block包含512个thread，也可以是3个block每个包含256个thread。sm上每个block尺寸必须相同，每个线程的调度与ID由该sm管理。
sm满负载工作效率是最高的，block的尺寸可以为8*8，16*16，32*32。
如果是8*8就有64个线程，因此需要768/64=12个block，但是由于sm最多8个block，因此一个sm实际执行了8*64=512个线程。
如果是16*16每个block就有256个线程，sm可同时接受3个block，3*256=768，满负载。
如果是32*32每个block就有1024个线程，sm无法处理。
warp实际上是一个32路的SIMD,是sm的基本调度单位。

关于减少latency的操作：
1.GlobalMemory通过coalesced的方法减少latency：
对于不是4bytes、8bytes、16bytes的结构体可以用__align__(X)补齐：
如struct __align__(16) vec3d{float x,float y,float z};
还有其他情况下的解决方法。

2.SharedMemory则是通过减少bank conflict问题去减少latency：
每个sm有16k的shared memory，shared memory分为16个bank，如果每个thread都区不同的bank就不会产生问题，而如果X个thread同时访问bank则会使效率降为1/X。对于__shared__ int data[128],它在SharedMemory中的布局为data[0]是bank0,data[1]是bank1,data[2]是bank2...data[15]是bank15，而data[16]又回到bank0。对于int number=data[base + tid];不会产生bank conflict,
而对于int number=data[base+4*tid];会产生bank0和bank4存取到同一bank，因此需要设计避免bank conflict。

纹理存储器：在Kernel中访问纹理存储器的操作成为纹理拾取(texture fetching)，将显
存中数据与纹理参照系关联的操作，称为将数据与纹理绑定(texture binding)。纹理存储
器适用于图像处理或查找表，并且对于数据量较大时的随机数据访问或非对齐访问有很好
的加速效果。GPU可以通过普通的线性存储器和cuda数组绑定到纹理的数据。
线性存储器只能与一维或二维纹理绑定，采用整型纹理拾取坐标，坐标值与数据在存储器
中的位置相同。
CUDA数组可以与一维、二维、三维纹理绑定，纹理拾取坐标为归一化或非归一化的浮点型
，并且支持许多特殊功能。

关于threadIdx、blockIdx、blockDim、grimDim
threadIdx是一个uint3类型，表示一个线程的索引。
blockIdx是一个uint3类型，表示一个线程块的索引，一个线程块中通常有多个线程。
blockDim是一个dim3类型，表示线程块的大小。
gridDim是一个dim3类型，表示网格的大小，一个网格中通常有多个线程块。

//thread 1D
__global__ void testThread1(int *c,const int *a,const int *b)
{
    int i = threadIdx.x;
    c[i] = b[i] - a[i];
}
//thread 2D
__global__ void testThread2(int *c,const int *a,const int *b)
{
    int i = threadIdx.x + threadIdx.y*blockDim.x;
    c[i] = b[i] - a[i];
}
//thread 3D
__global__ void testThread3(int *c,const int *a,const int *b)
{
    int i = threadIdx.x + threadIdx.y*blockDim.x +
threadIdx.z*blockDim.x*blockDim.y;
    c[i] = b[i] - a[i];
}