CUDA学习（四十五）

按位函数：
atomicAnd()：

int atomicAnd(int* address, int val);
unsigned int atomicAnd(unsigned int* address,
    unsigned int val);
unsigned long long int atomicAnd(unsigned long long int* address,
    unsigned long long int val);

读取位于全局或共享内存中地址地址的32位或64位字，计算（old和val），并将结果存回同一地址的存储器中。 这三个操作在一个原子事务中执行。 该函数返回旧的。
atomicAnd（）的64位版本仅受计算能力3.5及更高版本的设备支持。
atomicOr()：

int atomicOr(int* address, int val);
unsigned int atomicOr(unsigned int* address,
    unsigned int val);
unsigned long long int atomicOr(unsigned long long int* address,
    unsigned long long int val);

读取位于全局或共享内存中地址地址处的32位或64位字，计算（old | val），并将结果存回相同地址的存储器中。 这三个操作在一个原子事务中执行。 该函数返回旧的。
atomicOr（）的64位版本仅受计算能力3.5及更高版本的设备支持。
atomicXor()：

int atomicXor(int* address, int val);
unsigned int atomicXor(unsigned int* address,
    unsigned int val);
unsigned long long int atomicXor(unsigned long long int* address,
    unsigned long long int val);

读取位于全局或共享内存中地址地址的32位或64位字，计算（old ^ val），并将结果存回同一地址的存储器中。 这三个操作在一个原子事务中执行。 该函数返回旧的。
atomicXor（）的64位版本仅受计算能力3.5及更高版本的设备支持。
Warp投票函数：

int __all_sync(unsigned mask, int predicate);
int __any_sync(unsigned mask, int predicate);
unsigned __ballot_sync(unsigned mask, int predicate);
unsigned __activemask();

弃用声明：从CUDA 9.0开始，__any，__all和__ballot已被弃用。
warp投票功能允许给定warp的线程执行简化和广播操作。 这些函数将warp中每个线程的整数谓词作为输入，并将这些值与零进行比较。 通过以下方式之一将比较结果在变形的活动线程中进行组合（缩小），向每个参与线程广播单个返回值：
__all_sync(unsigned mask, predicate):
评估所有非退出线程在掩码中的形参并返回非零当且仅当谓词对所有线程计算为非零值时才返回非零值。
__any_sync(unsigned mask, predicate):
在掩码中评估所有非退出线程的形参并返回非零当且仅当形参对其中的任何线程计算为非零。
__ballot_sync(unsigned mask, predicate):
在掩码中评估所有非退出线程的谓词并返回一个整数，其第N位被设置当且仅当谓词对于第N个线程的变形计算为非零并且第N个线程处于活动状态时。
__activemask():
返回调用warp中所有当前活动线程的32位整数掩码。 如果在调用__activemask（）时变形中的第N条车道处于活动状态，则设置第N位。 非活动线程在返回的掩码中由0位表示。 退出程序的线程始终标记为非活动状态。 请注意，收敛于__activemask（）调用的线程不能保证在后续指令处收敛，除非这些指令正在同步warp-builtin函数。
注意：
对于__all_sync，__any_sync和__ballot_sync，必须传递一个指定参与调用的线程的掩码。 代表线程的通道ID的位必须为每个参与线程设置，以确保在硬件执行内部函数之前它们已正确收敛。 在掩码中命名的所有活动线程都必须使用相同的掩码执行相同的内部属性，否则结果是未定义的。