cudaStreamSynchronize 之后的任何 CUDA 操作都会阻塞，直到所有流都完成

Question

在使用 NVIDIA Visual Profiler 分析我的 CUDA 应用程序时，我注意到cudaStreamSynchronize 之后的任何操作都会阻塞，直到所有流完成为止。这是非常奇怪的行为，因为如果cudaStreamSynchronize 返回，则意味着流已完成，对吗？这是我的伪代码：

std::list<std::thread> waitingThreads;

void startKernelsAsync() {
    for (int i = 0; i < 200; ++i) {
        cudaHostAlloc(cpuPinnedMemory, size, cudaHostAllocDefault);
        memcpy(cpuPinnedMemory, data, size);
        cudaMalloc(gpuMemory);

        cudaStreamCreate(&stream);
        cudaMemcpyAsync(gpuMemory, cpuPinnedMemory, size, cudaMemcpyHostToDevice, stream);
        runKernel<<<32, 32, 0, stream>>>(gpuMemory);
        cudaMemcpyAsync(cpuPinnedMemory, gpuMemory, size, cudaMemcpyDeviceToHost, stream);

        waitingThreads.push_back(std::move(std::thread(waitForFinish, cpuPinnedMemory, stream)));
    }

    while (waitingThreads.size() > 0) {
        waitingThreads.front().join();
        waitingThreads.pop_front();
    }
}

void waitForFinish(void* cpuPinnedMemory, cudaStream_t stream, ...) {
    cudaStreamSynchronize(stream);
    cudaStreamDestroy(stream);  // <== This blocks until all streams are finished.
    memcpy(data, cpuPinnedMemory, size);
    cudaFreeHost(cpuPinnedMemory);
    cudaFree(gpuMemory);
}

如果我把cudaFreeHost放在cudaStreamDestroy之前，那么它就变成了阻塞操作。

这里有什么概念上的错误吗？

编辑：我发现了另一种奇怪的行为，有时它会在处理流的过程中解除阻塞，然后处理其余的流。

正常行为：

奇怪的行为（经常发生）：

EDIT2：我在 CUDA 6.0 上使用计算能力 3.5 的 Tesla K40c 卡进行测试。

正如 cmets 中所建议的，减少流的数量可能是可行的，但是在我的应用程序中，内存传输非常快，我想主要使用流来动态调度 GPU 的工作。问题是，在流完成后，我需要从固定内存中下载数据并清除分配的内存以用于似乎阻塞操作的更多流。

我为每个数据集使用一个流，因为每个数据集都有不同的大小，并且处理时间非常长。

有什么办法解决这个问题吗？

Answer 1

我还没有找到操作阻塞的原因，但我得出的结论是我对此无能为力，所以我决定实施内存和流池（如 cmets 中所建议的那样）以重用 GPU 内存、固定 CPU 内存和流，以避免任何类型的删除。

如果有人对此感兴趣，这是我的解决方案。启动内核的行为类似于调度内核的异步操作，并在内核完成后调用回调。

std::vector<Instance*> m_idleInstances;
std::vector<Instance*> m_workingInstances;

void startKernelAsync(...) {
    // Search for finished stream.
    while (m_idleInstances.size() == 0) {
        findFinishedInstance();
        if (m_idleInstances.size() == 0) {
            std::chrono::milliseconds dur(10);
            std::this_thread::sleep_for(dur);
        }
    }

    Instance* instance = m_idleInstances.back();
    m_idleInstances.pop_back();

    // Fill CPU pinned memory

    cudaMemcpyAsync(..., stream);
    runKernel<<<32, 32, 0, stream>>>(gpuMemory);
    cudaMemcpyAsync(..., stream);

    m_workingInstances.push_back(clusteringInstance);
}

void findFinishedInstance() {
    for (auto it = m_workingInstances.begin(); it != m_workingInstances.end();) {
        Instance* inst = *it;
        cudaError_t status = cudaStreamQuery(inst->stream);
        if (status == cudaSuccess) {
            it = m_workingInstances.erase(it);
            m_callback(instance->clusterGroup);
            m_idleInstances.push_back(inst);
        }
        else {
            ++it;
        }
    }
}

然后等待大家完成：

virtual void waitForFinish() {
    while (m_workingInstances.size() > 0) {
        Instance* instance = m_workingInstances.back();
        m_workingInstances.pop_back();
        m_idleInstances.push_back(instance);
        cudaStreamSynchronize(instance->stream);
        finalizeInstance(instance);
    }
}

这是一个图形分析器，非常有用！

Answer 2

查看工具包随附的 Cuda C 编程指南 PDF 中的“隐式同步”规则列表。 （我的副本中的第 3.2.5.5.4 节，但您可能有不同的版本。）

如果您的 GPU 是“计算能力 3.0 或更低”，则有一些适用的特殊规则。我的猜测是 cudaStreamDestroy() 正在达到这些限制之一。