用户
 找回密码
 立即注册
yuanwcj 该用户已被删除
发表于 2013-11-18 12:37:37
161107
我的n卡是GeForce Gtx TITAN,计算能力是3.5,14个SM(2688 cores),每个SM最大线程是2048,每个SM最大block数是16,每个warp线程数为32。

我现在有10*1K*1K个数据,想每个线程处理1个数据,我定义的thread大小是1024,然后block的大小是否可以定义成10*1000,这儿有几个问题想请教下:
1.   这里的1K实际上是1000,不是1024,我知道thread的大小最好是warp的倍数,所以设置的1024,这样的话,在kernel中是否需要做个if(threadIdx.x<1024)的判断,而这样的话最后一个warp中(993-1000)实际上只有8个线程,另外24个线程都浪费了,这样造成的浪费是否值得?
2.   因为block实际上是分配到不同的SM去执行的,而warp才是处理的最小单元,是否可以认为block的大小设置只要满足blockIdx的3维大小限制即可,除此之外,没有其他限制(不需要是32的倍数)?
3.   如果thread为1024实际上每个SM也只能分配2个block,对整个卡而言最多可以分配28个block,那如果这里的block大小超过了28的话该怎么理解?
4.   ice版主曾说过,“SM会从当前的resident blocks里面以warp为单位抓取线程来进行计算”。
    a,  这样是否可以理解成,因为warp才是处理的最小单元,所以哪怕block的大小足够大(不超过3维限制),也总是会被处理,所以这才是block大小能大于28的原因?
    b,是不是也是因为这个原因,当block的大小达到了一定程度,只要能充分利用warp切换掩盖各种处理延迟之后,block再增加也不会再提升加速性能?
    c,另外,可否认为cuda每次同时(每个处理时钟)可以处理14个warp(因为有14个SM),如果其中的1个或多个warp处理之后,会自动从其他就绪的warp组中再挑选warp进行处理?

问的比较啰嗦,希望不要介意,谢谢
使用道具 举报 回复
发表于 2013-11-18 12:44:50
楼主您好,关于您的几个问题:

(1)如果您的block shape是(1000,1,1)的话,您无需block内使用if (threadIdx.x < 1024)来判断,因为实际上将自动只有1000个存活的线程的。再度判断多此一举。
以及,的确最后一个warp里面有24个线程浪费了,但这个无所谓。适合您快速实现的才是最好的。一是24在1024里才导致2%的指令能力浪费。二是它目前不值得你优化,先实现了再说。

(2)您的理解正确!如上文所述,block里的线程数目最好是32的倍数,这样是为了避免执行能力浪费。但如果要较真说是否可以任意设置,那么只要不超过最大值,显然是可以的。

(3)如果总blocks超过了GPU一次能加载的blocks,则会分批加载上去。不会导致上不去的。请勿担心。

(4)关于ICE的内容,建议ICE上来亲口为您解释。

您的问题不怕啰嗦的,因为服务您是我们的荣幸。

感谢您的莅临。
使用道具 举报 回复 支持 反对
发表于 2013-11-18 12:55:22
本帖最后由 ice 于 2013-11-18 13:42 编辑

LZ您好:

您的这一句中“我定义的thread大小是1024,然后block的大小是否可以定义成10*1000”的thread应该是threads per block,而block应该是blocks per grid。

下面回答您的问题:

1:您上1024个线程,处理1000个数据,应该判断if(threadIdx.x<1000),您应该是笔误了。以及这里可以反推您使用的是一维的threads在block内部排列。这样一来最后一个warp确实排不满,但请勿忧虑,这个影响不大。

2:您这里的描述用语和一般习惯并不一致,我稍微说下。
一般的“block大小设置”是您block内线程的多少和形状安排,而您这里说的“block的大小设置只要满足blockIdx的3维大小限制即可”看上去像是在说grid中block的多少和形状安排。

如果我推测的确实是您的实际想法,那么您的理解是正确的,grid中安排block的多少和形状并无必须是32倍数的要求,只要满足grid中安排block的各个方向的限制即可。

3:如果您的threads per block为1024,那么您的每个SMX上同时可以最多resident 2个blocks。但这并不意味着您同时只能分配28个blocks,您可以分配blocks的数量远远多于这个,并且这个数量只和计算能力版本有关,与SM多少无关,具体到您的TITAN显卡,这个值为(2^31-1)*65535*65535个。
他们会按照一定顺序依次去SMX上执行的。

前三个问题回答如上。
--------------------------------------------------------------------------------------

修改了一处描述细节,使之更加贴切。
使用道具 举报 回复 支持 反对
发表于 2013-11-18 13:32:44
本帖最后由 ice 于 2013-11-18 13:43 编辑

本楼回复第四点:

a:您这里面又混淆了“block大小”的概念,我来详细说一下。
承袭前文,您这里的“block大小”指的是“blocks per grid”的大小,正如前面所说,这个值只和计算能力版本有关,而与SMX多少无关,哪怕是GK208核心这样的渣卡也和geforce titan这样的旗舰一样,能在一个grid中放置(2^31-1)*65535*65535个blocks的,这个是kepler架构决定的调度能力。以及这些blocks会依次使用SMX,一批结束了上下一批,直至完全结束。

而您引用的我的那段话说的是每个block在resident在SMX上的时候,SMX的scheduler调度的情况,这个只和当前resident在SMX上的block有关,和grid内一共上多少个blocks无关的。

因此您这里的因果关系并不成立,也不是我之前表达的意思。

b:您可以这样理解:当您的SMX上的resident threads数量/resident warps数量比较充足的时候,可以依靠硬件自动通过线程切换来掩盖处理延迟,达到较好的效果。这里与threads per block数量以及blocks per grid数量并非直接相关。

对于整个GPU而言,还需要blocks per grid数量足够大,才能保证每个SMX都有饱满的工作量,特别是您这种SMX数量较多的旗舰显卡。

当年的GPU跑满的时候,再增加工作量,将以接近线性的比例同步增加工作时间。

以及一般情况下,在设计算法实现的时候,需要合理安排threads per block的大小,以及其他资源的使用情况,以及保证有足够的工作量,以利GPU效能充分发挥。

c:这种认为是不正确的,虽然某特定GPU可能是14SMX的,但是同时(按照时钟周期角度)并不是只有14个warp在执行的,仅以SM3.5的titan为例,每个SMX,有4个scheduler,每周期最多可以发射4个warp的8条指令到SMX上的不同执行单元。

同时需要说明的是,warp不是一直在SP中全部执行完毕才退出的,当前的一条或者两条指令执行完毕以后,会接受scheduler的下次调度,可能继续执行后续的无关指令,也可能进入等待队列,具体调度方法并不公开。

大致如此,您的第四个问题详细解释如上。

欢迎您常来论坛讨论问题,尽管问即可,无须有任何顾虑。

祝您好运~
--------------------------------------------------------------------------
修改了一处因为拼音输入法造成的错别字。
使用道具 举报 回复 支持 反对
发表于 2013-11-18 15:56:20
本帖最后由 yuanwcj 于 2013-11-18 16:04 编辑
ice 发表于 2013-11-18 13:32
本楼回复第四点:

a:您这里面又混淆了“block大小”的概念,我来详细说一下。


非常谢谢玫瑰和ice两位版主的耐心答复。

您刚才提到了“仅以SM3.5的titan为例,每个SMX,有4个scheduler,每周期最多可以发射4个warp的8条指令到SMX上的不同执行单元。”
我还想再请教下这个和gpu的峰值处理能力之间的关系是怎样的。
我先前了解的gpu的峰值处理能力=cores*clock rate*2,即2688 * 875.5MHz * 2,而实际上我们在处理的时候和core的具体个数并没有任何关系,那这个公式是怎么得到的,还有这个公式和您提到的这句话是什么关系,以及从您这句话是否也可以推导出gpu的处理能力
使用道具 举报 回复 支持 反对
发表于 2013-11-18 16:21:14
yuanwcj 发表于 2013-11-18 15:56
非常谢谢玫瑰和ice两位版主的耐心答复。

您刚才提到了“仅以SM3.5的titan为例,每个SMX,有4个scheduler ...

LZ您好:

不客气的,下面将继续补充说明您5#的问题:

简单地说,我的那句话“仅以SM3.5的titan为例,每个SMX,有4个scheduler,每周期最多可以发射4个warp的8条指令到SMX上的不同执行单元。”是从发射指令的角度上讲的,此时只考虑指令发射能力,而不考虑具体的执行情况以及不考虑您kernel中实际的指令配比情况和双发射情况,也不考虑其他的一些损耗。
说这句的目的是为了向您说明一个SMX上同时并不是只有一个warp(的指令)在计算(发布到执行单元)。

而您说的GPU峰值的计算的公式实际上是从执行端考虑得出的,此时也不考虑指令发射的情况,以及指令配比的情况等。

所以两者是不相同的。

但是两者是互相配合的,因而是相关的。

一个SMX中,有192个SP/CUDA CORE,还有SFU和LSU等其他单元,scheduler负责将warp的当前指令或者当前两条指令(双发射的情况)发布到这些执行单元上。192个SP按照同时可以容纳的指令计算相当于6个warps每个线程一条指令,而4个scheduler在全部双发射的时候可以发射4个warps每个warp两条无关指令,这相当于8个warps每个warp一条指令。

因此发射能力相对于SP数量是够用的,考虑到SMX中还有其他的执行单元,以及实际也并非全部是一周期一条这样吞吐量的指令需要计算,SMX内部总体的发射能力和执行能力还是匹配的。

以及在计算峰值的时候,是和SP/CUDA CORE数量有关的,而不是您说的他们无关。我那句话里面的“执行单元包括SP但不限于SP”,希望没有对您造成误导。

关于GPU浮点峰值的详细讨论,请您参考本版近期的详细讨论帖:

http://cudazone.nvidia.cn/forum/ ... &extra=page%3D1

大致如此,未及之处,欢迎您继续提问。

祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-11-18 17:18:29
ice 发表于 2013-11-18 16:21
LZ您好:

不客气的,下面将继续补充说明您5#的问题:

再次谢谢ice版主的回复,我还会回来的。。。。
使用道具 举报 回复 支持 反对
发表于 2013-11-18 17:21:48
yuanwcj 发表于 2013-11-18 17:18
再次谢谢ice版主的回复,我还会回来的。。。。

不客气的,cudazone的大门永远为您敞开~
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册