CUDA学习(九十九)

高级主题:模块化程序和数据访问约束
在前面的示例中,cudaMallocManaged()指定cudaMemAttachHost标志,该标志创建对设备端执行最初不可见的分配。 (默认分配对于所有流上的所有GPU内核都是可见的。)这确保了在数据分配和针对特定流获取数据的时间间隔内不会与另一个线程执行的意外交互。
如果没有这个标志,如果另一个线程发起的内核碰巧正在运行,那么新的分配将被视为在GPU上使用。 这可能会影响线程在能够将其明确附加到私有流之前从CPU访问新分配的数据的能力(例如,在基类构造函数中)。 为了实现线程之间的安全独立,因此应该使用分配来指定这个标志。
另一种方法是在分配连接到流之后,在所有线程之间放置一个流程范围的屏障。 这将确保所有线程在启动任何内核之前完成其数据/流关联,从而避免危险。 在流被销毁之前需要第二个屏障,因为流破坏会导致分配恢复到其默认可见性。 cudaMemAttachHost标志既可以简化这个过程,也可以在需要时插入全局屏障。
带管理内存的Memcpy()/ Memset()行为:
由于托管内存可以从主机或设备访问,因此cudaMemcpy *()依赖于使用cudaMemcpyKind指定的传输类型来确定数据是作为主机指针还是设备指针访问。
如果指定了cudaMemcpyHostTo 并且管理源数据,则它将从主机访问(如果它可以从复制流中的主机一致访问)(1); 否则将从设备访问。 当指定cudaMemcpy ToHost并且目标是托管内存时,类似的规则适用于目标。
如果指定了cudaMemcpyDeviceTo 并且管理源数据,则将从设备访问它。 源必须可以从复制流中的设备一致访问(2); 否则,返回错误。 当指定cudaMemcpy ToDevice并且目标是托管内存时,类似的规则适用于目标。
如果指定了cudaMemcpyDefault,则将从主机访问托管数据(如果不能从复制流中的设备一致访问),或者数据的首选位置是cudaCpuDeviceId,并且可以从主机一致地访问 在复制流(1)中; 否则,它将从设备访问。
将cudaMemset *()与托管内存一起使用时,始终可以从设备访问数据。 数据必须能够在用于cudaMemset()操作(2)的流中与设备一致地访问; 否则,返回错误。
当通过cudaMemcpy 或cudaMemset 从设备访问数据时,操作流将被视为在GPU上处于活动状态。 在此期间,如果GPU对设备属性concurrentManagedAccess具有零值,则与该流相关联的数据的任何CPU访问或具有全局可见性的数据都将导致分段错误。 程序必须进行适当的同步,以确保在访问来自CPU的任何关联数据之前操作已完成。
(1)托管内存在给定流中可以从主机一致访问,at
必须满足以下至少一个条件:

  • 给定的流与设备属性concurrentManagedAccess具有非零值的设备相关联。
  • 内存既不具有全局可见性,也不与给定流关联。

(2)为了使管理内存在给定流中可以从设备一致访问,必须满足以下至少一个条件:

  • 设备对设备属性concurrentManagedAccess具有非零值。
  • 内存或者具有全局可见性或者与给定流相关联。
    timg
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值