CUDA学习（四十三）

只读数据缓存加载函数：
只读数据缓存加载功能仅受计算能力3.5和更高的设备支持：

T __ldg(const T* address);

返回位于地址地址处的类型T的数据，其中T为char，short，int，long long unsigned char，unsigned short，unsigned int，unsigned long long，int2，int4，uint2，uint4，float，float2，float4，double ，或double2。 该操作缓存在只读数据缓存中（请参阅全局内存）。
时间函数：

clock_t clock();
long long int clock64();

当在设备代码中执行时，返回每个时钟周期递增的每多处理器计数器的值。 在内核的开始和结束时对该计数器进行采样，取出两个采样的差值，并记录每个线程的结果，为每个线程提供一个度量，以便设备完成执行线程所花费的时钟周期数， 但不是设备实际执行线程指令所用的时钟周期数。 由于线程是时间分片的，因此前者的数字大于后者。
原子操作函数：
原子函数对驻留在全局或共享内存中的一个32位或64位字进行读 - 修改 - 写原子操作。 例如，atomicAdd（）在全局或共享内存中的某个地址处读取一个字，向其中添加一个数字，并将结果写回到相同的地址。 该操作是原子的，因为它保证在没有其他线程干扰的情况下执行。 换句话说，除非操作完成，否则其他线程无法访问此地址。 原子函数不会作为内存隔离，也不意味着内存操作的同步或排序约束（有关内存隔离的更多详细信息，请参阅内存隔离函数）。 原子功能只能用于设备功能。
在计算能力低于6.x的GPU架构上，从GPU完成的原子操作仅针对该GPU而言是原子的。 如果GPU尝试对对等GPU的内存进行原子操作，则该操作将显示为常规读取，然后写入对等GPU，并且这两个操作不会作为单个原子操作完成。 同样，从CPU启动的原子操作，从GPU到CPU存储器的原子操作不会是原子操作。
计算能力6.x引入了新的原子类型，允许开发者扩大或缩小原子操作的范围。 例如，atomicAdd_system保证该指令对于系统中的其他CPU和GPU是原子的。 atomicAdd_block意味着该指令只针对来自同一线程块中其他线程的原子的原子。 在以下示例中，CPU和GPU都可以自动更新地址addr处的整数值：

__global__ void mykernel(int *addr) {
    atomicAdd_system(addr, 10); // only available on devices with compute
    capability 6.x
}
void foo() {
    int *addr;
    cudaMallocManaged(&addr, 4);
    *addr = 0;
    mykernel << <... >> >(addr);
    __sync_fetch_and_add(addr, 10); // CPU atomic operation
}

原子的新范围版本可用于下面列出的所有原子，仅用于计算能力6.x及更高版本。
请注意，任何原子操作都可以基于atomicCAS（）（Compare和Swap）来实现。 例如，双精度浮点数的atomicAdd（）在计算能力低于6.0的设备上不可用，但可以按如下方式实现：

#if __CUDA_ARCH__ < 600
__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
        (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed,
            __double_as_longlong(val +
                __longlong_as_double(assumed)));
        // Note: uses integer comparison to avoid hang in case of NaN (since NaN !=
        NaN)
    } while (assumed != old);
        return __longlong_as_double(old);
}
#endif