1. 概述hip::Stream是 HIP 运行时ROCclr/hipamd 层中**命令流stream/queue**的核心实现类。用户通过公开句柄hipStream_t操作它而运行时内部将该句柄reinterpret_cast成hip::Stream*来访问真正的对象。一个 stream 表示一条按 FIFO 顺序执行的 GPU 命令队列同一 stream 内的 kernel 启动、内存拷贝、memset、event 等命令按提交顺序串行执行不同 stream 之间可并发。定义位置声明projects/clr/hipamd/src/hip_internal.hpp实现projects/clr/hipamd/src/hip_stream.cpp2. 类层次与底层映射classStream:publicamd::HostQueue{...};hip::Stream继承自amd::HostQueue这是关键的架构事实因此stream 在底层就是一个 ROCclr 命令队列。在hip_memory.cpp等处看到的command-enqueue()、command-queue()-finishCommand()、stream.finish()等调用本质上都是在操作这条HostQueue。队列再向下对接 HSA/ROCr由 GPU 的硬件队列AQL queue、SDMA 拷贝引擎或 compute shader 执行实际工作。3. 构造与生命周期3.1 构造函数Stream::Stream(hip::Device*dev,Priority p,unsignedintf,boolnull_stream,conststd::vectoruint32_tcuMask,hipStreamCaptureStatus captureStatus):amd::HostQueue(*dev-asContext(),*dev-devices()[0],0,amd::CommandQueue::RealTimeDisabled,convertToQueuePriority(p),cuMask,null_stream),device_(dev),priority_(p),flags_(f),null_(null_stream),cuMask_(cuMask),stream_id_(GenerateStreamId()),captureStatus_(captureStatus){device_-AddStream(this);}构造时初始化基类amd::HostQueue绑定到设备的amd::Context和amd::Device并传入优先级与 CU mask。生成进程内唯一的stream_id_单调递增。调用device_-AddStream(this)把自己注册到所属hip::Device。3.2 创建、销毁、终止方法作用Create()调用基类create()真正分配底层队列资源Destroy()从活动队列与设备注销设置强制销毁标志后release()terminate()调用HostQueue::terminate()清理注意析构函数~Stream() default被声明为private因为HostQueue是引用计数对象生命周期由retain()/release()管理不能直接delete。创建流程公开 API 入口成功失败hipStreamCreate / WithFlags / WithPriorityihipStreamCreatenew hip::Stream(...)hStream-Create()*stream reinterpret_cast(hStream)Stream::Destroy - hipErrorOutOfMemory4. 关键属性成员含义device_所属hip::Devicepriority_调度优先级High(-1)/Normal(0)/Low(1)flags_创建标志如hipStreamNonBlockingnull_是否为 null默认遗留流cuMask_CU 掩码限制可用的计算单元stream_id_进程内唯一流 IDcaptureStatus_等图捕获graph capture相关状态detached_green context 生命周期标志4.1 优先级Priority枚举定义为High -1, Normal 0, Low 1通过convertToQueuePriority()映射到amd::CommandQueue::Priority。优先级范围由hipDeviceGetStreamPriorityRange暴露。5. 特殊 stream 语义5.1 Null stream默认流传nullptr或hipStreamLegacy时运行时使用设备的 null stream。它与其他 stream 有隐式同步语义。isValid对 null/legacy 直接返回 true。5.2 Per-thread default streamhipStreamPerThread通过StreamPerThread为每线程、每设备惰性创建一个独立默认流hipStream_tStreamPerThread::Get(){constintcurrDevhip::getCurrentDevice()-deviceId();if(streams_.empty())streams_.resize(g_devices.size());if(streams_[currDev]nullptr||!hip::isValid(streams_[currDev])){ihipStreamCreate(streams_[currDev],hipStreamDefault,Priority::Normal);}returnstreams_[currDev];}它对hipDeviceReset做了防御性重建。解析入口为getStreamPerThread。5.3 句柄解析流程nullptr / hipStreamLegacyhipStreamPerThread普通句柄hipStream_t 句柄值?设备 null streamgetStreamPerThread - 线程局部流reinterpret_casthip::Stream*6. 执行与同步因为hip::Stream就是HostQueue命令执行依托基类接口提交command-enqueue()把命令放入队列尾部。同步整条流stream.finish()阻塞 host 直到队列清空见hipStreamSynchronize_common。同步单条命令command-queue()-finishCommand(command)。依赖串接getLastQueuedCommand()amd::Marker在跨队列时建立依赖见ihipMemcpy尾部逻辑。hipStreamSynchronize的处理hipError_thipStreamSynchronize_common(hipStream_t stream){getStreamPerThread(stream);if(streamnullptr){// null 流只同步阻塞流getCurrentDevice()-SyncAllStreams(false,/*blockingOnly*/true);returnhipSuccess;}CHECK_STREAM_DETACHED(stream);if(stream!hipStreamLegacyhip::Stream::StreamCaptureOngoing(stream))HIP_RETURN(hipErrorStreamCaptureUnsupported);// 捕获中不允许同步autohip_streamhip::getStream(stream,/*wait*/false);hip_stream-finish();hip_stream-GetDevice()-ReleaseFreedMemory();returnhipSuccess;}7. 流捕获Stream Capture / HIP Graphhip::Stream承担了把一串命令录制成hip::Graph的职责。相关状态机状态含义hipStreamCaptureStatusNone未捕获hipStreamCaptureStatusActive正在捕获hipStreamCaptureStatusInvalidated捕获被破坏非法操作触发关键机制SetCaptureGraph()开始捕获状态置为 Active。lastCapturedNodes_/removedDependencies_维护图节点依赖关系。parallelCaptureStreams_记录 fork 出的并行分支EndCapture()会递归结束它们。captureEvents_把 event 绑定到当前捕获。StreamCaptureOngoing实现了三种捕获模式下的跨流失效策略Relaxed不干扰其他流。Global使所有线程中正在捕获的流失效。ThreadLocal仅使当前线程的捕获流失效。正因为捕获期间不允许隐式同步像hipMemcpyWithStream这类“异步同步”的 API 会先检查StreamCaptureOngoing并在捕获中返回hipErrorStreamCaptureUnsupported。8. 执行上下文Green Context与 Detach为支持 green context 资源隔离stream 引入了detached_原子标志voidStream::Detach(){if(captureStatus_hipStreamCaptureStatusActive){captureStatus_hipStreamCaptureStatusInvalidated;for(autos:parallelCaptureStreams_)reinterpret_casthip::Stream*(s)-SetCaptureStatus(hipStreamCaptureStatusInvalidated);}detached_.store(true,std::memory_order_release);}当拥有该流的ExecutionCtx被销毁时调用Detach()正在进行的捕获被置为 invalidated含所有并行分支。之后对该流的提交/同步 API 须返回hipErrorStreamDetached由CHECK_STREAM_DETACHED宏强制。hipStreamDestroy对已 detach 的流仍可成功。9. 有效性与设备注册每个hip::Device维护其 stream 集合通过AddStream/RemoveStream/StreamExists管理。hip::isValid(stream)遍历所有设备确认句柄仍被某设备持有用于 API 入口的健壮性校验。10. 典型调用链示例以内存拷贝为例HSA/ROCr SDMA/CUamd::*MemoryCommandhip::Stream (HostQueue)hipMemcpyWithStream用户代码HSA/ROCr SDMA/CUamd::*MemoryCommandhip::Stream (HostQueue)hipMemcpyWithStream用户代码传入 hipStream_tisValid / StreamCaptureOngoing 校验hip::getStream(stream) 解析ihipMemcpy 构造拷贝命令command-enqueue()提交到底层硬件队列finishCommand (同步型) 等待完成数据搬运完成11. 小结维度实现要点本质继承amd::HostQueue的 FIFO 命令队列句柄hipStream_t与hip::Stream*直接 reinterpret_cast执行enqueue提交finish/finishCommand同步底层对接 HSA/ROCr特殊流null stream隐式同步、per-thread default线程设备惰性创建调度优先级High/Normal/Low CU mask高级特性Stream captureHIP Graph 录制、green context detach生命周期引用计数管理private 析构设备注册表跟踪hip::Stream是连接 HIP 公开 API 与 ROCclr 命令执行引擎的枢纽对上提供 CUDA 兼容的 stream 语义对下将所有 GPU 操作组织为有序、可并发、可录制的命令队列。 
长河君 · 2026/10/24 4:12:13 · 阅读 1.2k HIP hip::Stream 对象：功能与底层实现
1. 概述hip::Stream是 HIP 运行时ROCclr/hipamd 层中**命令流stream/queue**的核心实现类。用户通过公开句柄hipStream_t操作它而运行时内部将该句柄reinterpret_cast成hip::Stream*来访问真正的对象。一个 stream 表示一条按 FIFO 顺序执行的 GPU 命令队列同一 stream 内的 kernel 启动、内存拷贝、memset、event 等命令按提交顺序串行执行不同 stream 之间可并发。定义位置声明projects/clr/hipamd/src/hip_internal.hpp实现projects/clr/hipamd/src/hip_stream.cpp2. 类层次与底层映射classStream:publicamd::HostQueue{...};hip::Stream继承自amd::HostQueue这是关键的架构事实因此stream 在底层就是一个 ROCclr 命令队列。在hip_memory.cpp等处看到的command-enqueue()、command-queue()-finishCommand()、stream.finish()等调用本质上都是在操作这条HostQueue。队列再向下对接 HSA/ROCr由 GPU 的硬件队列AQL queue、SDMA 拷贝引擎或 compute shader 执行实际工作。3. 构造与生命周期3.1 构造函数Stream::Stream(hip::Device*dev,Priority p,unsignedintf,boolnull_stream,conststd::vectoruint32_tcuMask,hipStreamCaptureStatus captureStatus):amd::HostQueue(*dev-asContext(),*dev-devices()[0],0,amd::CommandQueue::RealTimeDisabled,convertToQueuePriority(p),cuMask,null_stream),device_(dev),priority_(p),flags_(f),null_(null_stream),cuMask_(cuMask),stream_id_(GenerateStreamId()),captureStatus_(captureStatus){device_-AddStream(this);}构造时初始化基类amd::HostQueue绑定到设备的amd::Context和amd::Device并传入优先级与 CU mask。生成进程内唯一的stream_id_单调递增。调用device_-AddStream(this)把自己注册到所属hip::Device。3.2 创建、销毁、终止方法作用Create()调用基类create()真正分配底层队列资源Destroy()从活动队列与设备注销设置强制销毁标志后release()terminate()调用HostQueue::terminate()清理注意析构函数~Stream() default被声明为private因为HostQueue是引用计数对象生命周期由retain()/release()管理不能直接delete。创建流程公开 API 入口成功失败hipStreamCreate / WithFlags / WithPriorityihipStreamCreatenew hip::Stream(...)hStream-Create()*stream reinterpret_cast(hStream)Stream::Destroy - hipErrorOutOfMemory4. 关键属性成员含义device_所属hip::Devicepriority_调度优先级High(-1)/Normal(0)/Low(1)flags_创建标志如hipStreamNonBlockingnull_是否为 null默认遗留流cuMask_CU 掩码限制可用的计算单元stream_id_进程内唯一流 IDcaptureStatus_等图捕获graph capture相关状态detached_green context 生命周期标志4.1 优先级Priority枚举定义为High -1, Normal 0, Low 1通过convertToQueuePriority()映射到amd::CommandQueue::Priority。优先级范围由hipDeviceGetStreamPriorityRange暴露。5. 特殊 stream 语义5.1 Null stream默认流传nullptr或hipStreamLegacy时运行时使用设备的 null stream。它与其他 stream 有隐式同步语义。isValid对 null/legacy 直接返回 true。5.2 Per-thread default streamhipStreamPerThread通过StreamPerThread为每线程、每设备惰性创建一个独立默认流hipStream_tStreamPerThread::Get(){constintcurrDevhip::getCurrentDevice()-deviceId();if(streams_.empty())streams_.resize(g_devices.size());if(streams_[currDev]nullptr||!hip::isValid(streams_[currDev])){ihipStreamCreate(streams_[currDev],hipStreamDefault,Priority::Normal);}returnstreams_[currDev];}它对hipDeviceReset做了防御性重建。解析入口为getStreamPerThread。5.3 句柄解析流程nullptr / hipStreamLegacyhipStreamPerThread普通句柄hipStream_t 句柄值?设备 null streamgetStreamPerThread - 线程局部流reinterpret_casthip::Stream*6. 执行与同步因为hip::Stream就是HostQueue命令执行依托基类接口提交command-enqueue()把命令放入队列尾部。同步整条流stream.finish()阻塞 host 直到队列清空见hipStreamSynchronize_common。同步单条命令command-queue()-finishCommand(command)。依赖串接getLastQueuedCommand()amd::Marker在跨队列时建立依赖见ihipMemcpy尾部逻辑。hipStreamSynchronize的处理hipError_thipStreamSynchronize_common(hipStream_t stream){getStreamPerThread(stream);if(streamnullptr){// null 流只同步阻塞流getCurrentDevice()-SyncAllStreams(false,/*blockingOnly*/true);returnhipSuccess;}CHECK_STREAM_DETACHED(stream);if(stream!hipStreamLegacyhip::Stream::StreamCaptureOngoing(stream))HIP_RETURN(hipErrorStreamCaptureUnsupported);// 捕获中不允许同步autohip_streamhip::getStream(stream,/*wait*/false);hip_stream-finish();hip_stream-GetDevice()-ReleaseFreedMemory();returnhipSuccess;}7. 流捕获Stream Capture / HIP Graphhip::Stream承担了把一串命令录制成hip::Graph的职责。相关状态机状态含义hipStreamCaptureStatusNone未捕获hipStreamCaptureStatusActive正在捕获hipStreamCaptureStatusInvalidated捕获被破坏非法操作触发关键机制SetCaptureGraph()开始捕获状态置为 Active。lastCapturedNodes_/removedDependencies_维护图节点依赖关系。parallelCaptureStreams_记录 fork 出的并行分支EndCapture()会递归结束它们。captureEvents_把 event 绑定到当前捕获。StreamCaptureOngoing实现了三种捕获模式下的跨流失效策略Relaxed不干扰其他流。Global使所有线程中正在捕获的流失效。ThreadLocal仅使当前线程的捕获流失效。正因为捕获期间不允许隐式同步像hipMemcpyWithStream这类“异步同步”的 API 会先检查StreamCaptureOngoing并在捕获中返回hipErrorStreamCaptureUnsupported。8. 执行上下文Green Context与 Detach为支持 green context 资源隔离stream 引入了detached_原子标志voidStream::Detach(){if(captureStatus_hipStreamCaptureStatusActive){captureStatus_hipStreamCaptureStatusInvalidated;for(autos:parallelCaptureStreams_)reinterpret_casthip::Stream*(s)-SetCaptureStatus(hipStreamCaptureStatusInvalidated);}detached_.store(true,std::memory_order_release);}当拥有该流的ExecutionCtx被销毁时调用Detach()正在进行的捕获被置为 invalidated含所有并行分支。之后对该流的提交/同步 API 须返回hipErrorStreamDetached由CHECK_STREAM_DETACHED宏强制。hipStreamDestroy对已 detach 的流仍可成功。9. 有效性与设备注册每个hip::Device维护其 stream 集合通过AddStream/RemoveStream/StreamExists管理。hip::isValid(stream)遍历所有设备确认句柄仍被某设备持有用于 API 入口的健壮性校验。10. 典型调用链示例以内存拷贝为例HSA/ROCr SDMA/CUamd::*MemoryCommandhip::Stream (HostQueue)hipMemcpyWithStream用户代码HSA/ROCr SDMA/CUamd::*MemoryCommandhip::Stream (HostQueue)hipMemcpyWithStream用户代码传入 hipStream_tisValid / StreamCaptureOngoing 校验hip::getStream(stream) 解析ihipMemcpy 构造拷贝命令command-enqueue()提交到底层硬件队列finishCommand (同步型) 等待完成数据搬运完成11. 小结维度实现要点本质继承amd::HostQueue的 FIFO 命令队列句柄hipStream_t与hip::Stream*直接 reinterpret_cast执行enqueue提交finish/finishCommand同步底层对接 HSA/ROCr特殊流null stream隐式同步、per-thread default线程设备惰性创建调度优先级High/Normal/Low CU mask高级特性Stream captureHIP Graph 录制、green context detach生命周期引用计数管理private 析构设备注册表跟踪hip::Stream是连接 HIP 公开 API 与 ROCclr 命令执行引擎的枢纽对上提供 CUDA 兼容的 stream 语义对下将所有 GPU 操作组织为有序、可并发、可录制的命令队列。