CUDA-GPU programming Introduction (3)

2024-06-20 05:58

本文主要是介绍CUDA-GPU programming Introduction (3),希望对大家解决编程问题提供一定的参考价值,需要的开发者们随着小编来一起学习吧!

关于提高performance的一些建议:
Important caveat:number of threads

并不是越多并行线程效率越高,因为每个线程都消耗一定的resource,主要是register和shared memory。所以开出再多的线程,GPU也只能在有限的资源下让一部分并行。优化应该根据资源需求。

unavoidable bottleneck: transfer between cpu and gpu

CPU和GPU之间的transfer是一个很大的bottleneck,而且长时间存在且无法避免,因为GPU能力有限,无法独立作业,必须依靠CPU来调用。简单的例子就是GPU无法access file system。 针对这个瓶颈,基本的两个处理就是,使用pinned memory,以及asynchronous transfers(overlapping computation and transfer)。

Optimizing access to global memory

GPU有很多的core来进行大量的计算,但是数据必须从global memory获取,如果当core的计算量少于和global memory之间的传输量,瓶颈就会出现,因为这时候大部分的时间都花在data的传输上,GPU的使用变得不划算。而且这种瓶颈在很多问题无法避免。
Utilizing the memory architecture effectively tends to be the biggest challenge in CUDA algorithms

GPU和global memory之间的传输可以有很高的带宽high bandwidth,但同时也有很高的延迟,high latency。所以读写内存的方式很重要。

Using many threads, latency can be overcome by hiding it among many threads. The pattern of global memory access is also very important, as cache size of the GPU is very limited.

Global memory access is fast when coalesced
It is best for adjacent threads belonging to the same warp (group of 32 threads) to be accessing locations adjacent in memory (or as close as possible)
• Good access pattern: thread i accesses global memory array member a[i]
• Inferior access pattern: thread i accesses global memory array member as a[i*nstride] where nstride >1
• Clearly, random access of memory is a particularly bad paradigm on the GPU

但有些问题本质上就有不连续的内存读写,使得优化变得困难。
典型案例:矩阵旋转 matrix transpose
A bandwidth-limited problem that is dominated by memory access。
这里写图片描述

一个基本的写法就是:

__global__ void transpose_naive(float *odata, float *idata, int width,int height)
{int xIndex, yIndex, index_in, index_out;xIndex = blockDim.x * blockIdx.x + threadIdx.x;yIndex = blockDim.y * blockIdx.y + threadIdx.y;if (xIndex < width && yIndex < height){index_in = xIndex + width * yIndex;index_out = yIndex + height * xIndex;odata[index_out] = idata[index_in];}
}

实际上数组在内存都是连续存储的,二维还是三维都只是直观上的表现,上述代码在内存中的实际表现如下:
这里写图片描述

这样看,不连续性就暴露的很彻底。这个问题还是可以解决的,因为我们还可以利用shared memory,因为shared memory不需要coalesced 读写,即使不是coalesced的也比global memory下的要快很多。尤其是需要多次访问一些数据的时候,放在shared memory比较好。但shared memory就是比较小,一般就是48kB或者16kB,而且必须要自己写代码的时候调整。

Each multiprocessor has some fast on-chip shared memory
• Threads within a thread block can communicate using the shared memory
• Each thread in a thread block has R/W access to all of the shared memory allocated to a block
• Threads can synchronize using the intrinsic __syncthreads();

这里写图片描述

具体代码如下:

__global__ void transpose(float *odata, float *idata,
int width, int height)
{__shared__ float block[BLOCK_DIM][BLOCK_DIM];unsigned int xIndex, yIndex, index_in, index_out;/* read the matrix tile into shared memory */xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;if ((xIndex < width) && (yIndex < height)){index_in = yIndex * width + xIndex;block[threadIdx.y][threadIdx.x] = idata[index_in];}__syncthreads();/* write the transposed matrix tile to global memory */xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;if ((xIndex < height) && (yIndex < width)){index_out = yIndex * height + xIndex;odata[index_out] = block[threadIdx.x][threadIdx.y];}
}

这里面,shared memory就是个中介的作用,让core和global memory之间不可避免的非连续读写在它这完成。单独说明一下odata的index计算。我们发现,xIndex和yIndex在计算block的位置的时候是符合转置关系的,但是在block内的位置依旧是原来的x,y关系,因为这里必须保证连续性,所以把转置关系留在shared memory里来做。同时补充一点,这里的x,y是一般意义上的x,y方向,即x是row方向,y是column方向。所以在一般内存里矩阵寻址的时候还是应该[y][x]。

这里写图片描述
这里写图片描述

同时,对于高度优化的代码还应该考虑bank conflict,因为shared memory中分了32个bank, 可以4 bytes和8 bytes,对应字长的。详细见另一篇博客。32是为了对应一个warp size,保证同时进来的warp里的thread不会落入同一个bank,不会有冲突。简单的解决方法就是把32x32大小的block加一个padding,32x33,这样就把可能有的冲突错开了。这个矩阵转置有冲突就是体现在从shared memory把内存拷到global memory的时候都是按列读取的,全在同一个bank,这样的话这些threads都从并行变成了serial处理,latency很高。简单处理后的结果是:

__global__ void transpose(float *odata, float *idata,
int width, int height)
{__shared__ float block[BLOCK_DIM][BLOCK_DIM + 1];unsigned int xIndex, yIndex, index_in, index_out;/* read the matrix tile into shared memory */xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;if ((xIndex < width) && (yIndex < height)){index_in = yIndex * width + xIndex;block[threadIdx.y][threadIdx.x] = idata[index_in];}__syncthreads();/* write the transposed matrix tile to global memory */xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;if ((xIndex < height) && (yIndex < width)){index_out = yIndex * height + xIndex;odata[index_out] = block[threadIdx.x][threadIdx.y];}
}

Higher dimensional coalesced access:
这里写图片描述
所以任何时候是否是coalesced access就看相邻的core的对应的相邻的thread是否是access相邻的内存位置。

这篇关于CUDA-GPU programming Introduction (3)的文章就介绍到这儿,希望我们推荐的文章对编程师们有所帮助!



http://www.chinasem.cn/article/1077295

相关文章

AI Toolkit + H100 GPU,一小时内微调最新热门文生图模型 FLUX

上个月,FLUX 席卷了互联网,这并非没有原因。他们声称优于 DALLE 3、Ideogram 和 Stable Diffusion 3 等模型,而这一点已被证明是有依据的。随着越来越多的流行图像生成工具(如 Stable Diffusion Web UI Forge 和 ComyUI)开始支持这些模型,FLUX 在 Stable Diffusion 领域的扩展将会持续下去。 自 FLU

如何用GPU算力卡P100玩黑神话悟空?

精力有限,只记录关键信息,希望未来能够有助于其他人。 文章目录 综述背景评估游戏性能需求显卡需求CPU和内存系统需求主机需求显式需求 实操硬件安装安装操作系统Win11安装驱动修改注册表选择程序使用什么GPU 安装黑神话悟空其他 综述 用P100 + PCIe Gen3.0 + Dell720服务器(32C64G),运行黑神话悟空画质中等流畅运行。 背景 假设有一张P100-

GPU 计算 CMPS224 2021 学习笔记 02

并行类型 (1)任务并行 (2)数据并行 CPU & GPU CPU和GPU拥有相互独立的内存空间,需要在两者之间相互传输数据。 (1)分配GPU内存 (2)将CPU上的数据复制到GPU上 (3)在GPU上对数据进行计算操作 (4)将计算结果从GPU复制到CPU上 (5)释放GPU内存 CUDA内存管理API (1)分配内存 cudaErro

PyInstaller问题解决 onnxruntime-gpu 使用GPU和CUDA加速模型推理

前言 在模型推理时,需要使用GPU加速,相关的CUDA和CUDNN安装好后,通过onnxruntime-gpu实现。 直接运行python程序是正常使用GPU的,如果使用PyInstaller将.py文件打包为.exe,发现只能使用CPU推理了。 本文分析这个问题和提供解决方案,供大家参考。 问题分析——找不到ONNX Runtime GPU 动态库 首先直接运行python程序

麒麟系统安装GPU驱动

1.nvidia 1.1显卡驱动 本机显卡型号:nvidia rtx 3090 1.1.1下载驱动 打开 https://www.nvidia.cn/geforce/drivers/ 也可以直接使用下面这个地址下载 https://www.nvidia.com/download/driverResults.aspx/205464/en-us/ 1.1.3安装驱动 右击,

Kubernetes的alpha.kubernetes.io/nvidia-gpu无法限制GPU个数

问题描述: Pod.yaml文件中关于GPU资源的设置如下: 然而在docker中运行GPU程序时,发现宿主机上的两块GPU都在跑。甚至在yaml文件中删除关于GPU的请求,在docker中都可以运行GPU。 原因: 上例说明alpha.kubernetes.io/nvidia-gpu无效。查看yaml文件,发现该docker开启了特权模式(privileged:ture): 而

GPU池化赋能智能制造

2023年3月10日,“第六届智能工厂高峰论坛”在杭州隆重揭幕。本次会议由e-works数字化企业网、浙江制信科技有限公司主办,中国人工智能学会智能制造专业委员会、长三角新能源汽车产业链联盟、长三角(杭州)制造业数字化能力中心、浙江省智能工厂操作系统技术创新中心协办。趋动科技作为钻石合作伙伴出席了本次峰会,与制造业精英企业以及行业专业人士共同分享制造业在智能工厂推进过程中的成功经验,探讨工厂改进中

【linux 常用命令】查看gpu、显卡常用命令

1.查看显卡基本信息 lspci | grep -i nvidia 2.查看显卡驱动版本 nvidia-smi -a 3.查看gpu使用情况 nvidia-smi (spam) [dongli@dt-gpu-1 train]$ nvidia-smi Fri Sep 27 16:42:33 2019 +----------------------------------------

CUDA:用并行计算的方法对图像进行直方图均衡处理

(一)目的 将所学算法运用于图像处理中。 (二)内容 用并行计算的方法对图像进行直方图均衡处理。 要求: 利用直方图均衡算法处理lena_salt图像 版本1:CPU实现 版本2:GPU实现  实验步骤一 软件设计分析: 数据类型: 根据实验要求,本实验的数据类型为一个256*256*8的整型矩阵,其中元素的值为256*256个0-255的灰度值。 存储方式: 图像在内存中