标签: uva

我熟悉一个配备Pascal P100 GPU + Nvlink的新集群.我写了一个ping-pong程序来测试gpu < - > gpu和gpu < - > cpu带宽和点对点访问.(我知道cuda样本包含这样的程序,但我想自己做以便更好地理解.)Nvlink带宽看似合理(双向约35 GB/s,理论最大值为40).然而,在调试乒乓球时,我发现了一些奇怪的行为.

首先,无论我指定什么cudaMemcpyKind,cudaMemcpyAsync都会成功,例如,如果cudaMemcpyAsync正在将内存从主机复制到设备,即使我将cudaMemcpyDeviceToHost作为类型传递,它也会成功.

其次,当主机内存没有页面锁定时,cudaMemcpyAsync会执行以下操作:

• 将内存从主机复制到设备似乎成功(没有segfaults或cuda运行时错误,数据似乎正确传输).
• 将内存从设备复制到主机会无声地失败:没有发生段错误,并且memcpy返回cudaSuccess后cudaDeviceSynchronize,但检查数据显示gpu上的数据没有正确传输到主机.

这种行为是期待的吗？我已经包含了一个在我的系统上演示它的最小工作示例代码(示例不是乒乓应用程序,它所做的只是使用各种参数测试cudaMemcpyAsync).

P100s启用了UVA,因此cudaMemcpyAsync简单地推断src和dst指针的位置并忽略cudaMemcpyKind参数是合理的.但是,我不确定为什么cudaMemcpyAsync无法为非页面锁定的主机内存引发错误.我的印象是严格的禁忌.

#include <stdio.h>
#include <cuda_runtime.h>
#include <stdlib.h>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

__global__ void checkDataDevice( int* current, int* next, int expected_current_val, int n )
{
  int tid = threadIdx.x + …