是否需要在 CUDA 程序中复制 DeviceToDevice？

Question

我正在做以下两个操作：

两个数组相加=> a + b = AddResult
两个数组相乘=> AddResult * a = MultiplyResult

在上面的逻辑中，AddResult是一个中间结果，作为下一次乘法运算的输入。

#define N 4096         // size of array

__global__ void add(const int* a, const int* b, int* c)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) 
    {
        c[tid] = a[tid] + b[tid];
    }
}

__global__ void multiply(const int* a, const int* b, int* c)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) 
    {
        c[tid] = a[tid] * b[tid];
    }
}

int main() 
{
    int T = 1024, B = 4;            // threads per block and blocks per grid
    int a[N], b[N], c[N], d[N], e[N];
    int* dev_a, * dev_b, * dev_AddResult, * dev_Temp, * dev_MultiplyResult;

    cudaMalloc((void**)&dev_a, N * sizeof(int));
    cudaMalloc((void**)&dev_b, N * sizeof(int));
    cudaMalloc((void**)&dev_AddResult, N * sizeof(int));
    cudaMalloc((void**)&dev_Temp, N * sizeof(int));
    cudaMalloc((void**)&dev_MultiplyResult, N * sizeof(int));

    for (int i = 0; i < N; i++) 
    {    
        // load arrays with some numbers
        a[i] = i;
        b[i] = i * 1;
    }

    cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_AddResult, c, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_Temp, d, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_MultiplyResult, e, N * sizeof(int), cudaMemcpyHostToDevice);

    //ADD
    add << <B, T >> > (dev_a, dev_b, dev_AddResult);
    cudaDeviceSynchronize();

    //Multiply
    cudaMemcpy(dev_Temp, dev_AddResult, N * sizeof(int), cudaMemcpyDeviceToDevice); //<---------DO I REALLY NEED THIS?
    multiply << <B, T >> > (dev_a, dev_Temp, dev_MultiplyResult);
    //multiply << <B, T >> > (dev_a, dev_AddResult, dev_MultiplyResult);
    
    //Copy Final Results D to H
    cudaMemcpy(e, dev_MultiplyResult, N * sizeof(int), cudaMemcpyDeviceToHost);


    for (int i = 0; i < N; i++) 
    {
        printf("(%d+%d)*%d=%d\n", a[i], b[i], a[i], e[i]);
    }

    // clean up
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_AddResult);
    cudaFree(dev_Temp);
    cudaFree(dev_MultiplyResult);

    return 0;
}

在上面的示例代码中，我将加法结果（即dev_AddResult）传输到另一个设备数组（即dev_Temp）以执行乘法运算。

问题： 由于加法结果数组（即 dev_AddResult）已经在 GPU 设备上，我真的需要将它转移到另一个数组吗？我已经尝试通过直接提供 dev_AddResult 作为输入来执行下一个内核，它产生了相同的结果。直接将一个内核的输出作为下一个内核的输入有什么风险吗？有什么最佳实践可以遵循吗？

Answer 1

是的，对于您所展示的情况，您可以使用一个内核的“输出”作为下一个内核的“输入”，而无需任何复制。您已经完成并确认它有效，所以我将省去任何示例。无论如何，更改都是微不足道的 - 消除中间的 cudaMemcpy 操作，并使用相同的 dev_AddResult 指针代替乘法内核调用中的 dev_Temp 指针。

关于“风险”，对于您给出的示例，我不知道有任何风险。从该示例转向可能更一般的用法，您需要确保在其他地方使用之前完成添加输出计算。

您的示例已经冗余地执行了此操作，至少使用了 2 种机制：

干预 cudaDeviceSynchronize() - 这会强制完成之前发布的工作
流语义 - 流语义的一个规则是发布到特定流中的工作将按发布顺序执行。发布到流 X 中的项目 B，直到先前发布到流 X 中的项目 A 完成后才会开始。

所以在这种情况下您并不需要 cudaDeviceSynchronize()。从功能的角度来看，它并没有“伤害”任何东西，但它可能会使整体执行时间增加几微秒。

更一般地说，如果您将加法和乘法内核发布到单独的流中，那么 CUDA 不提供执行顺序保证，即使您在加法内核之后“发布”乘法内核。

在那种情况下（不是你这里的那个）如果你需要乘法运算来使用之前计算的加法结果，你需要以某种方式强制执行（在乘法内核之前强制完成加法内核） .您已经在此处展示了一种方法，即使用同步调用。

是否需要在 CUDA 程序中复制 DeviceToDevice？

Is DeviceToDevice copy in a CUDA program needed?

cuda