Cuda学习笔记
文章平均质量分 74
以面向高性能计算为主,目标是写好多流异步双缓冲的GEMM的cuda学习笔记
库达ZT
这个作者很懒,什么都没留下…
展开
-
SPMV基础-1
稀疏矩阵*向量原创 2023-07-11 18:35:05 · 820 阅读 · 0 评论 -
TRT3-trt-basic - 3 动态shape
这一部分都是和之前一样是不变的。原创 2023-07-03 23:38:15 · 130 阅读 · 0 评论 -
TRT2-CUDA_Runtime_Api-2 流
流是一种基于context上的任务管道抽象,一个context可以创建多个流。流是异步控制的主要方式。nullptr表示默认流,每个线程都有自己的默认流。// CUDA运行时头文件#include <cuda_runtime.h>#include <stdio.h>#include <string.h>#define checkRuntime(op) __check_cuda_runtime((op), #op, __FILE__, __LI原创 2023-06-30 13:10:53 · 60 阅读 · 0 评论 -
TRT3-trt-basic - 2 inferance
这些都是上一节的代码,就不多说了。原创 2023-07-03 19:11:08 · 67 阅读 · 0 评论 -
TRT3-trt-basic - 1 basic
/第一件事,log类,获得msg。然后可以根据级别决定要不要输出public:第一件事,log类,获得msg。然后可以根据级别决定要不要输出// ----------------------------- 1. 定义 builder, config 和network -----------------------------// 这是基本需要的组件//形象的理解是你需要一个builder去build这个网络,网络自身有结构,这个结构可以有不同的配置。原创 2023-07-03 00:18:24 · 100 阅读 · 0 评论 -
SGEMM_test(持续性test)
【代码】SGEMM_test(持续性test)原创 2023-07-02 21:41:56 · 83 阅读 · 0 评论 -
TRT2-CUDA_Runtime_Api-4 错误处理
1. 若cuda核函数出错,由于他是异步的,立即执行cudaPeekAtLastError只会拿到对输入参数校验是否正确的状态,而不会拿到核函数是否执行正确的状态2. 因此需要等待核函数执行完毕后,才真的知道当前核函数是否出错,一般通过设备同步或者流同步进行等待block越界访问越界。原创 2023-07-02 21:10:03 · 50 阅读 · 0 评论 -
YOLOv5 后处理cuda实现
在mian函数中,运用load_file读取tensor文件这里用load_file打开图片, 这里是用二进制模式打开文件(ios::binary), 使用static std::vector存储数据。YOLOV5给出来的data是n x (5 + classes)的, 这里通过计算可以获得行数列数, 然后传入指向data的指针, nrows, ncols解码, 本案例提供cpu解码和GPU解码解码结束后返回的是。原创 2023-07-02 15:01:15 · 819 阅读 · 0 评论 -
TRT2-CUDA_Runtime_Api-3 加速Warpaffine
warpaffine最核心的kernel就是这一段了。原创 2023-06-30 22:09:22 · 85 阅读 · 0 评论 -
TRT2-CUDA_Runtime_Api-2 核函数、shared memory
在.vscode/settings.json中配置*.cu : cuda-cpp,可以使得代码被正确解析Makefile中,cu交给nvcc进行编译cu文件可以当做正常cpp写即可,他是cpp的超集,兼容支持cpp的所有特性1. 核函数是 cuda 编程的关键2. 通过 xxx.cu 创建一个 cudac 程序文件,并把 cu 交给 nvcc 编译,才能识别 cuda 语法3. __global__ 表示为核函数,由 host 调用。__device__ 表示为设备函数,由 device 调用。原创 2023-06-30 17:53:27 · 55 阅读 · 0 评论 -
TRT2-CUDA_Runtime_Api-1 驱动
GPU可以直接访问pinned memory,称之为(DMA Direct Memory Access)对于GPU访问而言,距离计算单元越近,效率越高,所以PinnedMemory原创 2023-06-30 00:39:41 · 81 阅读 · 0 评论 -
TRT1-CUDA_driver_Api-1 驱动、Check、Context、Memory
而由于高频的操作,所以对于设备而言,基本是一个线程固定访问一个显卡的,就是只是用一个context,很少会用到多个context。在这种情况下再一个个push、pop就很麻烦,所以推出了cuDevicePrimaryCtxRetain(),他是一个给他设备id,自动配置context,就是一个显卡对应一个primary context.不同的线程,只要设备id一样,那就是一样的primary context。函数用于增加设备的主上下文的引用计数,以便继续使用已经创建的上下文。的地址,和预期结果相符。原创 2023-06-28 23:59:11 · 180 阅读 · 1 评论 -
CUDA_GPU编程
•线程组分歧():尽量保证32个线程都进同样的分支,否则两个分支都会执行。•延迟隐藏():需要有足够的blockDim供SM在陷入内存等待时调度到其他线程组。•寄存器打翻():如果核函数用到很多局部变量(寄存器),则blockDim不宜太大。•共享内存():全局内存比较低效,如果需要多次使用,可以先读到共享内存。•跨步访问(coalescedacccess):建议先顺序读到共享内存,让高带宽的共享内存来承受跨步。•区块冲突():同一个warp。原创 2023-07-01 18:37:51 · 764 阅读 · 0 评论 -
CUDA_调整指令级原语3原子操作
在本节中,将学习如何使用原子操作,并学习在高并发环境下的共享数据上如何执 行正确的操作。通过使用一个原子函数,每个由CUDA提供的原子函数可以重复被执行:原子级比较并交换CAS运算符。原子级CAS是一个很重要的操作,不仅可以使你在CUDA中定义你自己的原子函数,还能帮助你更深层次地理解原子操作。CAS将3个内容作为输入:内存地址,存储在此地址中的期望值,已经实际想要存储在此位置的新值;然后执行以下几步;1. 读取目标地址并将该处地址的存储值与预期值进行比较。原创 2023-05-23 22:36:46 · 134 阅读 · 1 评论 -
CUDA_调整指令级原语2标准函数与内部函数
标准函数和内部函数在数值精确度和性能上的表现是不同的。标准函数支持大部分的数学运算。但是,许多等效的内部函数能够使用较少的指令、改进的性能和更低的数值精确度,实现相同的功能。结果如下图所示:可以看到对于标准函数powf或是内部函数__powf,与主机中所计算的pow都是不同的,但是标准函数的误差明显更小,但是内部函数的计算更快。即使使用数值稳定的CUDA函数,GPU上的运算结果仍与传统的只在CPU上运行的应用结果不同。原创 2023-05-23 21:10:51 · 145 阅读 · 1 评论 -
CUDA_调整指令级原语1浮点运算
讲了一下双精度单精度在实际应用情况中传入传出以及复杂运算的消耗。原创 2023-05-23 15:44:53 · 262 阅读 · 1 评论 -
CUDA 流——重叠GPUCPU基础
这张图左边的竖线就代表的是我们上一张图所标注的位置,其实可以看到,在提交完DtoH的命令之后,HtoD的数据传输可能才刚刚开始。再者就是对于这个代码还有很大的改进空间了,比如可以在写入的同时读取数据,这些就是双缓冲的内容了,等过两天再看看搞一搞。程序运行环境:对于同样的 CUDA 代码,在不同的计算机、不同的 GPU 或者不同的设备驱动程序版本下,可能会有所不同的执行时间,这取决于硬件配置、设备驱动程序质量等因素。但很明显,在HtoD的过程中,cudaeventquert就已经出现了不同的时长。原创 2023-05-18 17:27:16 · 188 阅读 · 1 评论