那些遇到过的问题

cudaMemcpyAsync奇怪的行为:1.cudaMemcpyKind没什么区别.2.复制失败,但是默默无闻

Mic*_*lli 0 cuda nvlink 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 + blockIdx.x*blockDim.x;
  for( int i = tid; i < n; i += blockDim.x*gridDim.x )
  {
    if( current[i] != expected_current_val )
      printf( "Error on device:  expected = %d, current[%d] = %d\n"
          , expected_current_val
          , i
          , current[i] );
    // Increment the data so the next copy is properly tested
    next[i] = current[i] + 1;
  }
}

void checkDataHost( int* current, int* next, int expected_current_val, int n )
{
  for( int i = 0; i < n; i++ )
  {
    if( current[i] != expected_current_val )
      printf( "Error on host:  expected = %d, current[%d] = %d\n"
          , expected_current_val
          , i
          , current[i] );
    // Increment the data so the next copy is properly tested
    next[i] = current[i] + 1;
  }
}

int main( int argc, char** argv )
{
  bool pagelocked = true;
  // invoking the executable with any additional argument(s) will turn off page locked memory, i.e.,
  // Run with pagelocked memory:  ./a.out
  // Run with ordinary malloc'd memory: ./a.out jkfdlsja
  if( argc > 1 )
    pagelocked = false;

  int copybytes = 1e8; // Ok to use int instead of size_t for 1e8.

  cudaStream_t* stream = (cudaStream_t*)malloc( sizeof(cudaStream_t) );
  cudaStreamCreate( stream );

  int* srcHost;
  int* dstHost;
  int* srcDevice;
  int* dstDevice;

  cudaMalloc( (void**)&srcDevice, copybytes );
  cudaMalloc( (void**)&dstDevice, copybytes );
  if( pagelocked )
  {
    printf( "Using page locked memory\n" );
    cudaMallocHost( (void**)&srcHost, copybytes );
    cudaMallocHost( (void**)&dstHost, copybytes );
  }
  else
  {
    printf( "Using non page locked memory\n" );
    srcHost = (int*)malloc( copybytes );
    dstHost = (int*)malloc( copybytes );
  }

  for( int i = 0; i < copybytes/sizeof(int); i++ )
    srcHost[i] = 1;

  cudaMemcpyKind kinds[4];
  kinds[0] = cudaMemcpyHostToDevice;
  kinds[1] = cudaMemcpyDeviceToHost;
  kinds[2] = cudaMemcpyHostToHost;
  kinds[3] = cudaMemcpyDeviceToDevice;

  // Test cudaMemcpyAsync in both directions,
  // iterating through all "cudaMemcpyKinds" to verify
  // that they don't matter.
  int expected_current_val = 1;
  for( int kind = 0; kind<4; kind++ )
  {
    // Host to device copy 
    cudaMemcpyAsync( dstDevice
        , srcHost
        , copybytes
        , kinds[kind]
        , *stream );
    gpuErrchk( cudaDeviceSynchronize() );

    checkDataDevice<<<56*8,256>>>( dstDevice
        , srcDevice
        , expected_current_val
        , copybytes/sizeof(int) );
    expected_current_val++;

    // Device to host copy
    cudaMemcpyAsync( dstHost
        , srcDevice
        , copybytes
        , kinds[kind]
        , *stream );
    gpuErrchk( cudaDeviceSynchronize() );

    checkDataHost( dstHost
        , srcHost
        , expected_current_val
        , copybytes/sizeof(int) );
    expected_current_val++;
  }

  free( stream );

  cudaFree( srcDevice );
  cudaFree( dstDevice );
  if( pagelocked )
  {
    cudaFreeHost( srcHost );
    cudaFreeHost( dstHost );
  }
  else
  {
    free( srcHost );
    free( dstHost );
  }

  return 0;
}
Run Code Online (Sandbox Code Playgroud)

Rob*_*lla 5

当遇到CUDA代码问题时,我强烈建议使用严格的(==检查每个调用返回代码)正确的CUDA错误检查.

您的错误检查存在缺陷,这些缺陷会导致您的一些困惑.

首先,在页面锁定的情况下,给定(映射)指针在主机和设备上都是可访问/有效的.因此,每个可能的方向枚举(H2D,D2H,D2D,H2H)都是合法且有效的.因此,不会返回任何错误,并且复制操作成功.

在非页面锁定的情况下,上述情况并非如此,因此一般而言,指示的传送方向与从指针检查的隐含传送方向更好地匹配.如果没有,cudaMemcpyAsync则返回错误代码(cudaErrorInvalidValue== 11).在您的情况下,您忽略此错误结果.你可以自己证明这一点,如果你有足够的耐心(如果你只是标记了第一个错误,而不是打印10M +元素中的每个不匹配会更好),运行你的代码cuda-memcheck(无论何时你都要做的另一件好事)使用CUDA代码时遇到问题)或者只是进行适当,严格的错误检查.

当cudaMemcpyAsync操作指示失败时,操作未成功完成,因此不会复制数据,并且数据检查表明不匹配.希望现在这并不奇怪,因为预期的复制操作实际上并没有发生(也没有"默默地"失败).

也许你很困惑,认为在任何类型的异步操作上捕获错误的方法是执行a cudaDeviceSynchronize然后检查错误.

这不正确cudaMemcpyAsync.在调用cudaMemcpyAsync操作时可以检测到的错误将由调用本身立即返回,并且由于此类错误是非粘性的,因此不会因后续CUDA调用而返回(显然).

这个故事的主旨:

  1. 做适当的CUDA错误检查.严格.
  2. 运行你的代码cuda-memcheck.

这是一个完整的示例,对代码进行了一些微不足道的修改,使输出在失败的情况下变得"理智",证明在失败的情况下指示了错误:

$ cat t153.cu
#include <stdio.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 + blockIdx.x*blockDim.x;
  for( int i = tid; i < n; i += blockDim.x*gridDim.x )
  {
    if( current[i] != expected_current_val )
      printf( "Error on device:  expected = %d, current[%d] = %d\n"
          , expected_current_val
          , i
          , current[i] );
    // Increment the data so the next copy is properly tested
    next[i] = current[i] + 1;
  }
}

void checkDataHost( int* current, int* next, int expected_current_val, int n )
{
  for( int i = 0; i < n; i++ )
  {
    if( current[i] != expected_current_val ){
      printf( "Error on host:  expected = %d, current[%d] = %d\n"
          , expected_current_val
          , i
          , current[i] );
      exit(0);}
    // Increment the data so the next copy is properly tested
    next[i] = current[i] + 1;
  }
}

int main( int argc, char** argv )
{
  bool pagelocked = true;
  // invoking the executable with any additional argument(s) will turn off page locked memory, i.e.,
  // Run with pagelocked memory:  ./a.out
  // Run with ordinary malloc'd memory: ./a.out jkfdlsja
  if( argc > 1 )
    pagelocked = false;

  int copybytes = 1e8; // Ok to use int instead of size_t for 1e8.

  cudaStream_t* stream = (cudaStream_t*)malloc( sizeof(cudaStream_t) );
  cudaStreamCreate( stream );

  int* srcHost;
  int* dstHost;
  int* srcDevice;
  int* dstDevice;

  cudaMalloc( (void**)&srcDevice, copybytes );
  cudaMalloc( (void**)&dstDevice, copybytes );
  if( pagelocked )
  {
    printf( "Using page locked memory\n" );
    cudaMallocHost( (void**)&srcHost, copybytes );
    cudaMallocHost( (void**)&dstHost, copybytes );
  }
  else
  {
    printf( "Using non page locked memory\n" );
    srcHost = (int*)malloc( copybytes );
    dstHost = (int*)malloc( copybytes );
  }

  for( int i = 0; i < copybytes/sizeof(int); i++ )
    srcHost[i] = 1;

  cudaMemcpyKind kinds[4];
  kinds[0] = cudaMemcpyHostToDevice;
  kinds[1] = cudaMemcpyDeviceToHost;
  kinds[2] = cudaMemcpyHostToHost;
  kinds[3] = cudaMemcpyDeviceToDevice;

  // Test cudaMemcpyAsync in both directions,
  // iterating through all "cudaMemcpyKinds" to verify
  // that they don't matter.
  int expected_current_val = 1;
  for( int kind = 0; kind<4; kind++ )
  {
    // Host to device copy
    cudaMemcpyAsync( dstDevice
        , srcHost
        , copybytes
        , kinds[kind]
        , *stream );
    gpuErrchk( cudaDeviceSynchronize() );

    checkDataDevice<<<56*8,256>>>( dstDevice
        , srcDevice
        , expected_current_val
        , copybytes/sizeof(int) );
    expected_current_val++;

    // Device to host copy
    cudaMemcpyAsync( dstHost
        , srcDevice
        , copybytes
        , kinds[kind]
        , *stream );
    gpuErrchk( cudaDeviceSynchronize() );

    checkDataHost( dstHost
        , srcHost
        , expected_current_val
        , copybytes/sizeof(int) );
    expected_current_val++;
  }

  free( stream );

  cudaFree( srcDevice );
  cudaFree( dstDevice );
  if( pagelocked )
  {
    cudaFreeHost( srcHost );
    cudaFreeHost( dstHost );
  }
  else
  {
    free( srcHost );
    free( dstHost );
  }

  return 0;
}
$ nvcc -arch=sm_61 -o t153 t153.cu
$ cuda-memcheck ./t153 a
========= CUDA-MEMCHECK
Using non page locked memory
========= Program hit cudaErrorInvalidValue (error 11) due to "invalid argument" on CUDA API call to cudaMemcpyAsync.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 [0x2ef423]
=========     Host Frame:./t153 [0x489a3]
=========     Host Frame:./t153 [0x2e11]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xf5) [0x21ec5]
=========     Host Frame:./t153 [0x2a49]
=========
Error on host:  expected = 2, current[0] = 0
========= ERROR SUMMARY: 1 error
$
Run Code Online (Sandbox Code Playgroud)

归档时间：

查看次数：

674 次

最近记录：

8 年，7 月 前
使用CUDA运行时API检查错误的规范方法是什么？ 252
更多相关链接
相关归档
CUDA内核的参数 8
opencl命令队列如何工作,我可以问什么 6
什么是“Visual Studio 配置文件”，为什么 nvcc 寻找它，为什么找不到它，如何修复？ 2
CUDA线程执行顺序 1
检查CUDA存在需要CUDA的程序？ 1
如何从C++头文件中调用CUDA文件？ 1
优化字节操作CUDA 1
使用cuRAND在函数内生成随机数,无需预先分配 1
为什么 cudaPointerGetAttributes() 为主机指针返回无效参数？ 1
通过 conda 安装 cudatoolkit 11.7 失败 -1
难疑归档
如何在Git中克隆所有远程分支？ 3987
如何在JavaScript中创建字符串大写的第一个字母？ 3565
如何列出目录的所有文件？ 3474
什么是依赖注入？ 2984
如何从列表列表中制作一个平面列表？ 2950
如何删除GitHub上的提交？ 1619
虚拟成员在构造函数中调用 1270
参考 - 这个错误在PHP中意味着什么？ 1071
jQuery的jquery-1.10.2.min.map触发了404(未找到) 1051
如何检查对象是否在JavaScript中有密钥？ 1047