CUDA编程
文章平均质量分 64
记录自己学习cuda的相关文章
s.feng
计算机视觉,C++
展开
-
clang 编译cuda原理
最近在看一门julia的语言,里面是原生支持cuda的,不过在国产卡上却无法适配,为了开展工作有必要了解非常清楚整个编译的机制。此外在研究过程中发现openai 的triton,以及tvm等一些ai框架对nvidia的支持原理都及其类似,所以了解原理更加有必要。原创 2024-07-31 17:43:35 · 492 阅读 · 0 评论 -
cuda中的cooperative_groups
以前block内部的同步是用syncthreads(), block之间没用提供同步的接口,因为也合理,毕竟如果block太多的话,block_n要等block_0算完退出后才会进入sm, 但是block_0为了同步又要等block_n,这样就是锁死了,因为gpu的逻辑和cpu不一样,gpu单个block寄存器的值不会暂存到显存里来切换block_0,那合格api咋用。最近看到一个代码cooperative_groups.this_grid().sync()很好奇,这里好好梳理一下。原创 2024-07-17 11:16:30 · 268 阅读 · 0 评论 -
cufftPlanMany参数说明
最近在看cufft这个库,传统的cufftPlan3d()这种plan接口逐渐被nvidia舍弃了,说是要用最新的cufftPlanMany,这个函数呢又依赖一个什么Advanced Data Layout(),最终把这个api搞得乌烟瘴气很难理解,为了理解自己写了一些测试来验证各个参数的意思,这里简单做一下总结。下面是函数声明以及对应的参数解释,看不懂的话可以结合后面的例子琢磨琢磨。图示如下,这里没有画输出,意思和输入一样,接下来做个实验试试。原创 2024-05-08 17:44:37 · 552 阅读 · 0 评论 -
cuda cache相关知识总结
I was reading about the L1 and L2 caches load and store and I have found that if there is a miss in L1 for a load instruction, L1 will get only the sector (32byte) of the 128 cache line from L2. But why do we say that the granularity of a fetching is 128by原创 2023-10-27 11:41:33 · 136 阅读 · 0 评论 -
single-passParallel Prefix Scan with Decoupled Look-back
最近在写基数排序,nvidia的基数排序依赖这个实现,所以有必要搞懂。原创 2023-08-01 13:37:31 · 100 阅读 · 0 评论 -
cuda中radix_sort
在second step中完全实在ScanCounters()函数中,具体分为upsweep, exclusivesum, downsweep.最终的目的是吧share memory的值修改成为下图中最右边的结果。这里的流程可能会和源码有出入,但是结果没问题。radix_sort排序是一种经典排序,在gpu上都有对其进行支持,这里主要参考cub中的实现,简单介绍一种单block的情形, 本文只适合看过源码但是没有看懂的同学。原创 2023-07-12 14:28:27 · 461 阅读 · 0 评论 -
内存对齐理解
内存对齐指的是对象首地址的位置是某个数值(alignment)的整数倍, 比如按4字节对齐,那么这个对象的首地址会是4的整数倍(0,4,8 …假如一个对象要求按alignment对齐,那么我们应该保证其分配的内存地址就是alignment的整数倍?原创 2023-05-03 20:52:10 · 341 阅读 · 0 评论 -
Tensor Core编程
这里的Tensor Core是指Nvidia的显卡中的计算单元。原创 2022-12-07 14:22:16 · 648 阅读 · 0 评论 -
CUDA中Occupancy相关知识
本篇文章适合稍微有一些cuda基础的朋友阅读。原创 2022-11-04 13:26:16 · 1671 阅读 · 0 评论 -
CUDA中的名词
在写kernel分析或者看一些博客的时候经常遇到一些名词,这里记录一下。原创 2022-10-24 00:33:27 · 426 阅读 · 1 评论 -
GPU中的内存相关概念
在缓存中有一个概念叫做cache line ,可以理解为一个内存单元大小,比如一个cache line是64字节的缓存L1, 如果L1的缓存大小是512字节,那么一共有8个单元(cache line)1个sector是global memory上的最小内存访问单元,也就是32字节。原创 2022-09-05 12:36:18 · 998 阅读 · 0 评论 -
bert深度优化
在当下cv, nlp领域,一般都会使用attention结构,对attention的优化就显得额外重要,这边文章就nvidia对attention的优化策略进行分析。原创 2022-08-15 19:39:14 · 596 阅读 · 0 评论 -
cuda的warp scheduler知识
在cuda中,每32个线程会被『捆』成一束–线程束,英文是warp, 一个warp执行一个指令,换句话说32个线程每次都是执行想用的指令。这里对指令发射的任务是有warp scheduler来完成的,具体工作原理如下:假如sm中有1个warp scheduler(线程束调度器),调度器最多可以管理8个warp。下图做了一个简单的展示,这里的slot个数就是调度器可以管理的warp最大值,为了后面解释方便,用涂有颜色的块表示某个slot被使用,空白则表示没有被使用。在上面的warp slot解释中,其实可原创 2022-07-07 20:15:58 · 2718 阅读 · 3 评论 -
内存事务和内存指令的区别
背景当一个warp去执行一个获取内存的指令时候,很重要一点是要考虑一下线程束中每个线程访问的内存位置。详细解释A memory “request” is an instruction which accesses memory, and a “transaction” is the movement of a unit of data between two regions of memory. 这段话意思是,一个内存请求是一个获取内存的指令,一个内存事务是在两个内存范围(比如L1和全局内存)之间的数原创 2022-05-09 00:31:54 · 512 阅读 · 0 评论 -
too many resources required for launch
背景当我们在使用一些低端的jetson设备的时候,比如nano, 偶尔会出现报错,报错显示:too many resources required for launch查资料可以发现,一般遇到这种情况就是两个问题,第一个就是寄存器不足,第二个就是共享内存不足。那么问题来了,我实现的一个kernel压根一点没用共享内存,那么就是寄存器不足,可是寄存器不足不是说可以使用显存的吗?我显存几个G为啥还说不够?以下内容仅仅适合高阶cuda玩家,普通小白直接看解决方案就行。分析什么时候kernel会由于原创 2022-04-25 14:19:53 · 2904 阅读 · 0 评论 -
CUDA Core理解
背景也接触一两年GPU了, 最近发现对最基础的cuda core 反而不认识了,具体的原因是:当在《CUDA C编程权威指南 》上看延迟隐藏的时候有很大的疑惑,为什么隐藏延迟需要的线程数=延迟*吞吐?当一个warp调度器发射一个除法指令给某一个cuda core的时候,假如需要5个时钟周期,那么应该是5个周期后才能接收下一个指令,但是根据书上的意思,cuda core 在下一个周期就可以接收另外有除法指令了,为什么?理解cuda core是可以执行 32 位浮点加法、32 位浮点乘法、32 位到 8原创 2022-04-22 18:11:09 · 2508 阅读 · 1 评论 -
share memory的bank conflict分析
背景在做高性能分析的时候,经常会出现一个什么bank conflict的名词,不仅是GPU的share memory会出现bank confict, 甚至连寄存器也会出现bank conflict, 那么这个是什么东西,下面进行一个系统的梳理。硬件从事计算机行业的同学肯定见过这个东西,没错这个玩意就是内存条,上面的小黑块就是内存颗粒,每一片称之为chip. 有兴趣的可以拆开自己的电脑看一下,这个chip大约有1mm左右的厚度,我们知道现在的半导体工艺一般都是纳米级别,所以这1mm后的内存,其实是由很原创 2022-04-21 17:01:17 · 3415 阅读 · 1 评论 -
CUDA矩阵乘法优化
前言纸上的来终觉浅,绝知此事要躬行。专业翻译为白话:https://zhuanlan.zhihu.com/p/410278370naive写法__global__ void matrixMul(const float *A, const float *B, float *C, int M, int N, int K) { int col = blockIdx.x * blockDim.x + threadIdx.x; int row = blockIdx.y * blockDim.y原创 2022-04-19 10:33:07 · 2205 阅读 · 0 评论 -
CUDA的指令发射与warp调度
概念下面的图是fermi的sm单元,可以看到里面有Warp Scheduler和Dispatch Unit,那么这两个单元是干什么的?通俗理解首先假如有n个warp是等待处理的,根据下面的图就可知道,在一个cycle中会有3个warp加载到warp_scheduler中,我们知道一个warp其实就是32个线程,说白了就是这个warp_scheduler同时可以加载32个线程的指令,接下来,每个线程可能都有很多事情要做,比如此刻有个warp需要做8个浮点数计算,从下面的图可以看到,一个warp_sch原创 2022-02-21 15:19:46 · 1695 阅读 · 0 评论 -
CUDA的L2缓存
理解当一个cuda kernel重复的获取一个global memory数据的时候,可以认为这种数据访问是持久的。如果一个数据只访问一次,那么这个数据被认为"过客"。在cuda11.0后和算力8.0以后的设备,都有能力去影响L2cache中的可持续化数据,进而潜在提高访问全局内存的带宽和降低延迟。L2 cache自留地持久化访问优先使用这块自留地,而普通或过客访问只能在持久化访问不用的时候用一下。自留地的设置如下:cudaGetDeviceProperties(&prop, device_原创 2022-02-18 15:44:23 · 1216 阅读 · 0 评论 -
CUDA--延迟隐藏
理解延迟隐藏的意思,其实就是充分利用好硬件的带宽,比如每秒的吞吐是10GB/s, 但是你就一个warp, 而warp的只是需要2GB/s, 这样就浪费了带宽,如果warp拆有比如5个,那么实际也是1s就解决了问题,好好体会。...原创 2022-02-17 15:05:54 · 825 阅读 · 0 评论 -
CUDA的合并访问和内存对齐
文章目录概述L1 cache 和L2 cachecache line情况1情况2情况3概述内存加载分为缓存加载和没有缓存加载,仔细研究这个说法有点奇怪,具体如下:缓存加载数据先经过L2 cache, 然后经过L1 cache, 最后到达线程的寄存器。非缓存加载数据先经过L2 cache, 到达线程的寄存器。L1 cache 和L2 cacheL1的取数窗口是每次可以取128bit数据,L2的取数窗口是每次可以取32bit的数据.当一个线程束要获取数据时,如果每个线程束需要3原创 2022-02-16 20:59:17 · 3445 阅读 · 0 评论 -
Welford算法解决layernorm问题
背景在利用框架做计算的时候,经常会遇到layernorm的问题,不知道有没有小伙伴发现,当fp32切到fp16的时候,有时候直接结果为nan或者为inf了,为此需要研究一下。原理其实layernorm的核心就是计算方差,定义的公式如下,但是实际上考虑到计算效率的问题,我们会采用FP32的公式来实现,具体可以节省多少计算量,有兴趣可以试一下,不过当把fp32强行切换到fp16的时候,就会出现误差,导致位置错误。welford算法之前很多框架采用的都是上面的fp32的算法,下面来看看一种新的计算方式原创 2022-02-10 12:28:22 · 2330 阅读 · 0 评论 -
CUDA的kernel并行
流并行原创 2022-01-11 20:34:49 · 4286 阅读 · 2 评论 -
CUDA Stream相关知识
流的分类流可以理解为一个管道,在GPU上运行的任何API都必须放在一个管道中,同一个管道中的API严格按照顺序执行。空流非空流我们一般写的cuda代码都是用的空流,也就是程序生成的一个默认流。...原创 2022-01-07 15:07:47 · 585 阅读 · 0 评论 -
1.了解NVIDIA显卡架构
架构排序(时间顺序):Tesla: 市面已经没有相关显卡Fermi:GeForce 400, 500, 600, GT-630Kepler:Tesla K40/K80, GeForce 700, GT-730Maxwell: Tesla/Quadro M series GeForce 900, GTX-970Pascal: Tesla p100,GTX 1080, GTX 1070, GTX 1060Votal: Tesla V100, GTX 1180Turing: T4,GTX 1660原创 2020-06-16 15:54:01 · 8460 阅读 · 2 评论 -
2. 查看自己显卡的硬件信息
下载NVIDIA_CUDA-**_Samples去github或者gitee上找一个最新仓库下载下来:git clone https://gitee.com/TimVerion/cuda-samples.git编译首先查看自己机器的cuda 版本,我自己的机器是10.2, 那么就切换到对应的tag:cd cuda-samplesgit checkout v10.2编译所有demo:make编译设备查询demo:cd Samples/deviceQuery/usr/local/cuda.原创 2021-09-06 21:38:20 · 379 阅读 · 0 评论 -
内存对齐问题
背景内存对齐可能很多程序员接触不到,也许只在面试的偶尔会被问到过,但是也只是背背固定的公式,大概知道怎么计算,也能知道大致的原理,就是数据不对齐,取数次数要变多,但是只是理解到这种程度还不够,目前intel cpu不需要对齐才能访问,但是对于一些新的arm芯片,自研芯片等等,自己在做hpc时,发现这个内存对齐问题还是一个比较严重的事情,这里好好的捋顺一下,避免自己以后遗忘。规则结构体中大小为 size 的字段,他的结构内偏移 offset 需要是其min(default, size)的整数倍 //d原创 2021-04-19 10:28:09 · 1335 阅读 · 2 评论 -
GPU的矩阵转置优化(transpose)
文章目录背景方案传统CPU转置一般GPU代码问题分析结论背景在图像处理,深度学习领域,有很多矩阵运算的工作,而伴随矩阵运算就存在大量的矩阵转置,转置不涉及计算,主要的工作都在数据的读取写入方面,所以如何加快数据搬移是一种很重要的优化点。方案传统CPU转置假设都是按行存储。int row = 1024int col = 512void transpose_CPU(vector<int>& in, vector<int>& out){ for (原创 2021-03-12 14:22:37 · 2696 阅读 · 2 评论 -
CUDA编译过程
首先当然是将.cu代码进行分离,利用是cudafe.exe 去分离cpu代码和GPU代码,在生成的中间文件可以看到test.cudafe1.cpp和test.cudafe1.gpu接下来nvopencc将根据编译 选项-arch将gpu 代码编译成对应-arch的test.ptx文件接下来ptxas.exe 编译 .ptx 到 .cubin,这个是根据编译选项-code定义的,比如test.sm_30.cubin,这一步叫做PTX离线编译,主要的目的是为了将代码编译成一个确定的计算能力和SM版本,对应.原创 2020-12-17 19:14:19 · 4675 阅读 · 4 评论 -
GPU中的SM和warp的关系
介绍我们在cpu上做好了准备工作,然后将指令发送给GPU, 在这里我们看到grid为10,block为128,也就说有10个block,每一block有128个thread,如何分配这些block?原则在GPU中一个block是不能拆分到两个SM中一个SM中的block要满足SM的条件限制SM有哪些限制?简单有这些,当然还有其他的没有展示,这里数值都是我瞎编的,具体的可以去官网文档查看参数:一个block需要哪些资源?一个SM能给的资源如上,那么一个block需要哪些资源呢?简单如原创 2020-09-04 13:10:21 · 4958 阅读 · 6 评论 -
GPU编程前的准备工作
1、看看自己计算机上有几块显卡int deviceCount;cudaGetDeviceCount(&deviceCount);int device;for(device = 0; device < deviceCount; ++device){ cudaDeviceProp deviceProp; cudaGetDeviceProperties(&deviceProp, device); printf("Device %d h...原创 2020-09-02 11:15:25 · 230 阅读 · 0 评论 -
cublasSgemmBatched使用说明
背景在GPU进行计算的时候,很多时候都需要利用cublas的API, 常用的API有两个:cublasSgemm 和cublasSgemmBatched, 使用过MKL的可能觉得很熟悉,连参数都是一样的,但是这里有一比较坑的地方是,在mkl的矩阵乘法中我们可以设置使用行优先或者列优先,考虑到很多代码底层都是c/c++写的,所以平时矩阵都是按照行优先来写的,不过mkl是支持列优先的矩阵乘法,但是cublas只支持列优先,也不知道英伟达公司是怎么想的,做成兼容的就那么难?不管怎么样,反正既然别人是制定规则的原创 2020-07-02 20:27:19 · 5488 阅读 · 2 评论 -
Nvidia显卡硬件与软件
下面这张与是GPU的硬件结构,这个图有点大,放大看可以看到很多细节,我们可以看到有各种硬件组成单元,比如GPC,TPC,SM, L2Cache,绿色小块就是cuda core也叫sp,后面这些都有专门对应的软件信息。再往里面深入一层就可以看到更加具体的细节,下面这种图就是SM的细节。可以看到sm里面又分为两个大块,每一个大块的名字叫SMP(SM Processing Block),对于浮点计算来说,CPU可以同时支持不同精度的浮点运算,但在GPU里针对单精度和双精度就需要各自独立的计算单元,一般在GPU原创 2020-06-17 15:54:40 · 917 阅读 · 0 评论 -
cuBLAS矩阵乘法
cuBLAS是cuda封装好的一个数学库,头文件为<cublas_v2.h>#define cublasSgemm cublasSgemm_v2CUBLASAPI cublasStatus_t CUBLASWINAPI cublasSgemm_v2( cublasHandle_t handle, cublasOperation_t transa, cublas...原创 2020-04-03 21:05:55 · 2464 阅读 · 2 评论