1

将 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");
}
4

1 回答 1

2

概括地说，在设备上：

  • 您需要将数据写入设备全局内存（之前使用 分配cudaMalloc）或直接写入主机内存（之前分配使用cudaHostAlloc）
  • 您可能希望从单个线程块将所有数据写入该区域，以确保在以下步骤之前写入所有数据
  • 然后，您需要在执行以下步骤之前发出threadfence()（如果您使用设备全局内存）或调用（如果使用主机内存）threadfence_system()
  • 接下来，您将写入设备全局内存或主机内存中的一个特殊位置，我们称其为邮箱位置，其中有一个特定值表示数据已准备好。
  • 可选择发出另一个 threadfence 或 threadfence_system 调用

在主机上：

  • 在启动内核之前，主机需要将邮箱位置设置为默认值。
  • 启动内核后，主机线程将需要“轮询”邮箱位置，寻找指示数据准备好的具体值
  • 一旦看到具体的值，说明数据准备好了，主机就可以消费数据了
  • 或者，如果您想重复此过程，主机可以将邮箱位置重置为默认值。设备可以在用新数据更新数据块之前检查这个默认值。

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

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

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

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

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

  • 而（更多帧）
    • 处理一帧
    • 信号主机
    • 递增帧指针

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

  • 而（更多帧）
    • 调用内核处理帧
    • 消费帧
    • 递增帧指针

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

于 2013-05-06T10:16:21.003 回答