CUDA学习（十三）

隐式同步：
如果主机线程在它们之间发出以下任一操作，则来自不同流的两个命令不能同时运行：

• 一个页面锁定的主机内存分配；
• 设备内存分配；
• 一个设备内存集；
• 两个地址之间的内存拷贝到相同的设备内存；
• 任何CUDA命令到NULL流，
• 在计算能力3.x和计算能力7.x中描述的L1 /共享存储器配置之间的切换

对于支持并发内核执行且计算能力为3.0或更低的设备，任何需要依赖性检查以查看流式内核启动是否完成的操作：

• 只有在CUDA上下文中任何流的所有先前内核启动的所有线程块都开始执行时才能开始执行;
• 阻止来自CUDA上下文中任何流的所有稍后的内核启动，直到检查完内核启动为止。
  需要依赖性检查的操作包括与正在检查的启动相同的流中的任何其他命令以及对该流的任何对cudaStreamQuery（）的调用。 因此，应用程序应该遵循这些指导方针来提高其并发内核执行的潜力：
• 所有的独立行动都应该在依赖行动之前发布，
  -任何类型的同步应尽可能延迟。

重叠行为：
两个流之间的执行重叠量取决于向每个流发出命令的顺序，以及设备是否支持数据传输和内核执行，并行内核执行或并发数据传输
例如，在不支持并发数据传输的设备上，创建和销毁代码示例的两个流根本不重叠，因为从设备的存储器复制之后从主机到设备的存储器复制被发布到流[1] 到主机被发送到流[0]，所以一旦从设备到主机的存储器拷贝发出到流[0]，它就只能启动完成。 如果代码按以下方式重写（并且假设设备支持数据传输和内核执行的重叠）

for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
        size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
    MyKernel << <100, 512, 0, stream[i] >> >
    (outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
        size, cudaMemcpyDeviceToHost, stream[i]);

在支持并发数据传输的设备上，两个创建和销毁代码示例流重叠：发送到流[1]的从主机到设备的存储器复制与从设备到主机的存储器复制重叠，发送到流[0 ]，甚至在内核启动时发送到流[0]（假设设备支持数据传输和内核执行的重叠）。然而，对于计算能力为3.0或更低的设备，内核执行不可能重叠，因为在从设备到主机的内存复制发送到流[0]之后，第二次内核启动发送到流[1]，所以直到发送到流[0]的第一次内核启动按照隐式同步完成。如果代码如上所述重写，则内核执行重叠（假定设备支持并发内核执行），因为在从设备到主机的内存复制发送到流[0]之前，第二次内核启动发布到流[1]。然而，在这种情况下，发送到流[0]的从设备到主机的存储器副本仅与根据隐式同步发布到流[1]的内核启动的最后线程块重叠，其可以仅表示总数的一小部分内核的执行时间。