语言整合:
使用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。
致敬冰蛙