​CUDA学习笔记(六)Warp解析

news/2024/7/24 14:15:18/

本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/,仅用于学习。

Warp

逻辑上,所有thread是并行的,但是,从硬件的角度来说,实际上并不是所有的thread能够在同一时刻执行,接下来我们将解释有关warp的一些本质。

Warps and Thread Blocks

warp是SM的基本执行单元。一个warp包含32个并行thread,这32个thread执行于SMIT模式。也就是说所有thread执行同一条指令,并且每个thread会使用各自的data执行该指令。

block可以是一维二维或者三维的,但是,从硬件角度看,所有的thread都被组织成一维,每个thread都有个唯一的ID(ID的计算可以在之前的博文查看)。

每个block的warp数量可以由下面的公式计算获得:

 

一个warp中的线程必然在同一个block中,如果block所含线程数目不是warp大小的整数倍,那么多出的那些thread所在的warp中,会剩余一些inactive的thread,也就是说,即使凑不够warp整数倍的thread,硬件也会为warp凑足,只不过那些thread是inactive状态,需要注意的是,即使这部分thread是inactive的,也会消耗SM资源。

Warp Divergence

控制流语句普遍存在于各种编程语言中,GPU支持传统的,C-style,显式控制流结构,例如if…else,for,while等等。

CPU有复杂的硬件设计可以很好的做分支预测,即预测应用程序会走哪个path。如果预测正确,那么CPU只会有很小的消耗。和CPU对比来说,GPU就没那么复杂的分支预测了(CPU和GPU这方面的差异的原因不是我们关心的,了解就好,我们关心的是由这差异引起的问题)。

这样我们的问题就来了,因为所有同一个warp中的thread必须执行相同的指令,那么如果这些线程在遇到控制流语句时,如果进入不同的分支,那么同一时刻除了正在执行的分之外,其余分支都被阻塞了,十分影响性能。这类问题就是warp divergence。

请注意,warp divergence问题只会发生在同一个warp中。

下图展示了warp divergence问题:

为了获得最好的性能,就需要避免同一个warp存在不同的执行路径。避免该问题的方法很多,比如这样一个情形,假设有两个分支,分支的决定条件是thread的唯一ID的奇偶性:

__global__ void mathKernel1(float *c) {int tid = blockIdx.x * blockDim.x + threadIdx.x;float a, b;a = b = 0.0f;if (tid % 2 == 0) {a = 100.0f;} else {b = 200.0f;}c[tid] = a + b;
}

一种方法是,将条件改为以warp大小为步调,然后取奇偶,如下: 

__global__ void mathKernel2(void) {int tid = blockIdx.x * blockDim.x + threadIdx.x;float a, b;a = b = 0.0f;if ((tid / warpSize) % 2 == 0) {a = 100.0f;} else {b = 200.0f;}c[tid] = a + b;
}

代码:

int main(int argc, char **argv) {
// set up device
int dev = 0;
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
printf("%s using Device %d: %s\n", argv[0],dev, deviceProp.name);
// set up data size
int size = 64;
int blocksize = 64;
if(argc > 1) blocksize = atoi(argv[1]);
if(argc > 2) size = atoi(argv[2]);
printf("Data size %d ", size);
// set up execution configuration
dim3 block (blocksize,1);
dim3 grid ((size+block.x-1)/block.x,1);
printf("Execution Configure (block %d grid %d)\n",block.x, grid.x);
// allocate gpu memory
float *d_C;
size_t nBytes = size * sizeof(float);
cudaMalloc((float**)&d_C, nBytes);
// run a warmup kernel to remove overhead
size_t iStart,iElaps;
cudaDeviceSynchronize();
iStart = seconds();
warmingup<<<grid, block>>> (d_C);
cudaDeviceSynchronize();
iElaps = seconds() - iStart;
printf("warmup <<< %4d %4d >>> elapsed %d sec \n",grid.x,block.x, iElaps );
// run kernel 1
iStart = seconds();
mathKernel1<<<grid, block>>>(d_C);
cudaDeviceSynchronize();
iElaps = seconds() - iStart;
printf("mathKernel1 <<< %4d %4d >>> elapsed %d sec \n",grid.x,block.x,iElaps );
// run kernel 3
iStart = seconds();
mathKernel2<<<grid, block>>>(d_C);
cudaDeviceSynchronize();
iElaps = seconds () - iStart;
printf("mathKernel2 <<< %4d %4d >>> elapsed %d sec \n",grid.x,block.x,iElaps );
// run kernel 3
iStart = seconds ();
mathKernel3<<<grid, block>>>(d_C);
cudaDeviceSynchronize();
iElaps = seconds () - iStart;
printf("mathKernel3 <<< %4d %4d >>> elapsed %d sec \n",grid.x,block.x,iElaps);
// run kernel 4
iStart = seconds ();
mathKernel4<<<grid, block>>>(d_C);
cudaDeviceSynchronize();
iElaps = seconds () - iStart;
printf("mathKernel4 <<< %4d %4d >>> elapsed %d sec \n",grid.x,block.x,iElaps);
// free gpu memory and reset divece
cudaFree(d_C);
cudaDeviceReset();
return EXIT_SUCCESS;
}

编译运行:

$ nvcc -O3 -arch=sm_20 simpleDivergence.cu -o simpleDivergence
$./simpleDivergence

输出:

$ ./simpleDivergence using Device 0: Tesla M2070
Data size 64 Execution Configuration (block 64 grid 1)
Warmingup elapsed 0.000040 sec
mathKernel1 elapsed 0.000016 sec
mathKernel2 elapsed 0.000014 sec

我们也可以直接使用nvprof(之后会详细介绍)这个工具来度量性能:

$ nvprof --metrics branch_efficiency ./simpleDivergence

输出为:

Kernel: mathKernel1(void)
1 branch_efficiency Branch Efficiency 100.00% 100.00% 100.00%
Kernel: mathKernel2(void)
1 branch_efficiency Branch Efficiency 100.00% 100.00% 100.00%

Branch Efficiency的定义如下:

到这里你应该在奇怪为什么二者表现相同呢,实际上当我们的代码很简单,可以被预测时,CUDA的编译器会自动帮助优化我们的代码。稍微提一下GPU分支预测(理解的有点晕,不过了解下就好),这里,一个被称为预测变量的东西会被设置成1或者0,所有分支都会得到执行,但是只有预测值为1时,才会得到执行。当条件状态少于某一个阈值时,编译器会将一个分支指令替换为预测指令,因此,现在回到自动优化问题,一份较长的代码就会导致warp divergence了。

可以使用下面的命令强制编译器不优化(貌似不怎么管用):

$ nvcc -g -G -arch=sm_20 simpleDivergence.cu -o simpleDivergence

Resource Partitioning

一个warp的context包括以下三部分:

  1. Program counter
  2. Register
  3. Shared memory

再次重申,在同一个执行context中切换是没有消耗的,因为在整个warp的生命期内,SM处理的每个warp的执行context都是on-chip的。

每个SM有一个32位register集合放在register file中,还有固定数量的shared memory,这些资源都被thread瓜分了,由于资源是有限的,所以,如果thread比较多,那么每个thread占用资源就叫少,thread较少,占用资源就较多,这需要根据自己的要求作出一个平衡。

资源限制了驻留在SM中blcok的数量,不同的device,register和shared memory的数量也不同,就像之前介绍的Fermi和Kepler的差别。如果没有足够的资源,kernel的启动就会失败。

当一个block或得到足够的资源时,就成为active block。block中的warp就称为active warp。active warp又可以被分为下面三类:

  1. Selected warp
  2. Stalled warp
  3. Eligible warp

SM中warp调度器每个cycle会挑选active warp送去执行,一个被选中的warp称为selected warp,没被选中,但是已经做好准备被执行的称为Eligible warp,没准备好要执行的称为Stalled warp。warp适合执行需要满足下面两个条件:

  1. 32个CUDA core有空
  2. 所有当前指令的参数都准备就绪

CUDA编程中应该重视对计算资源的分配:这些资源限制了active warp的数量。因此,我们必须掌握硬件的一些限制,为了最大化GPU利用率,我们必须最大化active warp的数目。

Latency Hiding

指令从开始到结束消耗的clock cycle称为指令的latency。当每个cycle都有eligible warp被调度时,计算资源就会得到充分利用,基于此,我们就可以将每个指令的latency隐藏于issue其它warp的指令的过程中。

和CPU编程相比,latency hiding对GPU非常重要。CPU cores被设计成可以最小化一到两个thread的latency,但是GPU的thread数目可不是一个两个那么简单。

当涉及到指令latency时,指令可以被区分为下面两种:

  1. Arithmetic instruction
  2. Memory instruction

顾名思义,Arithmetic  instruction latency是一个算数操作的始末间隔。另一个则是指load或store的始末间隔。二者的latency大约为:

  1. 10-20 cycle for arithmetic operations
  2. 400-800 cycles for global memory accesses

下图是一个简单的执行流程,当warp0阻塞时,执行其他的warp,当warp变为eligible时从新执行。

你可能想要知道怎样评估active warps 的数量来hide latency。Little’s Law可以提供一个合理的估计:

 

对于Arithmetic operations来说,并行性可以表达为用来hide  Arithmetic latency的操作的数目。下表显示了Fermi和Kepler相关数据,这里是以(a + b * c)作为操作的例子。不同的算数指令,throughput(吞吐)也是不同的。

这里的throughput定义为每个SM每个cycle的操作数目。由于每个warp执行同一种指令,因此每个warp对应32个操作。所以,对于Fermi来说,每个SM需要640/32=20个warp来保持计算资源的充分利用。这也就意味着,arithmetic operations的并行性可以表达为操作的数目或者warp的数目。二者的关系也对应了两种方式来增加并行性:

  1. Instruction-level Parallelism(ILP):同一个thread中更多的独立指令
  2. Thread-level Parallelism (TLP):更多并发的eligible threads

对于Memory operations,并行性可以表达为每个cycle的byte数目。

因为memory throughput总是以GB/Sec为单位,我们需要先作相应的转化。可以通过下面的指令来查看device的memory frequency:

$ nvidia-smi -a -q -d CLOCK | fgrep -A 3 "Max Clocks" | fgrep "Memory"

以Fermi为例,其memory frequency可能是1.566GHz,Kepler的是1.6GHz。那么转化过程为:

 

乘上这个92可以得到上图中的74,这里的数字是针对整个device的,而不是每个SM。

有了这些数据,我们可以做一些计算了,以Fermi为例,假设每个thread的任务是将一个float(4 bytes)类型的数据从global memory移至SM用来计算,你应该需要大约18500个thread,也就是579个warp来隐藏所有的memory latency。

 

Fermi有16个SM,所以每个SM需要579/16=36个warp来隐藏memory latency。

Occupancy

当一个warp阻塞了,SM会执行另一个eligible warp。理想情况是,每时每刻到保证cores被占用。Occupancy就是每个SM的active warp占最大warp数目的比例:

 

我们可以使用的device篇提到的方法来获取warp最大数目:

cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device);

然后用maxThreadsPerMultiProcessor来获取具体数值。

grid和block的配置准则:

  • 保证block中thrad数目是32的倍数。
  • 避免block太小:每个blcok最少128或256个thread。
  • 根据kernel需要的资源调整block。
  • 保证block的数目远大于SM的数目。
  • 多做实验来挖掘出最好的配置。

Occupancy专注于每个SM中可以并行的thread或者warp的数目。不管怎样,Occupancy不是唯一的性能指标,Occupancy达到当某个值是,再做优化就可能不在有效果了,还有许多其它的指标需要调节,我们会在之后的博文继续探讨。

Synchronize

同步是并行编程的一个普遍的问题。在CUDA的世界里,有两种方式实现同步:

  1. System-level:等待所有host和device的工作完成
  2. Block-level:等待device中block的所有thread执行到某个点

因为CUDA API和host代码是异步的,cudaDeviceSynchronize可以用来停住CUP等待CUDA中的操作完成:

cudaError_t cudaDeviceSynchronize(void);

因为block中的thread执行顺序不定,CUDA提供了一个function来同步block中的thread。

__device__ void __syncthreads(void);

当该函数被调用,block中的每个thread都会等待所有其他thread执行到某个点来实现同步。


http://www.ppmy.cn/news/1168951.html

相关文章

springMVC中统一异常处理@ControllerAdvice

1.在DispatcherServlet中初始化HandlerExceptionResolver 2.controller执行完成后执行processDispatchResult(processedRequest,response,mappedHandler,mv,dispatchException),有异常则处理异常 3.ExcepitonHandlerExceptionResolver中执行方法doResolveHandlerMethodExceptio…

STM32cubemx对FreeRTOS的适配(工程模板配置)

文章目录 前言一、工程的创建二、什么是CMSIS三、STM32cubemx生成的FreeRTOS工程分析总结 前言 本篇文章将带大家使用STM32cubemx对FreeRTOS进行工程模板的配置。 一、工程的创建 1.开始工程的创建&#xff1a; 2.芯片型号选择&#xff1a; 3.修改时钟为TIM8&#xff1a; …

Spring定时任务@Scheduled

在 Spring 框架中&#xff0c;可以使用定时任务来执行周期性或延迟执行的任务。Spring 提供了多种方式来配置和管理定时任务。有Java自带的java.util.Timer类&#xff0c;也有强大的调度器Quartz&#xff0c;还有SpringBoot自带的Scheduled。 在实际应用中&#xff0c;如果没有…

B-tree(PostgreSQL 14 Internals翻译版)

概览 B树(作为B树访问方法实现)是一种数据结构&#xff0c;它使您能够通过从树的根向下查找树的叶节点中所需的元素。为了明确地标识搜索路径&#xff0c;必须对所有树元素进行排序。B树是为有序数据类型设计的&#xff0c;这些数据类型的值可以进行比较和排序。 下面的机场代…

[每周一更]-(第68期):Excel常用函数及常用操作

日常工作&#xff0c;偶尔也会存在excel表格入库的情况&#xff0c;针对复杂的入库情况&#xff0c;一般都是代码编号&#xff0c;读文件-写db形式&#xff1b;但是有些简单就直接操作&#xff0c;但是 这些简单的入库不仅仅是直接入库&#xff0c;而是内容中有部分需要进行映射…

Cannot load from short array because “sun.awt.FontConfiguration.head“ is null

错误描述 在使用Easyexcel时发生了报错&#xff0c;请求返回空白 但是只在Linux上出现了该报错&#xff0c;在本地windows环境没有出现 JDK都使用的是17版本 错误原因 由于在linux上缺失Easyexcel使用的字体导致 解决办法 下载一个jdk1.8 在其jre/lib目录里复制fontconfi…

【Edabit 算法 ★☆☆☆☆☆】 Correct the Mistakes

【Edabit 算法 ★☆☆☆☆☆】 Correct the Mistakes bugs language_fundamentals algebra Instructions Fix the code in the code tab to pass this challenge (only syntax errors). Look at the examples below to get an idea of what the function should do. Examples…

学成在线第二天-查询课程、查询课程分类、新增课程接口实现以及跨域的处理思路和全局异常处理的使用以及面试题

目录 一、接口的实现 二、跨域的处理思路 三、全局异常处理 四、面试题 五、总结 一、接口的实现 1. 查询课程接口 思路&#xff1a; 典型的分页查询 按需查询 模糊查询的查询 controller&#xff1a; ApiOperation(value "课程列表", notes "课程…

DVWA(一)

环境搭建 搭建DVWA Web渗透测试靶场_dvwa 白屏-CSDN博客 Brute Force&#xff08;暴力破解&#xff09; LOW 输入账号密码 burp suite拦截请求 请求发送至intruder attack type:选择cluster bomb,将用户名和密码Add添加 payload 1 添加用户名字典&#xff0c;payload 2 添…

Python学习第3天-第一个Python程序

文章目录 前言一、创建项目二、创建程序总结 前言 下面给大家展示下经典的Hello World! 一、创建项目 二、创建程序 print("Hello World!")总结 回到顶部 学习网站 欢迎来到Python的世界&#xff01;

Elasticsearch 8.9 Master节点处理请求源码

大家看可以看ElasticSearch源码&#xff1a;Rest请求与Master节点处理流程&#xff08;1&#xff09; 这个图非常好&#xff0c;下午的讲解代码在各个类和方法之间流转&#xff0c;都体现这个图上 一、Master节点处理请求的逻辑1、节点(数据节点)要和主节点进行通讯&#xff0…

完整搭建Git服务器

文章目录 搭建Git服务器第一步&#xff0c;安装git&#xff1a;第二步&#xff0c;创建一个git用户&#xff0c;用来运行git服务第三步&#xff0c;添加SSH公钥第四步&#xff0c;初始化Git仓库&#xff1a;第五步&#xff0c;连接Git服务器 搭建Git服务器 GitHub就是一个免费…

《自控力》读后感

文章目录 前言意志力是可以锻炼的为何善行之后会有恶行“那有如何”&#xff1a;情绪低落会使人屈服于诱惑讽刺性反弹 前言 最近读完了《自控力》这本书&#xff0c;读之前&#xff0c;我抱的期望并不太高&#xff0c;本以为是碗鸡汤&#xff0c;没啥本质的效果&#xff0c;并且…

测试自动化框架

几代测试自动化框架 在开始自动化项目之前&#xff0c;您需要了解需求&#xff1b;您拥有哪些资源、需要自动化哪些内容、支持的平台等。但是&#xff0c;框架有不同类型&#xff0c;您需要仔细选择其中一种类型。您需要回答许多问题才能成功完成该过程。这里是其中的一些&…

将本地的项目上传到Gitee

目录 1.先在Gitee新建一个仓库,提交即可 2.进入到要上传的项目里面&#xff0c;右键选择 Git Bash Here 3.右键后就打开了Git命令窗口 4.配置你的用户名和邮箱(已经配置过则可跳过) 5.查看你的用户名和邮箱配置&#xff08;可不查看&#xff09; 6.输入git init指令&#…

LeetCode讲解篇之113. 路径总和 II

文章目录 题目描述题解思路题解代码 题目描述 题解思路 深度优先遍历二叉树&#xff0c;遍历的同时记录路径&#xff0c;直到遍历到叶节点&#xff0c;若路径和为targetSum则添加到结果集中 题解代码 func pathSum(root *TreeNode, targetSum int) [][]int {var res make([…

java实现多线程下载器

前言&#xff1a; &#x1f44f;作者简介&#xff1a;我是笑霸final&#xff0c;一名热爱技术的在校学生。 &#x1f4dd;个人主页&#xff1a;个人主页1 || 笑霸final的主页2 &#x1f4d5;系列专栏&#xff1a;项目专栏 &#x1f4e7;如果文章知识点有错误的地方&#xff0c;…

计算机网络第三章习题

1.假定1km长的CSMA/CD网络的数据率为1Gb/s。设信号在网络上的传播 速率为200000km/s。求能够使用此协议的最短帧长. 问题刨析: 逻辑链: 最短帧长数据传输率x争用期2τ(2倍端到端所需要的时间) 题目已经给出数据率为1Gb/s,所以我们要知道争用期2τ是多少. 端到端所需要的时间信…

【Java】迭代器的next方法

Collection 集合的遍历 概述&#xff1a;Iteration&#xff1a;迭代器&#xff0c;集合的专用遍历方式 Iterator<E> Iterator() 返回在此 collection 的元素上进行迭代的迭代器boolean hasNext() 如果返回仍有元素可以迭代&#xff0c;则返回 trueE next() 返回迭代的下一…

了解 Elasticsearch 自动生成的文档 _id:重复是一个问题吗?

Elasticsearch 中自动生成的文档 ID 当你在未指定 ID 的情况下对文档建立索引时&#xff0c;Elasticsearch 会自动为该文档生成唯一的 ID。 该 ID 是 Base64 编码的 UUID&#xff0c;由多个部分组成&#xff0c;每个部分都有特定的用途。 ID 生成过程针对索引速度和存储效率进…