流:
设备运行时可以使用有名和无名(NULL)流。 命名流可以被线程块内的任何线程使用,但是流句柄可能不会被传递给其他块或子/父内核。 换句话说,一个流应该被视为对它创建的块是私有的。 流句柄不保证在块之间是唯一的,因此在未分配块的块中使用流句柄将导致未定义的行为。
与主机端启动类似,启动到不同流中的工作可能会同时运行,但并不保证实际并发性。 CUDA编程模型不支持依赖子内核之间的并发性的程序,并且具有未定义的行为。
设备不支持主机端NULL流的跨流障碍语义(详情见下文)。 为了保持与主机运行时的语义兼容性,必须使用cudaStreamCreateWithFlags()API创建所有设备流,并传递cudaStreamNonBlocking标志。 cudaStreamCreate()调用是仅用于主机运行时的API,并且无法为设备编译。
由于设备运行时不支持cudaStreamSynchronize()和cudaStreamQuery(),因此当应用程序需要知道流启动的子内核已完成时,应该使用cudaDeviceSynchronize()。
隐式(NULL)流:
在主机程序中,未命名(NULL)流与其他流具有附加的屏障同步语义(详情请参阅默认流)。 设备运行时提供了一个隐含的,未命名的流在一个块中的所有线程之间共享,但是由于所有命名流必须使用cudaStreamNonBlocking标志创建,所以启动到NULL流中的工作不会在任何其他流中插入对未决工作的隐式依赖。
事件:
只支持CUDA事件的流间同步功能。 这意味着支持cudaStreamWaitEvent(),但cudaEventSynchronize(),cudaEventElapsedTime()和cudaEventQuery()不支持。 由于不支持cudaEventElapsedTime(),因此必须通过cudaEventCreateWithFlags()创建cudaEvents,并传递cudaEventDisableTiming标志。
对于所有设备运行时对象,事件对象可以在所有线程之间共享,而线程块则创建它们,但是对于该块是本地的,并且可能不会传递给其他内核或同一内核中的块之间。 不保证事件句柄在块之间是唯一的,因此在未创建块的块中使用事件句柄将导致未定义的行为。
同步:
cudaDeviceSynchronize()函数将同步线程块中任何线程启动的所有工作,直到调用cudaDeviceSynchronize()为止。 请注意,cudaDeviceSynchronize()可以从发散的代码中调用(请参阅阻止宽同步)。
如果调用线程用于与从其他线程调用的子网格同步,则由程序执行足够的附加线程间同步,例如通过调用__syncthreads()
。
块宽同步:
cudaDeviceSynchronize()函数并不意味着块内同步。 特别是,如果没有通过__syncthreads()
指令进行显式同步,则调用线程不会对任何其他线程发起的工作做出任何假设。 例如,如果一个块中的多个线程每次都启动工作,并且同时希望所有这些工作同步(可能是因为基于事件的依赖性),则由程序决定,在调用cudaDeviceSynchronize之前,所有线程都会提交此工作()。
因为允许实现从块中任何线程的启动进行同步,所以很有可能多个线程对cudaDeviceSynchronize()的同时调用将耗尽第一个调用中的所有工作,然后对以后的调用没有影响。
设备管理:
只有运行内核的设备才能从该内核控制。 这意味着设备运行时不支持设备API,如cudaSetDevice()。 从GPU看到的活动设备(从cudaGetDevice()返回)将具有与主机系统相同的设备编号。 cudaDeviceGetAttribute()调用可以请求关于另一个设备的信息,因为该API允许将设备ID指定为该调用的参数。 请注意,catch-all cudaGetDeviceProperties()API不是由设备运行时提供的 - 必须单独查询属性。