云服务器可以做两个网站,做网站维护挣钱吗,福建住房和建设网站密码忘记,无锡微盟网络科技有限公司精简CUDA教程——CUDA Runtime API tensorRT从零起步迈向高性能工业级部署#xff08;就业导向#xff09; 课程笔记#xff0c;讲师讲的不错#xff0c;可以去看原视频支持下。 Runtime API 概述
环境 图中可以看到#xff0c;Runtime API 是基于 Driver API 之上开发的…精简CUDA教程——CUDA Runtime API tensorRT从零起步迈向高性能工业级部署就业导向 课程笔记讲师讲的不错可以去看原视频支持下。 Runtime API 概述
环境 图中可以看到Runtime API 是基于 Driver API 之上开发的一套 API。之前提到过 Driver API 基本都是 cu 开头的而Runtime API 基本都是以 cuda 开头的。
Runtime API 的特点
Runtime API 与 Driver API 最大的区别是懒加载 即在真正执行功能时才自动完成对应的动作即 第一个 Runtime API 调用时会自动进行 cuInit 初始化避免 Driver API 未初始化的错误第一个需要 context 的 API 调用时会创建 context 并进行 context 关联和设置当前 context调用 cuDevicePrimaryCtxRetain 实现绝大部分 api 都需要 context例如查询当前显卡名称、参数、内存分配释放等 CUDA Runtime 是封装了 CUDA Driver 的更高级别、更友好的 APIRuntime API 使用 cuDevicePrimaryCtxRetain 为每个设备设置 context不再手动管理 context并且不提供直接管理 context 的 API可 Driver API 管理通常不需要可以更友好地执行核函数.cpp 可以与 .cu 文件无缝对接Runtime API 对应 cuda_runtime.h 和 libcudart.soRuntime API 随 cudatoolkit 发布主要知识点是核函数的使用、线程束布局、内存模型、流的使用主要是为了实现归约求和、放射变换、矩阵乘法、模型后处理就可以解决绝大部分问题
错误处理
类似于在介绍 Driver API 时的情况我们同样提出 Runtime API 的错误处理方式
#define checkRuntime(op) __check_cuda_runtime((op), #op, __FILE__, __LINE__)bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){if(code ! cudaSuccess){const char* err_name cudaGetErrorName(code);const char* err_message cudaGetErrorString(code);printf(runtime error %s:%d %s failed. \n code %s, message %s\n, file, line, op, err_name, err_message);return false;}return true;
}内存模型 pinned memory
内存模型是 CUDA 中很重要的知识点主要理解 pinned_memory、global_memory、shared_memory 即可其他的不太常用。pinned_memory 属于 host memory而 global_memory、shared_memory 属于 device memory。
下图是的 Device Memory 的分类 锁定性和性能
对于主机内存即整个 host memory 而言操作系统在逻辑上将其区分为两个大类
pageable memory可分页内存page lock memory (pinned memory)页锁定内存/锁页内存
可以理解为 page lock memory 是酒店的 vip 房间锁定给你一个人使用。而 pageable memory 是普通房间在酒店房间不够的时候选择性地将你的房间腾出来交换到硬盘上给其他人使用这样就能容纳更多人了。造成房间很多的假象代价是性能很低。pageable memory 就是常见的虚拟内存的特性。
基于前面的理解我们总结如下
锁定性 pinned memory 具有锁定特性是稳定不会被交换的这很重要相当于每次去这个房间都一定能找到你pageable memory 没有锁定特性对于第三方设备如 GPU去访问时因为无法感知内存是否被交换可能得到不到正确的数据相当于每次去房间找你说不定你的房间正好被交换了因此 GPU 可以直接访问 pinned memory 而不能访问 pageable memory 性能 pageable memory 的性能比 pinned memory 差因为我们的 pageable memory 很可能会被交换到硬盘上pageable memory 策略能使用内存假象比如实际只有 8G 内存却能使用 16G借助 swap 交换从而提高程序的运行数量pinned memory 也不能太多会导致操作系统整体性能变差可同时运行的程序变少而且 8G 内存最多就 8G 锁页内存。
数据传输到GPU pinned memory 可以直接传送数据到 GPU 而 pageable memory 由于并不锁定需要先传到 pinned memory。
关于内存其他几个点 GPU 可以直接访问 pinned memory称为 DMA (Direct Memort Access) 对于 GPU 访问而言距离计算单元越近效率越高所以 SharedMemory GlobalMemory PinnedMemory 代码中 由 new/malloc 分配的是 pageable memory由 cudaMallocHost 分配的是 PinnedMemory由 cudaMalloc 分配的是 GlobalMemory 尽量多用 PinnedMemory 储存 host 数据或者显式处理 Host 到 Device 时用 PinnedMemory 做缓存都是提高性能的关键
流 stream
流是一种基于 context 之上的任务管道任务队列抽象一个 context 可以创建 n 个流流是异步控制的主要方式nullptr 表示默认流每个线程都有自己的默认流。
生活中的例子
同步串行异步
在这个例子中男朋友的微信消息就是任务队列流的一种抽象女朋友发出指令之后她可以做任何事情无需等待指令执行完毕。即异步操作中执行的代码加入流的队列之后立即返回不耽误时间。女朋友发的指令被送到流中排队男朋友根据流的队列顺序执行。女朋友选择性在需要的时候等待所有的执行结果新建一个流就是新建一个男朋友给他发指令就是发微信可以新建很多个男朋友通过 cudaEvent 可以选择性等待任务队列中的部分任务是否就绪
注意
要十分注意指令发出后流队列中储存的是指令参数不能在任务加入队列后立即释放参数指针这会导致流队列执行该指令时指针失效而出错。应当在十分肯定流已经不需要这个指针之后才进行修改或释放否则会有非预期行为出现。
就比如女朋友让男朋友去卖西瓜并转给了他钱但是却在男朋友买瓜成功前将转账撤了回去这时就无法知道男朋友在水果店会发生什么比如会不会跟老板打起来之类的。因此要保证买瓜行为顺利完成行为符合预期在买瓜成功前就不能动买瓜的钱。
核函数
简介 核函数是 cuda 编程的关键 通过 xxx.cu 创建一个 cudac 程序文件并把 cu 文件交给 nvcc 编译才能识别 cuda 语法 __xxx__ 修饰 __global__ 表示为核函数由 host 调用__device__ 表示设备函数由 device 调用__host__ 表示主机函数由 host 调用__shared__ 表示变量为共享变量。可能存在上述多个关键字修饰同一个函数如 __device__ 和 __host__ 修饰的函数既可以设备上调用也可以在主机上调用 host 调用核函数 functiongridDim, blockDim, sharedMemorySize, stream(args, ...)gridDim 和 blockDim 的变量类型为 dim3是一个三维的值 function 函数总共启动的线程数目可以这样计算n_threads gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z 详细请参考线程束的相关知识 只有 __global__ 修饰的函数才可以用 的方式调用s 调用核函数是传值的不能传引用可以传递类结构体等核函数可以使模板 核函数的返回值必须是 void 核函数的执行是异步的也就是立即返回的 线程 layout 主要用到 blockDim、gridDim 和函数内访问线程索引主要用到 threadIdx、blockIdx、blockDim、gridDim 这些内置变量
线程索引计算
共涉及四个变量blockDim、gridDim、threadIdx、blockIdx 其中前两者可以认为是形状后两者可以认为是对应的索引。就像我们 PyTorch 中如果一个张量的形状为 (2,3)(2,3)(2,3) 那么对应的其两个维度上索引的取值范围就是0−1,0−20-1,0-20−1,0−2。 线程索引 id 计算方法左乘右加如上图所示。
共享内存 由 __shared__ 关键字修饰 共享内存因为更靠近计算单元所以访问速度更快 共享内存通常可以作为访问全局内存的缓存使用 可以利用共享内存实现线程间的通信 通常与 __syncthreads 同时出现这个函数是同步 block 内的所有线程全部执行到这一行才往下继续执行 如 __shared__ int shared_value1;
__shared__ int shared_value2;if (threadIdx.x 0) {if (blockIdx.x 0) {shared_value1 123;shared_value2 55;}else {shared_value1 331;shared_value2 8;}__syncthreads();printf(...)
}其他 threadIdx.x 不为 0 的线程不会进到判断语句但是会卡在 __syncthreads() 等待 threadIdx.x 为 0 的线程设置好共享内存再一起继续向下执行。 共享内存使用方式通常是在线程 id 为 0 的时候从 global memory 取值然后 __syncthreads 然后再使用 动态共享内存与静态共享内存 动态共享内存的声明需要加 extern 关键字不需要指定数组大小如 extern __shared__ char dynamic_shared_memory[];静态共享内存的声明需要指定数组大小如 const size_t static_shared_memory_size 6 * 1024; // 6KB
__shared__ char static_shared_memory[static_shared_memory_size];warp affine 实战
chapter: 1.6, caption: vector-add, description: 使用cuda核函数实现向量加法
chapter: 1.7, caption: shared-memory, description: 共享内存的操作
chapter: 1.8, caption: reduce-sum, description: 规约求和的实现利用共享内存高性能
chapter: 1.9, caption: atomic, description: 原子操作实现动态数组的操作
chapter: 1.10, caption: warpaffine, description: 仿射变换双线性插值的实现yolov5的预处理
chapter: 1.11, caption: cublas-gemm, description: 通用矩阵乘法的cuda核函数实现以及cublasSgemm的调用
chapter: 1.12, caption: yolov5-postprocess, description: 使用cuda核函数实现yolov5的后处理案例TODO
thrust
相当于 cuda 的 stl但并不常用
错误处理
若核函数出错由于它是异步的立即执行 cudaPeekAtLastError 只会拿到对输入参数校验是否正确的状态而不会拿到核函数是否正确执行的状态。
需要等待核函数真正执行完毕之后才知道当前核函数是否出错一般通过设备同步或者流同步进行等待
错误分为可恢复和不可恢复两种
可恢复 参数配置错误例如 block 越界一般最大值是 1024shared memory 超出大小范围一般是 64KB等通过 cudaGetlastError 可以获取错误代码同时把当前状态恢复为success该种错误可以在调用核函数之后立即通过 cudaGetLastError / cudaPeekAtLastError 拿到该种错误在下一个函数调用时会覆盖 不可恢复 核函数执行错误例如访问越界等该错误会传递到之后所有的 cuda 操作上错误状态通常需要等到核函数执行完毕才能够拿到也就是有可能在后续的任何流程中突然异常因为是异步的
