设备运行时组件仅可用于设备函数。
1.数学函数
设备运行时组件中存在准确性略低而速度更快的版本;其名称相同,但带有一个_前缀(如_sinf(x))。编译器有一个 (-use_fast_math) 选项,用于强制要求所有函数编译其准确性略低的版本(如果存在)。
2.同步函数
void __syncthreads();
同步块中的所有线程。一旦所有线程均达到此同步点,执行将正常恢复。
_syncthreads() 用于调整同一个块的线程之间的通信。在一个块内的某些线程访问共享或全局存储器中的相同地址时,部分访问操作可能存在写入后读取、读取后写入或写入后写入之类的风险。可通过在这些访问操作间同步线程来避免这些数据风险。
_syncthreads() 允许在条件代码中使用,但仅当条件估值在整个线程块中都相同时才允许使用,否则代码执行将有可能挂起,或者出现意料之外的副作用。
3.纹理函数
对于来自线性存储器的纹理,通过 tex1Dfetch() 系列函数访问纹理,示例如下:
template<class Type>
Type tex1Dfetch(
texture<Type, 1, cudaReadModeElementType> texRef,
int x);
float tex1Dfetch(
texture<unsigned char, 1, cudaReadModeNormalizedFloat> texRef,
int x);
float tex1Dfetch(
texture<signed char, 1, cudaReadModeNormalizedFloat> texRef,
int x);
float tex1Dfetch(
texture<unsigned short, 1, cudaReadModeNormalizedFloat> texRef,
int x);
float tex1Dfetch(
texture<signed short, 1, cudaReadModeNormalizedFloat> texRef,
int x);
这些函数会使用纹理坐标 x 获取绑定到纹理参考 texRef 的线性存储器区域。不支持纹理过滤和寻址模式。对于整型来说,这些函数可选择将整型转变为单精度浮点类型。
除了上述函数以外,还支持 2 元组和 4 元组,示例如下:
float4 tex1Dfetch(
texture<uchar4, 1, cudaReadModeNormalizedFloat> texRef,
int x);
以上示例将使用纹理坐标 x 获取绑定到纹理参考 texRef 的线性存储器。
来自 CUDA 数组的纹理:
对于来自 CUDA 数组的纹理,可通过 tex1D()、tex2D()、tex3D() 访问纹理:
template<class Type, enum cudaTextureReadMode readMode>
Type tex1D(texture<Type, 1, readMode> texRef,
float x);
template<class Type, enum cudaTextureReadMode readMode>
Type tex2D(texture<Type, 2, readMode> texRef,
float x, float y);
template<class Type, enum cudaTextureReadMode readMode>
Type tex3D(texture<Type, 3, readMode> texRef,
float x, float y, float z);
这些函数将使用纹理坐标 x、y 和 z 获取绑定到纹理参考 texRef 的 CUDA 数组。纹理参考的不变(编译时)和可变(运行时)属性相互结合,共同确定坐标的解释方式、在纹理获取过程中发生的处理以及纹理获取所提供的返回值。
4.原子函数
原子函数对位于全局或共享存储器内的一个 32 位或 64 位字执行读取-修改-写入原子操作。例如,atomicAdd() 将在全局或共享存储器内的某个地址读取 32 位字,将其与一个整型相加,并将结果写回同一地址。之所以说这样的操作是原子的,是因为它可在不干扰其他线程的前提下执行。换句话说,在操作完成中,其他任何线程都无法访问此地址。
原子操作仅适用于有符号和无符号整型(但 atomicExch() 是一个例外情况,它支持单精度浮点数字)。