为什么第一个 cuda 内核不能与以前的 memcpy 重叠？

Question

这是一个演示。内核不能与之前的 cudaMemcpyAsync 重叠，尽管它们在不同的流中。

#include <iostream>
#include <cuda_runtime.h>

__global__ void warmUp(){
    int Id = blockIdx.x*blockDim.x+threadIdx.x;
    if(Id == 0){
        printf("warm up!");
    }
}
__global__ void kernel(){
    int Id = blockIdx.x*blockDim.x+threadIdx.x;
    if(Id == 0){
        long long x = 0;
        for(int i=0; i<1000000; i++){
            x += i>>1;
        }
        printf("kernel!%d\n", x);
    }
}

int main(){
    //warmUp<<<1,32>>>();
    int *data, *data_dev;
    int dataSize = pow(10, 7);
    cudaMallocHost(&data, dataSize*sizeof(int));
    cudaMalloc(&data_dev, dataSize*sizeof(int));
    
    cudaStream_t stream1, stream2;
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);
    cudaMemcpyAsync(data_dev, data, dataSize*sizeof(int), cudaMemcpyHostToDevice, stream1);
    kernel<<<1, 32, 0, stream2>>>();
}

Visual Profiler show

经过一些尝试，我发现这是因为它是第一个内核调用。

取消注释 warmUp<<<1,32>>>();、Visual Profiler show、重叠！

为什么？

Answer 1

CUDA 使用 lazy initialization。因此，当您第一次执行特定操作或特定操作类型时，行为可能与您预期的不同。

操作 will/should 工作“正确”，但性能测量可能与您预期的不同。

与链接的文章相反，确实没有指定的公式来强制延迟初始化完成，而不执行您打算执行的实际工作。

如果您打算对应用程序做的唯一事情是启动单个内核，那么让该内核与先前的复制操作重叠对我来说似乎没有多大意义。无论如何，您应该预料到在所有操作以预期速度或以预期方式进行之前必须进行设备初始化。

延迟初始化行为可能因 CUDA 版本、平台（例如 OS）和 GPU 类型而异。

此外，内核启动是异步的。所以这个特定的编码模式：

int main(){
    ...
    kernel<<<1, 32, 0, stream2>>>();
}

在 CUDA 中一般不推荐，特别是在使用分析器时不推荐。您的代码应该为所有发布的工作提供正确完成的机会，以便探查器提供有用的结果。如果要分析这种类型的模式，您应该在代码末尾提供 cudaDeviceSynchronize() 或类似操作。

我也不建议对发出 printf 调用的内核进行性能分析。 printf 调用强加了 additional host/device synchronization behavior/needs，这可能会造成混淆；预测其对性能的影响并不容易。

为什么第一个 cuda 内核不能与以前的 memcpy 重叠？

why the first cuda kernel cannot overlap with previous memcpy?

cuda