主要学习核函数、线程布局、内存模型、流的使用
主要实现归约求和、放射变化、矩阵乘法、模型后处理
内存模型:pinned memory(cpu内存) global memory shared memory(GPU内存)
CPU有两个大类:pageable和page lock(vip房间)
房间不够用的时候pageable拿出来给别人用。容纳更多的数据。内存没有,但是跑更多的程序,性能降低。
pinned memory具有锁定性
pageable memory没有锁定性,可能数据不正确,但是提高程序运行数量
因为锁定性,所以GPU访问的是pinned memory
gpu直接访问pinned memory,叫DMA技术。
对GPU越近计算效率高
pinned memory<global memory <shared memory
流的使用
流是基于上下文之上的任务管理通道,一个contest可以创建多个流
流是异步控制
nullptr表示默认流
流就是:一个是管道,里面要执行的内容,任务队列(买西瓜),一个是要执行器(男朋友)
执行的代码加入流后,立即返回,不耽误时间。
带异步函数使用流,cudamemcpyAsync.
核函数:
__global__为核函数由host调用,定义在GPU上,可以在CPU上调用的函数;
__device__表示设备函数,由device调用,定义在GPU上,可以在CPU上调用的函数;
__host__为主机函数,由host调用,定义在GPU上,可以在CPU上调用的函数;
__share__表示变量为共享变量
host调用核函数,function<<<gridDim,blockDim,sharedMemorySize,stream>>>(args...)
gridDim,blockDim为多少个线程
cudaMalloc(**devPtr, byte_size),在GPU上申请空间:
int *gpu_int;
cudaMalloc((void**)&gpu_int, sizeof(int))
申请sizeof(int)个尺寸的gpu
cudaMemset(*devptr, value, byte_size),初始化
例如cudaMemset(gpu_int, 10, 10 * sizeof(int));
输出为10,10,10,10,10,10,10,10,10,10
cudaMemcpy(*dst, *src, byte_size, 类型),GPU和CPU参数传递函数,有4中:
CPU2CPU:cudaMemcpyHostToHost
CPU2GPU:cudaMemcpyHostToDevice
GPU2CPU:cudaMemcpyDeviceToHost
GPU2GPU:cudaMemcpyDeviceToDevice
cudaMemcpy(cpu_int, gpu_int, 10 * sizeof(int), cudaMemcpyDeviceToHost);
核函数参数
jobs总共需要多少个线程
每个线程快中放多少个线程
有多少个块
int jobs = dst_height * dst_width;
int threads = 256;
int blocks = ceil(jobs / (float)threads);
warpaffine_kernel <<<blocks, threads>>> (img_buffer_device, src_width * 3, src_width,src_height, dst, dst_width,dst_height, 128, d2s, jobs);
int position = blockDim.x * blockIdx.x + threadIdx.x;
if (position >= edge) return;