同步功能:void __syncthreads()
一直等待,直到线程块中的所有线程都已达到此点为止,并且__syncthreads()
之前的这些线程所做的全局和共享内存访问对块中的所有线程均可见。__syncthreads()
用于协调同一个块的线程之间的通信。 当一个块中的某些线程访问共享或全局内存中的相同地址时,对于这些内存访问中的一些存在潜在的read-after-write, write-after-read, or write-after-write危险。 通过同步这些访问之间的线程可以避免这些数据危害。__syncthreads()
在条件代码中是允许的,但前提条件是整个线程块的条件相同,否则代码执行可能会挂起或产生意想不到的副作用。
计算能力2.x及更高版本的设备支持下面描述的__syncthreads()
的三种变体:int __syncthreads_count(int predicate);
与__syncthreads()
具有相同的附加功能,它可以评估块的所有线程的形参并返回形参求值为非零的线程数。int __syncthreads_and(int predicate);
与__syncthreads()
具有相同的附加特性,它可以评估块的所有线程的形参并返回非零当且仅当形参对所有线程计算为非零值时才返回非零值。int __syncthreads_or(int predicate);
与__syncthreads()
相同,额外的功能是它为块的所有线程评估x形参并返回非零当且仅当形参对其中任何线程计算为非零值时才返回非零值。void __syncwarp(unsigned mask=0xffffffff);
将导致正在执行的线程等待,直到在掩码中命名的所有warp通道在恢复执行之前执行__syncwarp()
(具有相同的掩码)。 在掩码中命名的所有未指定的线程都必须使用相同的掩码执行相应的__syncwarp()
,否则结果未定义。
执行__syncwarp()
可确保参与障碍的线程之间的内存排序。 因此,希望通过内存进行通信的warp中的线程可以存储到内存中,执行__syncwarp()
,然后安全地读取其他线程存储在warp中的值。
对于.target sm_6x或更低版本,掩码中的所有线程必须在收敛中执行相同的__syncwarp()
,且掩码中所有值的并集必须与活动掩码相等。 否则,行为是不确定的。
数学函数:
参考手册列出了设备代码中支持的所有C / C ++标准库数学函数以及仅在设备代码中支持的所有固有函数。
数学函数在相关时为这些函数提供准确性信息。
纹理对象API:
tex1Dfetch():
template<class T>
T tex1Dfetch(cudaTextureObject_t texObj, int x);
使用整数纹理坐标x从由一维纹理对象texObj指定的线性存储器区域取回。 tex1Dfetch()仅适用于非标准化坐标,因此只支持边界和钳位寻址模式。 它不执行任何纹理过滤。 对于整数类型,它可以选择将整数提升为单精度浮点。