cuda基础

一：核函数调用与参数传递
1：设备指针
1）可以将cudaMalloc()分配的指针传递给在设备上执行的函数
2）可以用cudaMalloc()分配的指针在设备上进行内存读写操作
3）可以将设备指针传递给在主机上执行的函数
4）不能在主机代码中使用设备指针对内存进行读写操作

二：设备属性
1：使用
1）当编写支持双精度浮点数的应用程序时需要查询支持该功能的设备

三：共享内存与同步
1：对于共享变量，编译器为每个线程块生成一个副本，只需更具线程块中线程数量分配数量
2：归约算法
每个线程将cache[]中的两个数据想加，再放回cache[]。得到的结果为原始数据的一半，执行log2(threadsPerBlock)个步骤，得到所有cache[]的总和
3：线程块选择
min（32,(N+threadPerBlock)/thredPerBlock）
4:线程发散
某些线程执行一条指令，其他线程不执行。
当__synthreads()位于发散分支中，会使设备一直等待而出错。

四：常量内存
1:常量内存使用：
当所有线程束访问相同的只读数据时，常量内存可以大大提高性能。
常量内存可以广播半个线程束。
2：事件与性能分析
cudaEvent_t start,stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);
//do somthing
cudaEventRecord(stop,0);
cudaEventSynchrnize(stop);//使CUP在某个事件上同步，因为GPU执行异步函数调用时，CPU会执行下一条语句。//运行时阻塞
int time;
cudaEventElapsedTime(&time,start,stop);//统计时间
cudaEventDestroy(start);
cudaEventDestory(stop);

五：纹理内存
1：使用--分配内存，绑定内存，核函数调用，解除绑定，释放内存。
1）将输入变量声明为texture类型引用
texture<DataType,type,readmode> texout;//type--参考系类型（1，2，3维）
2)将变量绑定到内存缓冲区
cudaBindTexture(NULL,texout,data.dev_outSrc,imageSize)

六：流
1：页锁定主机内存--固定内存
操作系统不会将固定内存映射到磁盘且不可分页，始终驻留在物理内存中（可以通过直接内存访问DMA进行主机与GPU之间的复制）
1）仅针对cudaMemcpy()调用中的目标内存或源内存，才使用页锁定内存。且在不使用时以及释放（cudaFreeHost）
2）分配流使用的固定内存
cudaHostAlloc(host_a,N*sizeof(int),cudaHostAllocDefault);
3）只能以异步方式对固定内存进行复制操作
cudaMemcpyAsync() ---- cudaHostAlloc()

异步方式不能保证好久执行操作，但可以保证在下一个操作之前执行完。

2:cuda流
1）设备支持重叠功能--prop.deviceOverlap
例如：支持核函数的同时还支持主机与设备之间的复制
int main()
{
cudaDeviceProp prop;
int whichDevice;
HANDLE_ERROR(cudaGetDevice(&whichDevice));
HANDLE_ERROR(cudaGetDeviceProperties(&prop,whichDevice));
if (!prop.deviceOverlap)
printf("Device not support overlap; ");
return 0;
}

2)高效的使用多个cuda流
将操作放入队列时采用宽度优先方式：将流之间的操作交叉添加在队列中。

//由于GPU内存远小于主机内存，因此使用分块的方式执行计算。
for(int i = 0;i < FULL_DATA_SIZE;i += 2*N)
{
（1）深度优先方式
HANDLE_ERROR(cudaMemcpyAsync(dev_a0,host_a+i,N*sizeof(int),cudaMemcpyHostToDevice,stream0));
HANDLE_ERROR(cudaMemcpyAsync(dev_b0,host_b+i,N*sizeof(int),cudaMemcpyHostToDevice,stream0));

kernel<<<N/256,256,dev_a0,dev_b0,dev_c0>>>;

HANDLE_ERROR(cudaMemcpyAsync(host_c+i,dev_c0,N*sizeof(int),cudaMemcpyDeviceToHost,stream0));

（2）宽度优先方式
HANDLE_ERROR(cudaMemcpyAsync(dev_a0,host_a+i+N,N*sizeof(int),cudaMemcpyHostToDevice,stream1));
HANDLE_ERROR(cudaMemcpyAsync(dev_b0,host_b+i+N,N*sizeof(int),cudaMemcpyHostToDevice,stream1));

kernel<<<N/256,256,dev_a1,dev_b1,dev_c1>>>;

HANDLE_ERROR(cudaMemcpyAsync(host_c+i+N,dev_c0,N*sizeof(int),cudaMemcpyDeviceToHost,stream1));
}

计算a中三个值与b中三个值的平均值

#define N (1024*1024)
#define FULL_DATA_SIZE (n*1024)

__global__ void kernel(int *a,int *b,int *c)
{
int idx = threadIdx.x + blockDim.x*blockIdx.x;
if(idx <　Ｎ)
｛
int idx1 = (idx + 1) % 256;
int idx2 = (idx + 2) % 256;
float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
c[idx] = (as + bs) / 2;
｝

}
int main()
{
cudaDeviceProp prop;
int whichDevice;
HANDLE_ERROR(cudaGetDevice(&whichDevice));
HANDLE_ERROR(cudaGetDeviceProperties(&prop,whichDevice));
if (!prop.deviceOverlap)
printf("Device not support overlap; ");

//创建启动计时器
cudaEvent_t start,stop;
float elapsedtime;

HANDLE_ERROR(cudaEventCreate(&start));
HANDLE_ERROR(cudaEventCreate(&stop));
HANDLE_ERROR(cudaEventRecord(start,0));

//创建流
cudaStream_t stream0,stream1;
HANDLE_ERROR(cudaStreamCreate(&stream0));
HANDLE_ERROR(cudaStreamCreate(&stream1));
return 0;
}

七：零拷贝主机内存

首先应该查询设备：
cudaDeviceProp prop;
int whichDevice;
//查看设备支持零拷贝内存吗
HANDLE_ERROR(cudaGetDevice(&whichDevice));
HANDLE_ERROR(cudaGetDeviceProperties(&prop,whichDevice));

if(prop.cudaMapHostMemory != 1)
{
printf("cannot map memory ");
return 0;
}

//标识希望设备映射主机内存
HANDLE_ERROR(cudaSetDeviceFlags( cudaDeviceMapHost));




1：获取零拷贝主机内存
cudaHostAlloc() ---- cudaHostAllocMapped
与固定内存有相同的属性，但可以在核函数中直接访问这种内存
cudaHostAlloc((void **)&a,size,cudaHostWriteCombined | cudaHostMapped) //合并式写入
2：将主机指针映射到设备指针
cudaHostGetDevicePointer(&dev_a,a,0);

3：性能
1）集成的GPU
与主机共享系统内存，使用零拷贝内存可提高性能，但是该内存与固定内存性能相同
2）单独的GPU
一次性读写操作可使用零拷贝内存，但主机不会缓存零拷贝内存的内容，多次读写时应该避开。

八：使用多GPU
在需要使用多GPU之前应查询设备数
应为每个GPU分配一个不同的线程来控制。
1：线程函数routine()
routine(&(data[1]));
可在应用程序的默认线程中调用。
2：线程辅助函数和start_routine()
CUTThread thread = start_routine(routine,&(data[0]));
3：线程阻塞函数end_thread(thread)
主应用线程将等待其他线程运行完。


九：可移动的固定内存
对于分配固定内存的线程来说是锁定内存，对于其他的来说是可分页的固定内存。要使所有线程都将这块内存视为固定内存，就需要使用可移动的固定内存机制
1）分配
cudaHostAlloc() --- cudaHostPortable