CUDA学习（一百）

语言整合：
使用nvcc编译主机代码的CUDA Runtime API用户可以通过<<< ... >>>运算符访问其他语言集成功能，例如共享符号名称和内联内核启动。 Unified Memory为CUDA的语言集成增加了一个额外元素：用__managed__关键字注释的变量可直接从主机和设备代码中引用。
以下简化GPU编程的例子说明了__managed__全局声明的一个简单使用：

// Managed variable declaration is an extra annotation with __device__
__device__ __managed__ int x;
__global__ void kernel() {
    // Reference "x" directly - it's a normal variable on the GPU.
    printf("GPU sees: x = %d\n", x);
}
int main() {
    // Set "x" from Host code. Note it's just a normal variable on the CPU.
    x = 1234;
    // Launch a kernel which uses "x" from the GPU.
    kernel << < 1, 1 >> >();
    cudaDeviceSynchronize();
    return 0;
}

具有__managed__变量的可用功能是符号在设备代码和主机代码中均可用，无需取消引用指针，数据由所有人共享。 这使得在主机和设备程序之间交换数据变得特别容易，而不需要明确的分配或复制。
在语义上，__managed__变量的行为与通过cudaMallocManaged（）分配的存储的行为相同。 数据托管在物理GPU存储中，并且系统中的所有GPU以及CPU都可以看到数据。 流可见性默认为cudaMemAttachGlobal，但可能会受到cudaStreamAttachMemAsync（）的限制。
有效的CUDA上下文对于__managed__变量的正确操作是必需的。 如果尚未创建当前设备的上下文，则访问__managed__变量可以触发CUDA上下文创建。 在上面的示例中，在内核启动之前访问x会触发设备0上的上下文创建。如果没有该访问，内核启动将触发上下文创建。
声明为__managed__的C ++对象受到某些特定的约束，特别是在涉及静态初始化器的情况下。 请参阅CUDA C编程指南中的C / C ++语言支持以获取这些约束的列表。
__managed__变量发生主机程序错误:
__managed__变量的使用取决于底层的统一内存系统是否正常运行。 例如，如果CUDA安装失败或CUDA上下文创建失败，则可能会发生错误的功能。
当特定于CUDA的操作失败时，通常会返回一个错误，指示失败的来源。 如果统一内存系统运行不正常，则使用__managed__变量会引入一种新的失败模式，即非CUDA操作（例如，CPU访问应该成为有效主机内存地址的内容）可能会失败。 这种无效的内存访问不能轻易归因于底层的CUDA子系统，尽管像cuda-gdb这样的调试器会指出托管内存地址是失败的根源。
查询统一内存支持:
设备属性:
统一内存仅在计算能力为3.0或更高的设备上受支持。 程序可以通过使用cudaGetDeviceProperties（）查询GPU设备是否支持托管内存并检查新的managedMemory属性。 该功能还可以使用具有属性cudaDevAttrManagedMemory的单个属性查询函数cudaDeviceGetAttribute（）来确定。
如果在GPU上和当前操作系统下允许托管内存分配，则任一属性都将设置为1。 请注意，即使GPU具有足够的功能，32位应用程序也不支持统一内存（除非在Android上）。
支持平台上的计算能力6.x的设备可以访问可分页的内存，而无需在其上调用cudaHostRegister。 应用程序可以通过检查新的pageableMemoryAccess属性来查询设备是否支持连贯地访问可分页内存。
使用新的页面错误机制，全局数据一致性通过统一内存得以保证。 这意味着CPU和GPU可以同时访问统一内存分配。 这在计算能力低于6.x的设备上是非法的，因为如果CPU在GPU内核处于活动状态时访问统一内存分配，则无法保证一致性。 程序可以通过检查concurrentManagedAccess属性来查询并发访问支持。
指针属性：
要确定给定的指针是否引用托管内存，程序可以调用cudaPointerGetAttributes（）并检查isManaged属性的值。 如果指针指向托管内存，则该属性设置为1，否则设置为0。

致敬冰蛙