玉足吧 国内大厂GPU CUDA高频口试问题汇总一

发布日期:2024-10-29 11:58    点击次数:200

玉足吧 国内大厂GPU CUDA高频口试问题汇总一

1. 你不错省略说下CUDA的线程组织结构吗?

相关于GPU硬件结构来说,CUDA是一个并行想到的编程模子,它是基于GPU的体绑缚构瞎想的。为了高效地运用GPU的并行想到才略,简化并行编程玉足吧,晋升要领的性能和可膨大性,CUDA提供了一种概述的档次模子,即CUDA的线程组织结构。

CUDA的线程组织结构包括网格(grid)、线程块(block)和线程(thread)。

率先,GPU中枢要领 kernel 在 device上 实行时启动好多 Thread,而一个 kernel 所启动的扫数线程称为一个网格 grid。

其次,合并个 grid 上的线程分享相似的全局内存空间,而 grid 又不错分为好多线程块 block,线程块是向GPU进行调养的最小单元,GPU同期支持多个线程块的实行,达到上限后,惟一旧的线程块内的线程一说念实行完成后,新的线程块才会被调养入GPU。

终末,一个 block 内部包含好多 Thread,每个线程Thread都是孤独实行的,况兼不错探访全局内存和分享内存,线程之间不错通过同步原语进行同步操作。

在进行CUDA编写gpu kernel函数时,需要使用__global__修饰,复返值必须为void 。在核函数的调用时需要使用三括号的神色来指明核函数中的线程数量以及陈列情况的,如<<<1, 1>>>

三括号中的第一个数字不错看作线程块的个数,第二个数字不错看作每个线程块中的线程数,三括号中的两个数字分别等于网格大小和线程块大小,即 <<<网格大小, 线程块大小>>> 。而核函数的总线程数即为网格大小乘以线程块大小。

值得可贵的是从开普勒架构出手,最大允许的线程块大小是1024。

图片

此外,grid和block在CUDA的编程模子中也支持多维组织。

dim3 grid(2, 2);dim3 block(4, 2, 2);kernel_fun<<< grid, block >>>(prams...);

如上所示,grid 和 block 都是界说为 dim3 类型的变量。

CUDA的这种 <<<grid,block>>> 其实等于一个多级索引的表率,第一级索引是 (grid.xIdx, grid.yIdy),对应上图选中的 Block 等于(1, 1),通过它咱们就能找到了这个线程块的位置,然后咱们启动二级索引 (block.xIdx, block.yIdx, block.zIdx) 来定位到指定的 Thread。这等于CUDA的线程组织结构。

图片

图片

从硬件上来看,CUDA编程模子中的网格、线程块和线程梗概与硬件结构中的GPU、SM(流式多处理器)和SP(流式处理器)是逐一双应的。

图片

从上图不错看出,一个Grid不错包括多个SM,也不错探访Global Memory和Constant Memory;

一个Block只可在一个SM中,且一个SM包含多个Block,Block不错探访Shared Memory;

一个Block中有多个Thread,而一个Thread只可探访Registers或local Memory。

此外,一个线程块Block还不错细分红多个线程束,一个线程束(也等于一束线程)是一个线程块内部相邻的warpSize个线程,目下warpSize 都为32。

一般来说,但愿线程块的大小是warpSize的整数倍,不然系统会自动为剩下的n个线程补都32-n个线程,形成一个圆善的线程束,而这32-n个线程并不会被核函数调用,从而闲置。

2. 不错省略谈谈CUDA的存储体绑缚构,每一种存储的优短处,该何如合理使用。

CUDA的存储体绑缚构包括全局内存(Global Memory)、分享内存(Shared Memory)、常量内存(Constant Memory)/纹理内存(Texture Memory)和腹地内存(Local Memory)。

图片

全局内存(Global Memory)

这是GPU中最大的内存,即咱们常说的HBM内存,不错被扫数块上的扫数线程探访,当咱们在GPU中出手化一个值而不指定其存储位置时,它会自动存储在全局内存中。联系词,探访全局内存昔日比其他内存类型慢,因此需要进行优化以幸免性能下落,不错通过合并内存探访和使用分享内存来优化性能。

分享内存(Shared Memory)

合并个Block内的线程不错通过分享内存分享数据。比较探访全局内存至少快个10倍,但分享内存的容量有限(昔日为几十KB),无法被其他线程块探访。由于分享内存暖热存内存提供快速的探访速率,因此咱们昔日在想到过程中使用它们来存储数据。典型的表率是率先将所罕有据从 CPU 复制到 GPU 并将其存储在全局内存中。然后,咱们将数据瓦解成更小的部分(块)并将它们推送到分享内存中进行想到。想到完成后,摈弃将被推回全局内存。

纹理内存和常量内存(Texture and Constant Memory):

这些是GPU 中的特殊内存类型,针对探访特定数据类型(举例玉足吧纹理或常量值)进行了优化。扫数块中的扫数线程都不错探访这些内存类型。

举例,常量内存极端只可用于存储只读数据,纹理内存只可用于存储二维图像数据,这两种内存类型的探访速率都卓著快,不错与分享内存相比好意思。

因此,使用纹理内存和常量内存的目的是优化数据探访并减少分享内存的想到负载。咱们不错将一部分数据分派给纹理内存和常量内存,而不是将所罕有据推送到分享内存中。这种分派计策通过运用纹理内存和常量内存的优化探访功能来匡助增强内存性能。

腹地内存(Local Memory)

每个线程都不错使用我方的腹地内存,不错在其中存储临时变量。它具有最小的范围况兼专用于每个单独的线程。

3. 你了解CUDA stream吗?为什么要使用多个stream?

stream卓著于是GPU上的任务队伍,用官方的话叫作念一条敕令活水线,它允许多个CUDA操作在不同的stream中并行实行,从而晋升GPU的运用率和性能。

每个kernel调用或大多数CUDA API都不错指定联系到某一个stream,合并个stream的任务是严格保证规矩的,上一个敕令实行完成才会实行下一个敕令。

不同stream的敕令不保证任何实施规矩,部分优化技巧需要用到多个stream才能罢了。如在实行kernel的同期进行数据拷贝,需要一个stream实行kernel,另一个stream进行数据拷贝,或者针对大数据集进行切分,然后不错选择多个stream实行并行加速拷贝。

此外,多个stream还大略轻便地鉴别和科罚不同的任务,晋升应用要领的机动性和可膨大性。

图片

4. GPU和CPU分别合适实行哪些要领?联结它们的硬件架构证明一下为什么它们有各自的上风。

图片

CPU 将大宗芯单方面积极端用于可减少辅导延长的功能,举例大缓存、更少的 ALU 和更多的限度单元。

比较之下,GPU 使用大宗 ALU 来最大化其想到才略和婉曲量,它们使用相等一丝的芯片区域算作缓存和限度单元,使得其具有很高的延长。

是以关于他们的瞎想主张来说,CPU 被瞎想为规矩实行辅导,为了晋升规矩实行性能,多年来 CPU 瞎想中引入了许多功能,包括辅导活水线、乱序实行、推测实行和多级缓存等。CPU的要点是减少辅导实行延长,以便 CPU 大略尽快实行辅导序列。

GPU 专为大限度并行性和高婉曲量而瞎想,但代价是较高的辅导延长。这一瞎想场所受到了它们在视频游戏、图形、数值想到和目下深度学习中的使用的影响。扫数这些应用要领都需要以相等快的速率实行大宗线性代数和数值想到,因此东说念主们对晋升这些开辟的婉曲量干预了大宗可贵力。

此外,GPU领有高带宽的内存和极端用于并行想到的辅导集,大略更高效地处理大限度数据的并行想到任务。

5. 证明一下神经蚁集加速器与CPU、GPU的区别,他们各自有何上风?

在CPU中70%晶体管用来构建Cache,还有一部分限度单元,想到单元少,是以说CPU的中枢擅长完成多相通杂任务,重在逻辑,重在串行要领;

GPU的想到模子是单辅导、多数据SIMT处理,晶体管大部分构建想到单元,运算复杂度低,合适大限度并行想到GPU的中枢擅长完成具有省略的限度逻辑的任务,重在想到,重在并行,合适深度学习,图像处理,大数据规模。但GPU无法单独责任,必须由CPU进行限度调用才能责任。

图片

NPU,即神经蚁集加速器,在电路层模拟神经元,比较于CPU和GPU,NPU通过隆起权重罢了有储和想到一体化,从而晋升运行遵守,一条辅导完成一组神经元的处理,晋升运行遵守。NPU是师法生物神经蚁集而构建的,CPU、GPU处理器需要用数千条辅导完成的神经元处理,NPU只消一条或几条就能完成,因此在深度学习的处理遵守方面上风昭着,但它需要专用定制化罢了,而且不可松驰膨大,并欠亨用。

6. 半精度浮点数FP16各个部分的具体位数,为什么要有半精度浮点数?

半精度浮点数(FP16)是一种二进制浮点数风光,其具体位数如下:

1位标识位:用于示意正负号。

5位指数位:用于示意指数部分,取值范围为-14至15(本体上是通过偏移值来示意,即果然的指数值减去15)。

10位余数位:用于示意余数部分。

float16 最大范围是 [-65504 - 66504]

float16 能示意的精度范围是 2^−24 ,进取这个数值的数字会被平直置0;

float16和float比较恰里,总结下来等于两个原因:内存占用更少,想到更快。

内存占用更少:这个是显然可见的,通用的模子 fp16 占用的内存只需蓝本的一半。memory-bandwidth 减半所带来的克己:

模子占用的内存更小,老到的时候不错用更大的batchsize。

模子老到时,通讯量(出奇是多卡,或者多机多卡)大幅减少,大幅减少恭候时辰,加速数据的畅通。

想到更快:

目下的不少GPU硬件都有针对 fp16 的想到进行优化,举例TensorCore,半精度的想到婉曲量不错是单精度的 2-8 倍;

那既然fp16像上头说的那么好,那么是否一说念都使用 fp16 即可了呢?天然不是,全用fp16主要 存在两个问题:1. 数据溢出问题;2. 舍入错误。

7. 不错谈下TensorCore的加速旨趣吗?

率先,当谈到Tensor Core的想到速率时,与CUDA Core比较,它大略在一个时钟周期内实行多个操作。Tensor Core不错同期对两个4×4的FP16张量进行矩阵乘积想到,并将摈弃累加到另一个4×4的张量上(即D = A * B + C)。

其次,TensorCore的加速也基于夹杂精度矩阵乘法。夹杂精度并不是蚁集层面既有 FP16 又有 FP32,它指的是在底层硬件算子层面,使用半精度(FP16)算作输入和输出,使用全精度(FP32)进行中间摈弃想到从而不耗费过多精度的时代。

通过硬件上的特殊瞎想,Tensor Core 表面上不错罢了 8 倍于 FP32 Core 的想到婉曲量(Volta 和 Turing 架构),况兼莫得昭着的占用面积和功耗增多。

图片

既然Tensor Core这样好,为什么Nv不坐褥一个一说念是Tensor Core的深度学习专用GPU架构?

天然说Tensor Core是极端为加速深度学习和 AI 责任负载(举例矩阵运算)而瞎想,但目下深度学习也不可离开Cuda Core。

CUDA Core 针对多样并行想到任务进行了优化,更合适于通用并行想到任务。

率先,深度学习任务其不单是是矩阵运算,还有好多的并行想到,这就看瓶颈在那儿了。如果瓶颈时在并行想到,那么这种类型的深度学习任务可能更合适CUDA Core。

其次,Tensor Core的使用是有些限度的,关于GEMM想到其恶果很好,其次其输入和输出数据类型需如若半精度或单精度,矩阵的维度最佳是 8 的倍数。

天然Tensor Core也可用于通用并行想到任务,但它们可能不如 CUDA Core 高效。

8. MPI,OpenMP以及CUDA各自适用的加速场景。

MPI

MPI是一个跨谈话的通讯公约,支持高效轻便的点对点、播送和组播。从成见上讲,MPI应该属于OSI参考模子的第五层或者更高,他的罢了可能通过传输层的sockets和TCP掩饰大部分的层。

MPI是基于音问传递的并行编程的,用户要领运用这些接口进行进度之间的数据挪动、麇集、规约和同步。MPI表率章程了这些接口的调用表率和语义,不同的罢了(举例mpich或者openmpi)可能选择不同的优化计策。

其中,点对点通讯指的是两个进度之间的通讯,可用于限度同步或者数据传输,举例MPI_Send和MPI_Recv。

聚贯通信包括了一双多、多对一和多对多的通讯神色,常用于一组进度之间的数据交换,举例AlltoAll,Allreduce等。

它适用于需要在大限度漫衍式系统上进行高性能想到的场景,如集群想到等。

图片

既然谈到了聚贯通信,不错谈一下Ring-Allreduce吗,为什么深度学习需要Ring-Allreduce?

假定有5块GPU,每一块GPU领有圆善的模子参数不错进行forward pass和backward pass,整个的老到数据大小为K,咱们需要凭据老到数据想到出所需要的梯度进行一次迭代。

探究一个省略的同步通讯计策。率先,每张GPU领有同样的出手参数,咱们将大小为K的老到数据分为N块,也等于5块,分给每张GPU。每个GPU基于我方那一部分的数据,想到得到腹地的local gradients,然后N-1块(4块)GPU将想到所得的local gradients 发送给GPU 0,让GPU 0对扫数的local gradients进行reduce(聚集操作)得到全局的梯度,然后再将该全局梯度复返给每块GPU进行back propagation来更新每个GPU上的模子参数。

那么咱们不错想到下扫数这个词通讯过程的communication cost。率先要记取实践中,network的 bandwidth是有限的,假定每张GPU需要发送给GPU 0的通讯数据大小是1GB,咱们的network bandwidth是1GB每秒(GPU 0最多每秒收受1GB的数据),那么咱们需要4秒才不错将数据一说念发送到GPU 0上,然后想到出全局的平均梯度。咱们不错看到其通讯资本是 C * N,由于受GPU 0的network bandwidth的影响,通讯资本跟着开辟数的增多,而线性增长。

图片

比较之下,Ring Allreduce的通讯资本恒定,和开辟数量无关,十足由系统中GPU之间最慢的相连决定。

咱们将扫数开辟安排在一个逻辑环中,每个GPU应该有一个左邻和一个右邻,开辟只会从它的右邻居发送数据,并从它的左邻居吸收数据,扫数这个词想到过程通过Scatter reduce和Allgather两个通讯原语完成。

OpenMP

OpenMP 是基于分享内存时势的一种并行编程模子。

OpenMP 是以线程为基础的,其实行时势选择fork-join的神色,其中fork创建新线程或者叫醒已有的线程,join将多个线程合并。

在要领实行的时候,惟一干线程在运行,当遭逢需要并行想到的区域,会派生出线程来并行实行, 在并行实行的时候, 干线程和派生线程共同责任, 在并行代码扫尾后, 派生线程退出或者挂起,不再责任,限度经过回到单独的线程中。

图片

伦理小说在线阅读

OpenMP适用于单台想到机上的多核并行想到。通过在代码中插入辅导,开发者不错相通并行实行,并将任务分派给多个处理器中枢。OpenMP适用于需要在单个想到节点上进行并行想到的场景,如多核处理器、多线程编程等。

CUDA

CUDA(Compute Unified Device Architecture)是一种用于GPU加速想到的并行想到平台和编程模子。CUDA适用于运用GPU进行并行想到的场景。

通过编写CUDA C/C++代码,开发者不错将想到任务分派给GPU上的成百上千个并行想到单元(CUDA中枢),以罢了高效的并行想到。CUDA适用于需要大限度并行想到的科学想到、机器学习、深度学习等规模。

9. 不错说下DMA和RDMA是什么吗?以及有哪些硬件上的罢了。

率先,DMA(平直内存探访)是一种才略,允许在想到机主板上的开辟平直把数据发送到内存中去,数据搬运不需要CPU的参与。

图片

传统内存探访需要通过CPU进行数据copy来挪动数据,通过CPU将内存中的Buffer1挪动到Buffer2中。DMA时势:不错同DMA Engine之间通过硬件将数据从Buffer1挪动到Buffer2,而不需要操作系统CPU的参与,大大缩短了CPU Copy的支拨。

其次,RDMA其实从名字上就不错看出,其多了一个R,即Remote。指的是在两个或者多个想到机进行通讯的时候使用DMA, 从一个主机的内存平直探访另一个主机的内存。

图片

RDMA是一种新的平直内存探访时代,RDMA让想到机不错平直存取其他想到机的内存,而不需要经过处理器的处理。RDMA将数据从一个系统快速挪动到汉典系统的内存中,而不合操作系统形成任何影响。

在罢了上,RDMA本体上是一种智能网卡与软件架构充分优化的远端内存平直高速探访时代,通过将RDMA公约固化于硬件(即网卡)上,以及支持Zero-copy和Kernel bypass这两种路线来达到其高性能的汉典平直数据存取的主张。使用RDMA的上风如下:

零拷贝(Zero-copy) - 应用要领大略平直实行数据传输,在不波及到蚁集软件栈的情况下。数据大略被平直发送到缓冲区或者大略平直从缓冲区里吸收,而不需要被复制到蚁集层。

内核旁路(Kernel bypass) - 应用要领不错平直在用户态实行数据传输,不需要在内核态与用户态之间作念高下文切换。

不需要CPU烦闷(No CPU involvement) - 应用要领不错探访汉典主机内存而不耗尽汉典主机中的任何CPU。汉典主机内存大略被读取而不需要汉典主机上的进度(或CPU)参与。汉典主机的CPU的缓存(cache)不会被探访的内存内容所填充。

音问基于事务(Message based transactions) - 数据被处理为冲破音问而不是流,摈弃了应用要领将流切割为不同音问/事务的需求。

支持分散/团聚条件(Scatter/gather entries support) - RDMA原生态支持分散/团聚。也等于说,读取多个内存缓冲区然后算作一个流发出去或者吸收一个流然后写入到多个内存缓冲区里去。

在具体的汉典内存读写中,RDMA操作用于读写操作的汉典造谣内存地址包含在RDMA音问中传送,汉典应用要领要作念的只是在其腹地网卡中注册相应的内存缓冲区。汉典节点的CPU除在相连建筑、注册调用等以外,在扫数这个词RDMA数据传输过程中并不提供行状,因此莫得带来任何负载。

终末,RDMA 三种不同的硬件罢了。RDMA算作一种host-offload, host-bypass时代,使低延长、高带宽的平直的内存到内存的数据通讯成为了可能。目下支持RDMA的蚁集公约有:

InfiniBand(IB): 从一出手就支持RDMA的新一代蚁集公约。由于这是一种新的蚁集时代,因此需要支持该时代的网卡和交换机。

RDMA过交融以太网(RoCE): 即RDMA over Ethernet, 允许通过以太网实行RDMA的蚁集公约。这允许在表率以太网基础架构(交换机)上使用RDMA,只不外网卡必须是支持RoCE的特殊的NIC。

互联网广域RDMA公约(iWARP): 即RDMA over TCP, 允许通过TCP实行RDMA的蚁集公约。这允许在表率以太网基础架构(交换机)上使用RDMA,只不外网卡要求是支持iWARP(如果使用CPU offload的话)的NIC。不然,扫数iWARP栈都不错在软件中罢了,然则失去了大部分的RDMA性能上风。

10. 平时何如进行kernel的优化,会用到哪些器具?

率先,要优化kernel函数需要先了解GPU硬件的构造。其次,需要熟悉常见的profiler器具,主要包括Nsight System和Nsight Compute。

图片

在优化的妙技和方朝上主要祥和几个点:

1. 使用异步API

使用异步API如cudaMemcpyAsync可让GPU操作与CPU操作并行,CPU忙完后调用cudaStreamSynchronize,cudaEventWait等操作恭候GPU任务完成。

2. 优化内存与显存传输遵守

使用Pinned(page-locked) Memory晋升传输速率

通过在不同的Stream里同期分别实行kernel调用及数据传输,使数据传输与运算并行。(可贵default stream的坑)

尽量将小的数据在GPU端合成大块数据后传输

3. 优化Kernel访存遵守

晋升Global Memory访存遵守

对Global Memory的访存需要可贵合并访存(coalesced )。

warp的访存合并后,肇始地址及访存大小对都到32字节

尽量幸免跨步访存

CUDA 8.0及以上的开辟不错通过编程限度L2的访存计策晋升L2掷中率。

晋升Shared Memory的访存遵守

shared memory由32个bank构成

每个bank每时钟周期的带宽为4字节

连气儿的4字节单元映射到连气儿的bank。如0-3字节在bank0,4-7字节在bank1……字节128-131字节在bank0

若warp中不同的线程探访相似的bank,则会发生bank冲突(bank conflict),bank冲突时,warp的一条访存辅导会被拆分为n条不冲突的访存央求,缩短shared memory的有用带宽。是以需要尽量幸免bank冲突。

CUDA 11.0以上不错使用async-copy feature

4. 优化线程级并行

在SMSP责任时,某些warp会由于访存依赖、寄存器依赖等原因stall。此时warp scheduler不错选中另一个eligible warp,实行其辅导,以荫藏前一个warp的stall,使SMSP中的各个硬件资源尽量保合手重荷。但假如SMSP中扫数的warp都不在eligible情状,则硬件只可空转恭候某个warp从stall中规复(如从global中央求的数据终于回想了)。

Occupancy筹谋用来揣测SM现时activate warp数量与表面上最多支持的activate warp数量的比值。Occupancy数量越高,代表SMSP厚爱的activate warp越多,当某个warp stall时,有更多的备选warp,有更大的概率不错找到一个eligible warp。极点情况Occupancy为1/8时,SM仅4个warp,每个SMSP 1个warp,当该warp stall时,smsp莫得其它warp不错遴荐,硬件势必空转恭候。

影响Occupancy筹谋的包括以下身分:

Thread Block 线程块的大小。

每个线程块的Shared Memory使用量

每个线程使用的Register(寄存器数量)

高的Occupancy不一定代表较高的性能,如某些算法照实需要每线程128寄存器时,保合手0.5的Occupancy反而是最优遴荐。但过低的Occupancy会对性能带来较大的负面影响。

5. 辅导级优化

晋升想到访存比

GPU实行想到时,需要LDS、LDG等辅导先将数据读入寄存器,再进行想到,终末通过STS、STG等辅导将数据保存下来。

以矩阵乘法为例,先进行矩阵分块,最终拆解为每个线程想到MxK,KxN的两个小矩阵的乘法:

若两小矩阵为M=2,N=2,K=1,即2x1;1x2,终末得到2x2的矩阵算作摈弃。则读入4个float需4条辅导,想到辅导亦然4条,想到访存比4/4=1;

若两小矩阵为M=8,N=8,K=1,即8x1;1x8,终末得到8x8的矩阵算作摈弃。则读入16个float,需读取辅导16条,想到辅导8x8=64条,想到访存比64/16=4;若使用向量读(float4)每条辅导读入4个float,则读取辅导仅4条,想到访存比64/4=16

晋升想到访存比,不错让GPU的更多时钟周期用于进行想到,相对的进行数据IO占用的时钟周期更少。

晋升辅导级并行

辅导级并行基应允趣:

当代岂论是CPU如故GPU,辅导的实行都是通过活水线进行的,活水线分为多个stage,即一条辅导实行完成需要每个stage的责任都实行完成。而一个时钟周期并不是完成一条辅导实行的扫数时辰,而是每一个stage完成现时责任的时辰。活水线不错同期实行多条辅导的不同阶段。

当后续辅导的实行需要依赖前边辅导的摈弃写回寄存器,咱们说出现了寄存器依赖。此时后续辅导需要恭候第前边辅导摈弃写回寄存器才能实行,若后续辅导实行时前边辅导摈弃尚未写回寄存器,活水线会失速(stall),此时warp scheduler出手切换到其它eligible warp,若无eligible warp,则SMSP将会空转。

若后续辅导不依赖前边辅导的摈弃,则即使前边辅导未实行完毕,后续辅导也不错出手实行。出奇的,即使前序辅导是一条耗时几百周期的LDG(全局内存读取)辅导或耗时几十周期的LDS(分享内存读取)辅导,只消后续一系列辅导不依赖读取回想的数据,后续一系列辅导不错正常实行而无须恭候该LDG/LDS辅导执写回寄存器。

通过以下神色,不错晋升辅导级并行,在线程级并行够不上较好恶果的情况下,进一步晋升要领性能:

数据预取(Prefetch):数据1已读取到寄存器,使用该数据1想到前,先将后续数据2的读取辅导辐照,再实行一系列数据1的处理辅导;这样数据1的处理和数据2的读取在活水线上同期实行着。当数据1处理完成,需要处理数据2时,不错确保数据2还是存在于寄存器中,此时雷同的将数据3的读取和数据2的处理同步实行起来。

辅导重排:在存在寄存器依赖的辅导间插入富余的其它辅导,使得后续辅导实行时,前边想到辅导的摈弃已写回到寄存器。从CUDA C层面有阻塞地提供一些语句间的并行性,nvcc编译器不错一定程度上自动进行辅导重排。若对nvcc重排摈弃不闲静需要我方重排时,官方尚未怒放SASS汇编器,目下只存在一些第三方SASS汇编器器具。

晋升Register的遵守

Register File也存在bank冲突,但在CUDA C层面上莫得平直手法进行物理寄存器限度。

不错通过SASS汇编器,东说念主工进行辅导寄存器分派,以尽量摈弃register bank conflict。

不错通过SASS汇编器,为寄存器探访添加reuse标记,以尽量摈弃register bank conflict。

6. 使用TensorCore进一步加速矩阵运算

TensorCore不错用来快速进行D=A*B+C矩阵运算,提供load_matrix_sync, store_matrix_sync, mma_sync 等API。

可参考这篇著作:https://zhuanlan.zhihu.com/p/570795544

11. CPU上哪些并行优化表率?

图片

线程级并行:将任务瓦解成多个线程,运用多核处理器同期实行这些线程,从而加速任务完成速率。

SIMD辅导集:使用单辅导多数据的辅导集,同期处理多个数据,晋升向量运算的遵守。

OpenMP和MPI:这是一些常用的并行编程框架,不错匡助开发东说念主员罢了并行想到,充分运用多核处理器的性能。

数据并行:将数据分割成小块,分派给不同的处理器中枢并行处理,加速数据处理速率。

图片

12. ptx 是什么,不错深度解析下吗?

PTX(Parallel Thread Execution)是由NVIDIA推出的一种GPU要领谈话,用于编写在GPU上实行的并行要领。它是一种初级要领谈话,雷同于汇编谈话,但比汇编谈话更易于编写和阅读。

图片

PTX要领谈话的基本单元是线程块(thread block),而线程是在GPU上孤独实行的最小单元。线程块内的线程不错通过分享内存进行通讯和互助。

PTX提供了一系列基本辅导,如加载内存、存储内存、线程同步和分支跳转等。此外,PTX还提供了一些针对GPU架构的优化辅导,举例cooperative groups、shared memory和atomic operations等,这些辅导大略更好地运用GPU的想到才略。

开发者不错使用NVIDIA提供的CUDA Toolkit进行PTX要领的开发,也不错使用其他第三方器具。在编写PTX要领时,需要探究GPU的并行架构和线程调养计策,以充分进展GPU的想到才略。

13. roofline模子有什么用?何如细目最优的BLOCK_SIZE。

Roofline模子是一种用于分析想到平台性能的表面模子,由加州理工大学伯利克实验室淡薄。该模子基于想到强度(Operational Intensity,OI)和内存带宽两个筹谋,不错预计不同想到平台在不同想到强度下的表面想到上限。想到强度是指算法的想到量和数据量之比,昔日以FLOPS为单元。

在建筑Roofline模子时,需要获取想到平台的硬件参数,包括CPU频率、内存带宽、AVX512和FMA等。这些参数不错通过查询CPU手册或使用lscpu等器具获取。

Roofline模子中的表面性能峰值和表面内存带宽不错通过以下公式想到:

表面性能峰值 = 频率 * 512 * AVX数量 * FMA32/64

表面内存带宽 = 内存带宽

想到强度(OI)是指算法的想到量和数据量之比,昔日以FLOPS为单元。关于访存密集型算法,不错通过soft prefetch等时代优化;关于想到密集型算法,不错通过SIMD等时代优化。

细目最优的BLOCK_SIZE需要联结具体的硬件环境和要领性情来进行评估。一般来说,较大的BLOCK_SIZE不错晋升想到遵守,但也会增多内存探访冲突的可能性。因此,在细目最优的BLOCK_SIZE时,需要衡量想到和内存带宽之间的关系,并探究要领的并行度、数据探访时势等身分。昔日不错通过实验和性能分析来寻找最优的BLOCK_SIZE。

14. 寥落矩阵的存储风光有哪些?寥落矩阵的应用场景?寥落矩阵想到与粘稠矩阵想到有何不同?

寥落矩阵的存储风光有以下几种:

COO(Coordinate)风光:将非零元素存储为三元组 (i, j, value),其中 i 和 j 分别示意元素所在行和列的下标,value 示意元素的值。

CSR(Compressed Sparse Row)风光:将矩阵按行压缩存储为三个数组,分别示意非零元素的值、列下标和每行第一个非零元素在上述两个数组中的位置。

CSC(Compressed Sparse Column)风光:与 CSR 雷同,但按列压缩存储。

DIA(Diagonal)风光:将矩阵压缩存储为一个主对角线和些许个副对角线,每个对角线用一个数组存储。

寥落矩阵想到与粘稠矩阵想到有以下不同之处:

寥落矩阵想到昔日需要使用特定的算法和数据结构,举例寥落矩阵乘法、最小割算法等。这些算法和数据结构昔日需要探究非零元素的位置和数量,以及存储风光的遴荐等身分。

寥落矩阵想到与粘稠矩阵想到在想到复杂度上有所不同。由于寥落矩阵中大部分元素为零,因此寥落矩阵乘法等想到的复杂度昔日远远低于粘稠矩阵乘法。

寥落矩阵想到与粘稠矩阵想到在硬件罢了上也有所不同。举例,在GPU等加速器上,粘稠矩阵想到昔日不错通过SIMD等时代获取较高的想到遵守,而寥落矩阵想到需要使用特定的算法和数据结构来充分运用硬件资源。

15. 何如想到CPU辅导的婉曲量和时延?

在先容何如想到CPU辅导的婉曲量和时延前,咱们先来了解下一些基本的成见:

时钟频率:这是CPU的基本责任频率,以赫兹(Hz)为单元。举例,一个500MHz的CPU,那时钟频率为500,000,000Hz。

时钟周期:这是CPU完成一个基本操作所需的时辰,昔日等于时钟频率的倒数。举例,关于500MHz的CPU,那时钟周期为1 / 500,000,000 = 2纳秒。

机器周期:这是CPU完成一个辅导所需的时辰,昔日包含多个时钟周期。举例,一个辅导可能需要3个机器周期才能完成。

辅导婉曲量:这是CPU每秒钟大略实行的辅导数,昔日以每秒百万条辅导(MIPS)为单元。想到公式为:辅导婉曲量 = 时钟频率 / 辅导周期。

辅导时延:这是CPU实行一个辅导所需的时辰,昔日以时钟周期为单元。想到公式为:辅导时延 = 机器周期 * 时钟周期。

图片

举例,假定一个500MHz的CPU,每4个时钟周期构成一个想到机周期,实行一条辅导平均需要三个机器周期,那么该处理器的一个机器周期为8纳秒,平均实行速率为41.67 MIPS。

如果以为这篇著作对你有所匡助,请点一下赞或者在看,是对我的坚信和支持~ 本站仅提供存储行状,扫数内容均由用户发布,如发现无益或侵权内容,请点击举报。