));
对于用模板定义的线程块片(thread block tile)，还可以使用如下函数（类似线程束内的基本函数）：
int any(int predicate);
int all(int predicate);
unsigned int ballot(int predicate);
T shfl(T var, int srcLane);
T shfl_down(T var, unsigned int delta);
T shfl_up(T var, unsigned int delta);
T shfl_xor(T var, unsigned int laneMask);
与对应的线程束内基本函数相比，上述线程块片的函数主要有两点不同：①、少了代表掩码的参数，因为线程组内所有线程都必须参与函数的计算；②、洗牌函数少了最后一个代表宽度的参数，因为宽度就等于线程块片的大小，即模板参数。
4. 原子函数
原子函数对它第一个参数指向的数据进行一次“读-改-写”的原子操作，第一个参数可以指向全局内存，也可以指向共享内存。原子操作是一个线程一个线程轮流做的，但没有明确的次序，另外，原子函数没有同步功能。所有原子函数都是__device__函数，只能在核函数中使用。
下表列出了所有原子函数的原型，address所指变量的值在执行原子函数前为old，执行原子函数后为new。对于每个原子函数，返回值都是old。
| 运算 |
功能 |
函数原型 |
| 加法 |
new = old + val |
T atomicAdd(T* address, T val); |
| 减法 |
new = old - val |
T atomicSub(T* address, T val); |
| 自增 |
new = (old >= val) ? 0 : (old + 1) |
T atomicInc(T* address, T val); |
| 自减 |
new = ((old == 0) || (old > val)) ? val : (old - 1) |
T atomicDec(T* address, T val); |
| 最小值 |
new = (old < val) : old : val |
T atomicMin(T* address, T val); |
| 最大值 |
new = (old > val) : old : val |
T atomicMax(T* address, T val); |
| 交换 |
new = val |
T atomicExch(T* address, T val); |
| 比较-交换(compare and swap) |
new = (old == compare) ? val : old |
T atomicCAS(T* address, T compare, T val); |
| 按位与 |
new = old & val |
T atomicAnd(T* address, T val); |
| 按位或 |
new = old | val |
T atomicOr(T* address, T val); |
| 按位异或 |
new = old ^ val |
T atomicXor(T* address, T val); |
上面所列的函数中，我们用T表示相关变量的数据类型。各个原子函数对数据类型的支持情况见下表。注：atomicAdd对double和__half2的支持始于Pascal架构，对__half的支持始于Volta架构。
| 原子函数 |
int |
unsigned |
unsigned long long |
float |
double |
__half2 |
__half |
| atomicAdd |
yes |
yes |
yes |
yes |
yes |
yes |
yes |
| atomicSub |
yes |
yes |
no |
no |
no |
no |
no |
| atomicInc |
no |
yes |
no |
no |
no |
no |
no |
| atomicDec |
no |
yes |
no |
no |
no |
no |
no |
| atomicMin |
yes |
yes |
yes |
no |
no |
no |
no |
| atomicMax |
yes |
yes |
yes |
no |
no |
no |
no |
| atomicExch |
yes |
yes |
yes |
yes |
no |
no |
no |
| atomicCAS |
yes |
yes |
yes |
no |
no |
no |
no |
| atomicAnd |
yes |
yes |
yes |
no |
no |
no |
no |
| atomicOr |
yes |
yes |
yes |
no |
no |
no |
no |
| atomicXor |
yes |
yes |
yes |
no |
no |
no |
no |