CUDA学习（九十七）

显式同步和逻辑GPU活动：
请注意，即使内核在上述示例中快速运行并在CPU触及y之前完成，也需要显式同步。 Unified Memory使用逻辑活动来确定GPU是否空闲。 这与CUDA编程模型保持一致，该模型指定内核可以在启动后随时运行，并且不能保证在主机发出同步调用之前完成。
逻辑上保证GPU完成其工作的任何函数调用都是有效的。 这包括cudaDeviceSynchronize（）; cudaStreamSynchronize（）和cudaStreamQuery（）（只要它返回cudaSuccess而不是cudaErrorNotReady）指定的流是仍然在GPU上执行的唯一流; cudaEventSynchronize（）和cudaEventQuery（）在指定事件没有被任何设备工作跟踪的情况下; 以及被记录为与主机完全同步的cudaMemcpy（）和cudaMemset（）的用法。
在流之间创建的依赖关系将通过同步流或事件来推断其他流的完成。 可以通过cudaStreamWaitEvent（）创建依赖关系，或者在使用默认（NULL）流时隐式创建依赖关系。
如果没有其他可能正在访问托管数据的流在GPU上处于活动状态，则CPU从流回调中访问托管数据是合法的。 另外，任何设备工作都没有遵循的回调可以用于同步：例如，通过在回调中发信号通知条件变量; 否则，CPU访问仅在回调期间有效。
有几点值得注意：

• 在GPU处于活动状态时，CPU始终允许访问非管理的零拷贝数据。
• GPU在运行任何内核时都被认为是活动的，即使该内核不使用托管数据。 如果内核可能使用数据，则禁止访问，除非设备属性concurrentManagedAccess为1。
• 除了那些适用于非管理内存的多GPU访问的内存之外，托管内存的并发GPU间访问没有限制。
• 并发GPU内核访问托管数据没有限制。

注意最后一点如何允许GPU内核之间的竞争（race），就像目前非管理GPU内存的情况一样。 如前所述，从GPU的角度来看，托管内存的功能与非托管内存相同。 以下代码示例说明了这些要点：

int main() {
    cudaStream_t stream1, stream2;
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);
    int *non_managed, *managed, *also_managed;
    cudaMallocHost(&non_managed, 4); // Non-managed, CPU-accessible memory
    cudaMallocManaged(&managed, 4);
    cudaMallocManaged(&also_managed, 4);
    // Point 1: CPU can access non-managed data.
    kernel << < 1, 1, 0, stream1 >> >(managed);
    *non_managed = 1;
    // Point 2: CPU cannot access any managed data while GPU is busy,
    // unless concurrentManagedAccess = 1
    // Note we have not yet synchronized, so "kernel" is still active.
    *also_managed = 2; // Will issue segmentation fault
                       // Point 3: Concurrent GPU kernels can access the same data.
    kernel << < 1, 1, 0, stream2 >> >(managed);
    // Point 4: Multi-GPU concurrent access is also permitted.
    cudaSetDevice(1);
    kernel << < 1, 1 >> >(managed);
    return 0;
}

使用流管理数据可视性和并行CPU + GPU访问：
到目前为止，假设对于6.x之前的SM体系结构：1）任何活动内核都可以使用任何托管内存，以及2）在内核处于活动状态时使用来自CPU的托管内存无效。 在这里，我们提出了一个更好的粒度控制管理内存的系统，该系统设计用于支持托管内存的所有设备，包括concurrentManagedAccess等于0的早期体系结构。
CUDA编程模型提供流作为程序机制来指示内核启动之间的依赖性和独立性。 启动到同一个流中的内核保证连续执行，而启动到不同流中的内核可以同时执行。 流描述了工作项目之间的独立性，因此可以通过并发来提高效率。
统一内存通过允许CUDA程序将托管分配与CUDA流明确关联，建立在流独立模型上。 通过这种方式，程序员根据内核是否被启动到指定的流中来指示内核使用数据。 这使基于特定程序数据访问模式的并发机会成为可能。 控制这种行为的功能是：

cudaError_t cudaStreamAttachMemAsync(cudaStream_t stream,
        void *ptr,
        size_t length = 0,
        unsigned int flags = 0);

cudaStreamAttachMemAsync（）函数将从ptr开始的内存的长度字节与指定的流相关联。 （目前，长度必须始终为0，表示应该连接整个区域）。由于这种关联，只要流中的所有操作都已完成，统一内存系统就允许CPU访问此内存区域，无论其他流 活跃。 实际上，这限制了活动GPU对托管内存区域的独占所有权，以限制每个流的活动而不是整个GPU的活动。
最重要的是，如果分配不与特定的流相关联，则对所有正在运行的内核都可见，而不管它们的流如何。 这是cudaMallocManaged（）分配或__managed__变量的默认可见性; 因此，在任何内核运行时，CPU可能不会触及数据的简单规则。
通过将分配与特定的流相关联，该程序保证只有启动到该流中的内核才会触及该数据。 统一内存系统不执行错误检查：程序员有责任确保担保得到履行。
除了允许更高的并发性之外，使用cudaStreamAttachMemAsync（）可以（并且通常会）在统一内存系统内启用数据传输优化，这可能会影响延迟和其他开销。