CUDA笔记

1.cudaDeviceSynchronize()用于CPU和GPU同步，即cpu和GPU均运行至cudaDeviceSynchronize()后再继续；CPU多线程时，会阻止所有线程;

2._syncthreads()用于核函数内线程块线程同步，即同一block内所有线程执行至__syncthreads()处等待全部线程执行完毕后再继续（因此_syncthreads()不要位于发散分支中（if）,因为如果一些线程无法执行_syncthreads()，硬件将会使这些线程一直保持等待）；

3.cudaMemcpy与cpu是同步的，即在GPU中运行完后才会继续执行cpu;

4.核函数是异步的，cpu不会等待核函数内运行完再继续；

5.在device或global中，可以使用malloc和new申请显存,也可以使用memcpy拷贝数据；在设备中可以使用printf打印值；

6.要为device或global申请全局变量，变量前声明__device__,且只能通过cudaMemcpyToSymbol申请内存和赋值，例：

__device__  float* dev_tmp;//为显存申明全局变量;
//为dev_tmp申请了内存并赋值
float *tmp=new float[20];
for(int i=0;i<20;++i)
     tmp[i]=i;
float *ptr;
cudaMalloc((void**)&ptr,20*sizeof(float));
cudaMemcpy(ptr,tmp,20*sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(dev_tmp,&ptr,sizeof(ptr));

__global__ void func()
{
    for(int i=0;i<20;++i)
       printf("%f ",dev_tmp[i]);
}

7.动态共享内存

extern __shared__ float cache[];动态共享内存的大小通过核函数第三个参数输入；

https://www.cnblogs.com/cofludy/p/7622254.html

8.把线程块分解为线程的目的：

(1).线程块数量的硬件限制；

(2).共享内存的使用（共享内存始终驻留在物理GPU上，不是驻扎在GPU之外的系统内存上）

（将线程块和线程分为二维（dim3）似乎只是为了某种索引方便？）

9.常量内存：

常量内存用于保存在核函数执行期间不会发生变化的数据，使用常量内存在一些情况下，能有效减少内存带宽，降低GPU运算单元的空闲等待；常量内存是只读内存，声明时要静态地分配空间;

__constant__ float s[1600];

将数据从CPU拷贝到常量内存中时用cudaMemcpyToSymbol；

__constant__ float dev_s[1600];
float s[1600];
cudaMemcpyToSymbol(dev_s,s,1600*sizeof(float));