在 CUDA9 中，"cudaMemcpyAsync()"既是设备又是主机功能吗？

In CUDA9, is "cudaMemcpyAsync()" both a device and a host function?

本文关键字：主机 功能 CUDA9 cudaMemcpyAsync      更新时间：2023-10-16

根据官方 CUDA 文档，我们有

__host__ ​ __device__ ​cudaError_t cudaMemcpyAsync ( void* dst, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0 )

这意味着它既是主机功能又是设备功能。 但是，在我的本地 Linux 盒子上的实际安装中，我看到/usr/local/cuda/include/cuda_runtime_api.h：

/** CUDA Runtime API Version */
#define CUDART_VERSION  9000
// Many lines away...
extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream __dv(0));

这似乎暗示它严格是一个主机函数。

我试图编译一个调用cudaMemcpyAsync()的简单内核，但得到了错误

streaming.cu(338)： 错误： 调用__host__ 函数("cudaMemcpyAsync") from a __global__ 不允许使用函数("loopy_plus_one")

这是另一个证据。

所以我真的很困惑：文档不正确，还是我的 CUDA 安装已过时？

编辑：更新 - 如果我更改我的编译命令以显式指定sm_60，即nvcc -arch=sm_60 -o out ./src.cu，则编译错误消失，但弹出一个新错误：

ptxas 致命：未解决的外部函数"cudaMemcpyAsync">

CUDA设备运行时 API 中有一个cudaMemcpyAsync的设备实现，您可以在此处的编程指南中看到该实现。在那里，在动态并行性的介绍部分中，它指出

动态并行性仅受具有计算能力的设备支持 3.5 及更高

版本

在文档中，它还记录了设备运行时 API 内存函数的用法：

关于所有memcpy/memset函数的说明：

  • 仅支持异步内存/设置功能
  • 仅允许设备到设备的内存
  • 不能传入本地或共享内存指针

您还可以找到有关如何编译和链接使用设备运行时 API 的代码的确切说明：

CUDA 程序自动与主机运行时库链接 使用 NVCC 编译时，但设备运行时作为静态运行时提供 库必须与希望 使用它。

设备运行时作为静态库提供(cudadevrt.lib on Windows，libcudadevrt.a 在 Linux 和 MacOS 下)，其中 GPU 必须链接使用设备运行时的应用程序。的链接 设备库可以通过NVCC和/或NVlink完成。

因此，要完成这项工作，您必须做三件事：

  1. 在编译时选择至少具有计算能力 3.5 的物理目标体系结构
    2. 编译
  2. 时对设备代码使用单独的编译
  3. 链接 CUDA 设备运行时库

正是由于这三个原因(即不做任何原因)，您在尝试在内核代码中使用cudaMemcpyAsync时看到了编译和链接错误。

一旦我正确指定计算能力，它似乎就可以工作，

nvcc -arch=compute_60 -o out src.cu
相关文章：