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

Question

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

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

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

1. 一旦计算出的数据不再与设备相关的，所以没有点保持最后。
2. 数据大小太大而无法放在设备内存上，并等到结束。
3. 主机不应该等到内核执行结束才开始处理数据。

你能指出我我必须采取的路径和可能的CUDA概念和功能，我必须使用达到我的要求？简而言之，如何写入主机并通知主机块数据已准备好进行主机处理？

N.B.每个线程不会与任何其他线程共享任何生成的数据，它们将独立运行。所以，据我所知（如果我错了，请纠正我），块，线和经纱的概念不会影响问题。换句话说，如果他们帮助答案，我可以自由改变他们的组合。

下面是一个示例代码，显示我试图做的事：

#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）
• 您可能希望从单个线程块执行写入该区域的所有数据，以确保在所有的数据在以下步骤之前被写
• 然后你会想发出threadfence()（如果您使用的设备的全局存储器）或threadfence_system()通话（如果使用主机内存）以下步骤之前
• 接下来，您将写入设备全局内存或主机内存中的特殊位置，我们称之为邮箱位置，并指定数据已准备就绪的特定值。
• 可选问题的另一个threadfence或threadfence_system叫

在主机：

• 之前启动内核，主机需要将邮箱位置设置为默认值。
• 启动内核后，主机线程将需要“轮询”邮箱位置，查找指示数据已准备好的特定值
• 一旦看到特定值，表示数据已准备好，主机就可以消费数据
• （可选）如果要重复此过程，主机可以将邮箱位置重置为默认值。在用新数据更新数据块之前，设备可以检查此默认值。

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

你给的理由并不真的暗示我的代码不能被重构来在内核上创建数据 - 在内核启动的基础上启动，这将很好地解决这些问题，并且不需要上述过程也是如此。

编辑：回应评论中的问题。 没有一个具体的例子，关于如何重构代码来为每个内核调用提供一个数据块很困难。

让我们来看一个图像处理的情况，我有一个30帧的视频序列存储在全局内存中。内核将根据某种算法处理每个帧，然后将处理后的数据提供给主机。

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

• 而（more_frames）
• • 过程框架
• • 信号主机
• • 增量帧指针

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

• 而（more_frames）
• • 调用内核处理框架
• • 消耗帧
• • 增量帧指针

通过这样做，内核标志着知道该帧何时处理完成所需的显式同步，并且可以对数据进行消耗。