CUDA3.0新特性简览一
新的局域b[/b]同步指令：
int __syncthreads_count( int pred ) ( PTX Instruction : bar.red.popc.u32
除具有与__syncthreads同样的功能外，同时返回block内匹配pred的thread的数量。
int __syncthreads_and( int pred ) ( PTX Instruction : bar.red.and.pred )
与__syncthreads具有同样的功能外，但同时返回一个整型syncline标志，这个值只有在block内的所有thread都匹配pred时才为非零。
int __syncthreads_or( int pred ) ( PTX Instruction : bar.red.or.pred )
与__syncthreads具有同样的功能，但同时返回一个整型syncline标志，且只要block内的任一thread匹配pred时就为非零。
实际上在上述三个同步函数中的条件谓词在底层的硬件层次是Multiprocessor的所有32个CUDA core共享且在到达syncpoint之前对block内所有thread可设置的一个32位整数寄存器(为叙述方便，将此类寄存器命名为syncmask。如果使用上述任何一个同步指令，则当前block内的每个当前warp会累积当前和之前的syncmask值，假设block尺寸为64，使用指令__syncthreads_orb,[/b]则：
syncmask
warp0 : 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
Op à | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | |
warp1 : 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1
__syncthreads_or :
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1
于是返回值为非零b[/b]
或许这些指令将可以为我们提供很多有趣的实现及更复杂的同步算法，比如：
while( __syncthreads_andb ){[/b]
change pred
}
while( __syncthreads_and( __syncthreads_countb==blockDim.x ) ){[/b]
change pred
}
全速的常量内存：
在compute capability<2.0的设备上，只有一个warp内的所有thread访问同一个地址，才不会发生warp serlizaed.而在>=2.0的设备上则没有这个限制。这个功能可以有可能提高大尺寸矩阵乘法的效率。
支持递归：
Fermi 通过程序二进制接口( ABI : Application Binary Interface )机制支持设备函数参数列表堆栈，从而首次在GPU上实现了设备函数的递归调用
大幅更改的PTX ISA 2.0:
太多了，不是简单几句话可以说的完的，留待以后单独写个系列
更大的可配置共享内存/L1 Cache:
相关内容大家早已知道就不多说了，重点说下CUDA3.0重的共享内存限制。
在具备compute capability2.0的设备上，共享内存的限制在某些角度看放的更宽了，而另外的角度看确有限制的更紧了：有32个bank(和我以前的一篇文章推断的一样^^).不同线程顺序访问8位或者64位数据不会产生bank conflict,但是当前warp内的2个half-warp不再是冲突无关了，亦有可能发生bank conflict.让我们回想一下fermi 架构下新的线程调度器,每个Multiprocessor具有32个CUDA Core,有两个独立的warp scheduler，如果没有双精度和”重浮点函数“操作，则每两个时钟周期一个warp scheduler发射一组指令到线程号为偶数的thread,另一个scheduler发射一组指令到线程号为奇数的thread,相当于将bank冲突的限制“宽度”延长至1.x之前的两倍，从而可以使用类似cc1.x之前那样的技巧避免访问64位共享内存时的bank conflict,e.g:
shared double b64_smem[ 512+512/warpSize ];
b64_smem[ threadIdx.x+( threadIdx.x>>5 )+1 ]=…
附件:
External Media
g [时间:2010-3-27 09:59]
