当前位置:网站首页>CUDA用户对象
CUDA用户对象
2022-07-02 06:12:00 【扫地的小何尚】
CUDA用户对象
CUDA 用户对象可用于帮助管理 CUDA 中异步工作所使用的资源的生命周期。 特别是,此功能对于 CUDA 图和流捕获非常有用。
各种资源管理方案与 CUDA 图不兼容。 例如,考虑基于事件的池或同步创建、异步销毁方案。
// Library API with pool allocation
void libraryWork(cudaStream_t stream) {
auto &resource = pool.claimTemporaryResource();
resource.waitOnReadyEventInStream(stream);
launchWork(stream, resource);
resource.recordReadyEvent(stream);
}
// Library API with asynchronous resource deletion
void libraryWork(cudaStream_t stream) {
Resource *resource = new Resource(...);
launchWork(stream, resource);
cudaStreamAddCallback(
stream,
[](cudaStream_t, cudaError_t, void *resource) {
delete static_cast<Resource *>(resource);
},
resource,
0);
// Error handling considerations not shown
}
由于需要间接或图更新的资源的非固定指针或句柄,以及每次提交工作时需要同步 CPU 代码,这些方案对于 CUDA 图来说是困难的。如果这些注意事项对库的调用者隐藏,并且由于在捕获期间使用了不允许的 API,它们也不适用于流捕获。存在各种解决方案,例如将资源暴露给调用者。 CUDA 用户对象提供了另一种方法。
CUDA 用户对象将用户指定的析构函数回调与内部引用计数相关联,类似于 C++ shared_ptr
。引用可能归 CPU 上的用户代码和 CUDA 图所有。请注意,对于用户拥有的引用,与 C++ 智能指针不同,没有代表引用的对象;用户必须手动跟踪用户拥有的引用。一个典型的用例是在创建用户对象后立即将唯一的用户拥有的引用移动到 CUDA 图。
当引用关联到 CUDA 图时,CUDA 将自动管理图操作。克隆的 cudaGraph_t
保留源 cudaGraph_t
拥有的每个引用的副本,具有相同的多重性。实例化的 cudaGraphExec_t
保留源 cudaGraph_t
中每个引用的副本。当 cudaGraphExec_t
在未同步的情况下被销毁时,引用将保留到执行完成。
这是一个示例用法。
cudaGraph_t graph; // Preexisting graph
Object *object = new Object; // C++ object with possibly nontrivial destructor
cudaUserObject_t cuObject;
cudaUserObjectCreate(
&cuObject,
object, // Here we use a CUDA-provided template wrapper for this API,
// which supplies a callback to delete the C++ object pointer
1, // Initial refcount
cudaUserObjectNoDestructorSync // Acknowledge that the callback cannot be
// waited on via CUDA
);
cudaGraphRetainUserObject(
graph,
cuObject,
1, // Number of references
cudaGraphUserObjectMove // Transfer a reference owned by the caller (do
// not modify the total reference count)
);
// No more references owned by this thread; no need to call release API
cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0); // Will retain a
// new reference
cudaGraphDestroy(graph); // graphExec still owns a reference
cudaGraphLaunch(graphExec, 0); // Async launch has access to the user objects
cudaGraphExecDestroy(graphExec); // Launch is not synchronized; the release
// will be deferred if needed
cudaStreamSynchronize(0); // After the launch is synchronized, the remaining
// reference is released and the destructor will
// execute. Note this happens asynchronously.
// If the destructor callback had signaled a synchronization object, it would
// be safe to wait on it at this point.
子图节点中的图所拥有的引用与子图相关联,而不是与父图相关联。如果更新或删除子图,则引用会相应更改。如果使用 cudaGraphExecUpdate
或 cudaGraphExecChildGraphNodeSetParams
更新可执行图或子图,则会克隆新源图中的引用并替换目标图中的引用。在任何一种情况下,如果先前的启动不同步,则将保留任何将被释放的引用,直到启动完成执行。
目前没有通过 CUDA API 等待用户对象析构函数的机制。用户可以从析构代码中手动发出同步对象的信号。另外,从析构函数调用 CUDA API 是不合法的,类似于对 cudaLaunchHostFunc
的限制。这是为了避免阻塞 CUDA 内部共享线程并阻止前进。如果依赖是一种方式并且执行调用的线程不能阻止 CUDA 工作的前进进度,则向另一个线程发出执行 API 调用的信号是合法的。
用户对象是使用 cudaUserObjectCreate
创建的,这是浏览相关 API 的一个很好的起点。
边栏推荐
- The difference between session and cookies
- Google Play Academy 组队 PK 赛,正式开赛!
- Lucene Basics
- Compte à rebours de 3 jours pour l'inscription à l'accélérateur de démarrage Google Sea, Guide de démarrage collecté à l'avance!
- 日志(常用的日志框架)
- JWT tool class
- 如何使用MITMPROXy
- Linear DP (split)
- 深入学习JVM底层(五):类加载机制
- LeetCode 47. Full arrangement II
猜你喜欢
Support new and old imperial CMS collection and warehousing tutorials
Shenji Bailian 3.54-dichotomy of dyeing judgment
Shenji Bailian 3.53-kruskal
递归(迷宫问题、8皇后问题)
链表(线性结构)
线性dp(拆分篇)
借力 Google Cloud 基础设施和着陆区,构建企业级云原生卓越运营能力
Monitoring uplink of VRRP
亚马逊aws数据湖工作之坑1
In depth understanding of JUC concurrency (I) what is JUC
随机推荐
Shenji Bailian 3.52-prim
The difference between session and cookies
【每日一题】—华为机试01
Google Go to sea entrepreneurship accelerator registration countdown 3 days, entrepreneurs pass through the guide in advance collection!
BGP中的状态机
LeetCode 78. subset
深入学习JVM底层(五):类加载机制
一起学习SQL中各种join以及它们的区别
Mock simulate the background return data with mockjs
锐捷EBGP 配置案例
MySQL的10大经典错误
Cglib代理-代码增强测试
Contest3147 - game 38 of 2021 Freshmen's personal training match_ E: Listen to songs and know music
Leverage Google cloud infrastructure and landing area to build enterprise level cloud native excellent operation capability
Sumo tutorial Hello World
WLAN相关知识点总结
程序员的自我修养—找工作反思篇
找到页面当前元素z-index最高的数值
LeetCode 47. 全排列 II
Format check JS