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 的一个很好的起点。