天天看点

CUDA学习(一百)

语言整合:

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

以下简化GPU编程的例子说明了__managed__全局声明的一个简单使用:

具有<code>__managed__</code>变量的可用功能是符号在设备代码和主机代码中均可用,无需取消引用指针,数据由所有人共享。 这使得在主机和设备程序之间交换数据变得特别容易,而不需要明确的分配或复制。

在语义上,<code>__managed__</code>变量的行为与通过cudaMallocManaged()分配的存储的行为相同。 数据托管在物理GPU存储中,并且系统中的所有GPU以及CPU都可以看到数据。 流可见性默认为cudaMemAttachGlobal,但可能会受到cudaStreamAttachMemAsync()的限制。

有效的CUDA上下文对于<code>__managed__</code>变量的正确操作是必需的。 如果尚未创建当前设备的上下文,则访问<code>__managed__</code>变量可以触发CUDA上下文创建。 在上面的示例中,在内核启动之前访问x会触发设备0上的上下文创建。如果没有该访问,内核启动将触发上下文创建。

声明为<code>__managed__</code>的C ++对象受到某些特定的约束,特别是在涉及静态初始化器的情况下。 请参阅CUDA C编程指南中的C / C ++语言支持以获取这些约束的列表。

<code>__managed__</code>变量发生主机程序错误:

<code>__managed__</code>变量的使用取决于底层的统一内存系统是否正常运行。 例如,如果CUDA安装失败或CUDA上下文创建失败,则可能会发生错误的功能。

当特定于CUDA的操作失败时,通常会返回一个错误,指示失败的来源。 如果统一内存系统运行不正常,则使用<code>__managed__</code>变量会引入一种新的失败模式,即非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。

致敬冰蛙

继续阅读