cudaMalloc改为异步吗？

2019年7月27日 228次阅读

我在其他地方读到cudaMalloc将跨内核同步.

(例如
will cudaMalloc synchronize host and device?)

但是,我刚刚测试了这段代码并根据我在visual Profiler中看到的内容,似乎cudaMalloc没有同步.如果你将cudaFree添加到循环中,它会同步.我正在使用CUDA 7.5.有谁知道cudaMalloc是否改变了它的行为？还是我错过了一些微妙之处？非常感谢！

__global__ void slowKernel()
{
  float input = 5;
  for( int i = 0; i < 1000000; i++ ){
    input = input * .9999999;
  }
}

__global__ void fastKernel()
{
  float input = 5;
  for( int i = 0; i < 100000; i++ ){
    input = input * .9999999;
  }
}

void mallocSynchronize(){
  cudaStream_t stream1, stream2;
  cudaStreamCreate( &stream1 );
  cudaStreamCreate( &stream2 );
  slowKernel <<<1, 1, 0, stream1 >>>();
  int *dev_a = 0;
  for( int i = 0; i < 10; i++ ){
    cudaMalloc( &dev_a, 4 * 1024 * 1024 );
    fastKernel <<<1, 1, 0, stream2 >>>();
    // cudaFree( dev_a ); // If you uncomment this, the second fastKernel launch will wait until slowKernel completes
  }
}

最佳答案 你的方法是有缺陷的,但你的结论对我来说是正确的(如果你看一下你的个人资料数据,你应该看到长核和短核都花费相同的时间并且运行得非常快,因为积极的编译器优化正在消除所有的代码在这两种情况下).

我把你的榜样变成了更合理的东西

#include <time.h>
__global__ void slowKernel(float *output, bool write=false)
{
    float input = 5;
#pragma unroll
    for( int i = 0; i < 10000000; i++ ){
        input = input * .9999999;
    }
    if (write) *output -= input;
}

__global__ void fastKernel(float *output, bool write=false)
{
    float input = 5;
#pragma unroll
    for( int i = 0; i < 100000; i++ ){
        input = input * .9999999;
    }
    if (write) *output -= input;
}

void burntime(long val) {
    struct timespec tv[] = {{0, val}};
    nanosleep(tv, 0);
}

void mallocSynchronize(){
    cudaStream_t stream1, stream2;
    cudaStreamCreate( &stream1 );
    cudaStreamCreate( &stream2 );
    const size_t sz = 1 << 21;
    slowKernel <<<1, 1, 0, stream1 >>>((float *)(0));
    burntime(500000000L); // 500ms wait - slowKernel around 1300ms
    int *dev_a = 0;
    for( int i = 0; i < 10; i++ ){
        cudaMalloc( &dev_a, sz );
        fastKernel <<<1, 1, 0, stream2 >>>((float *)(0));
        burntime(1000000L); // 1ms wait - fastKernel around 15ms
    }
}

int main()
{
    mallocSynchronize();
    cudaDeviceSynchronize();
    cudaDeviceReset();
    return 0;
}

[注意需要POSIX时间函数,所以这不会在Windows上运行]

在相当快的Maxwell设备(GTX970)上,我看到循环中的cudaMalloc调用与配置文件跟踪中仍在执行的slowKernel调用重叠,然后在另一个流中运行fastKernel调用.我愿意接受最初的结论,即微小的时间变化可能会导致你在破碎的例子中看到的效果.但是,在此代码中,主机和设备跟踪之间0.5秒的同步时移似乎非常不可能.您可能需要更改burntime调用的持续时间以获得相同的效果,具体取决于GPU的速度.

所以这是一个很长的说法,是的,它看起来像是在Linux上使用CUDA 7.5和Maxwell设备的非同步调用.我不相信一直都是这样,但是再一次,文档从来没有,尽我所知,说是否应该阻止/同步.我无法访问较旧的CUDA版本和支持的硬件,以查看此示例对旧驱动程序和Fermi或Kepler设备的作用.

点赞