相干于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之间最慢的连络决定。
app咱们将统统开垦安排在一个逻辑环中，每个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. 优化内存与显存传输服从
1. 横滨水手最早成立于1972年，球队与鹿岛鹿角是仅有的两支一直在顶级联赛踢球的队伍。俱乐部历史上获得过4次联赛冠军、2次天皇杯冠军、1次联赛杯冠军等荣誉。
1. 川崎前锋最早成立于1955年，球队前身为富士通足球俱乐部，是日本足球联盟元老俱乐部之一。由于部分原因俱乐部在上世纪遭到降级，直到2000年才重返顶级联赛。在2017年之后，队伍的整体表现日渐强大，先后在2017、2018、2020以及2021年获得联赛冠军。随后还夺得了日本天皇杯冠军、日本联赛杯冠军、以及日本超级杯。
使用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。
如果以为这篇著述对你有所匡助，请点一下赞或者在看，是对我的笃信和相沿~ 本站仅提供存储处事，统统内容均由用户发布，如发现无益或侵权内容，请点击举报。