CUDA学习之六（设备运行时组件）

设备运行时组件仅可用于设备函数。

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() 是一个例外情况，它支持单精度浮点数字）。