从设备写入主机并通知主机

Question

将 CUDA 5 与 VS 2012 和功能 3.5（Titan 和 K20）一起使用。

在内核执行的特定阶段，我想将生成的数据块发送到主机内存并通知主机数据已准备好，以便主机对其进行操作。

我不能等到内核执行结束才从设备读回数据，因为：

1. 数据一经计算就不再与设备相关，因此将其保留到最后没有意义。
2. 数据量太大，设备内存无法容纳，请等到最后。
3. 主机不必等到内核执行结束才开始处理数据。

您能否指出我必须采取的路径以及我必须使用哪些 cuda 概念和功能来实现我的要求？简而言之，如何写入主机并通知主机一个块数据已准备好供主机处理？

注意每个线程不与任何其他线程共享任何生成的数据，它们独立运行。所以，据我所知（如果我错了，请纠正我），块、线程和扭曲的概念不会影响问题。或者换句话说，如果它们有助于答案，我可以随意更改它们的组合。

以下是显示我正在尝试做的示例代码：

#pragma once
#include <conio.h>
#include <cstdio>
#include <cuda_runtime_api.h>

__global__ void Kernel(size_t length, float* hResult) 
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // Processing multiple data chunks
    for(int i = 0;i < length;i++)
    {
        // Once this is assigned, I don't need it on the device anymore.
        hResult[i + (tid * length)] = i * 100;
    }

}

void main()
{
    size_t length = 10;
    size_t threads = 2;
    float* hResult;
    // An array that will hold all data from all threads
    cudaMallocHost((void**)&hResult, threads * length * sizeof(float));
    Kernel<<<threads,1>>>(length, hResult);
    // I DO NOT want to wait to the end and block to get the data
    cudaError_t error = cudaDeviceSynchronize();
    if (error != cudaSuccess) { throw error; }
    for(int i = 0;i < threads * length;i++)
    {
        printf("%f\n", hResult[i]);;
    }
    cudaFreeHost(hResult);
    system("pause");
}

Answer 1

这是一种可能的方法。概括地说，在设备上：

• 您需要将数据写入设备全局内存（之前使用cudaMalloc 分配）或直接写入主机内存（之前使用cudaHostAlloc 分配）。应通过volatile 指针访问此内存。
• 您可能希望将所有数据从单个线程块写入此区域，以确保在以下步骤之前写入所有数据
• 然后，您需要在执行以下步骤之前发出threadfence()（如果您使用设备全局内存）或threadfence_system() 调用（如果使用主机内存）
• 接下来，您将写入设备全局内存或主机内存中的一个特殊位置，我们将其称为邮箱位置，其中的特定值表示数据已准备就绪。还应使用volatile 指针访问此位置。
• 可选择发出另一个 threadfence 或 threadfence_system 调用
• 对于接收端的设备内存使用情况，同样应使用volatile 指针访问两个区域（有效负载和“邮箱”）。

在主机上：

• 在启动内核之前，主机需要将邮箱位置设置为默认值。
• 启动内核后，主机线程将需要“轮询”邮箱位置，寻找指示数据准备就绪的具体值
• 一旦看到具体值，说明数据准备好了，主机就可以消费数据了
• （可选）如果您想重复此过程，主机可以将邮箱位置重置为默认值。设备可以在使用新数据更新数据块之前检查此默认值。
• 主机线程应使用volatile 指针访问邮箱位置和有效负载区域。

请注意，即使使用上述过程，如果数据是从多个线程块生成/创建的，仍然需要隐含的设备范围同步。唯一可用的直接设备范围同步是内核启动（特别是内核的完成）。从单个线程块复制数据只是将设备范围同步的要求移出这个特定序列（移到这个序列之前的某个地方）。

您给出的原因并没有真正向我表明无法重构代码以在内核启动的基础上创建数据，这将巧妙地解决这些问题并消除对上述过程的需要也是。

编辑： 回答 cmets 中的问题。 如果没有具体示例，很难更具体地说明如何重构代码以在每次内核调用时传递一个数据块。

让我们以一个图像处理案例为例，其中我有一个 30 帧的视频序列存储在全局内存中。内核会根据某种算法对每一帧进行处理，然后将处理后的数据提供给主机。

在您的建议中，内核处理完一帧后，它可以向主机发出数据准备就绪的信号，然后继续处理下一帧。问题是，如果帧由多个线程块处理，则没有简单的方法可以知道所有线程块何时完成处理该帧。设备范围的同步屏障可能是需要的，但它并不方便存在，除非通过内核调用机制。但是，大概在这样的内核中，我们可能会有这样的序列：

• 同时（更多帧）
• • 处理帧
• • 信号主机
• • 增加帧指针

在重构的方法中，我们会将循环移到内核之外，以托管代码：

• 同时（更多帧）
• • 调用内核处理帧
• • 消费框架
• • 增加帧指针

通过这样做，内核标记需要知道帧处理何时完成的显式同步，并且可以消费数据。