我的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进行处理?
问的比较啰嗦,希望不要介意,谢谢
|