CUDA3.0新特性简览一

新的局域( CTA/block )同步指令：

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_or( threadIdx.x>=63 ),则：

                                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

于是返回值为非零(不一定是1，而是syncmask的值)

或许这些指令将可以为我们提供很多有趣的实现及更复杂的同步算法，比如：

while( __syncthreads_and( pred ) ){

change pred

}

while( __syncthreads_and( __syncthreads_count( pred )==blockDim.x ) ){

change pred

}