已将cudaMalloc更改为异步？

我读过其他地方cudaMalloc将在内核之间进行同步。（例如will cudaMalloc synchronize host and device?）但是，我只是测试了这个代码，并根据我在视觉分析器中看到的情况，似乎cudaMalloc不同步。如果将cudaFree添加到循环中，那么确实会同步。我正在使用CUDA 7.5。有谁知道cudaMalloc是否改变了它的行为？或者我错过了一些微妙之处？非常感谢！已将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 
    } 
}

来源

2016-03-15 hildy

@RobertCrovella：在Maxwell设备上我正在测试它，它看起来像cudaMalloc不是同步调用。当我在上面的代码中修复内核中的明显缺陷并添加一些明智的nanosleep调用来扩展配置文件时间轴时，我看到cudaMalloc正在运行，而慢速和快速内核的两个实例在GPU上都处于活动状态。这可能是一个探查器问题，但它看起来不像我的眼睛：http://pastebin.com/rC8XxKmT – talonmies

我想我应该仔细看看。我也见证了它。更令人好奇的是，如果我让两个内核实际上都依赖于'dev_a'并且使得这两个内核实际上触及'dev_a'中的全局内容，我*仍然*见证了前两个内核的重叠，'cudaMalloc'在两者之间操作。我无法解释这一点。 –

你的方法是有缺陷的，但你的结论看起来是正确的，我（如果你看一下您的个人资料，你应该看到多空双方的内核正在采取相同的时间量和运行非常很快，因为积极的编译器优化消除了这两种情况下的所有代码）。

我把你的例子为更合理

#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上运行]

在一个相当快的麦克斯韦设备（GTX970），我见回呼中的cudaMalloc调用与配置文件跟踪中仍然执行的slowKernel调用重叠，然后在其他流中调用fastKernel调用。我愿意接受最初的结论，即时间微小变化可能会导致您在破碎示例中看到的效果。但是，在此代码中，主机和设备跟踪之间0.5秒的时间同步时间似乎非常不可能。您可能需要更改burntime调用的持续时间才能获得相同的效果，具体取决于GPU的速度。

所以这是一个非常长的说法，是的，它看起来像是Linux上的一个非同步调用，使用CUDA 7.5和Maxwell设备。我并不认为情况总是如此，但是再次，文件从来没有，尽我所知，说是否应该阻止/同步与否。我无法访问较旧的CUDA版本和支持的硬件，以查看此示例使用较旧的驱动程序和Fermi或Kepler设备的效果。

来源

2016-03-16 16:01:04 talonmies

感谢您对它进行测试。我会注意到，在Windows上使用CUDA 7.5时，原始内核不会得到优化，因为它是值得的。除非CUDA视觉轮廓仪向我说谎他们正在拍摄多少时间。 – hildy

@hildy：当我使用VS2012在Windows上编译内核时，我得到了两个空的存根，就像在linux上一样。你使用的是调试还是发布版本？无论如何，如果不直接与NVIDIA谈论这件事，我不会在这里看到更多。如果你能够将这个问题从无人答复的队列中提取出来（这是一个维基条目，我没有为你赢得任何声誉），那么这将对你有所帮助。 – talonmies

这可能是我使用调试版本。我相信应该仍然正确地测试libcuda的底层同步行为。 – hildy

已将cudaMalloc更改为异步？

回答

相关问题