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

Question

根据official CUDA doc，我们有

__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__ 来自 __global__ 的函数（“cudaMemcpyAsync”） 不允许使用函数（“loopy_plus_one”）

这是另一个证据。

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

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

ptxas 致命：未解析的外部函数“cudaMemcpyAsync”

Answer 1

在 CUDA device 运行时 API 中有一个cudaMemcpyAsync 的设备实现，您可以在编程指南here 中查看该实现。在那里，在动态并行的介绍部分中，notes

只有具有计算能力的设备支持动态并行 3.5及更高版本

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

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

• 仅支持异步 memcpy/set 函数
• 只允许设备到设备的 memcpy
• 不能传入本地或共享内存指针

您还可以找到确切的instructions，了解您必须如何编译和链接使用设备运行时 API 的代码：

CUDA 程序会自动与主机运行时库链接 使用 nvcc 编译时，但设备运行时以静态形式提供 必须显式链接到希望的程序的库 使用它。

设备运行时以静态库的形式提供（cudadevrt.lib on Windows、Linux 和 MacOS 下的 libcudadevrt.a），针对它的 GPU 必须链接使用设备运行时的应用程序。的链接 设备库可以通过 nvcc 和/或 nvlink 来完成。

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

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

正是由于这三个原因（即不执行其中任何一个），您在尝试在内核代码中使用 cudaMemcpyAsync 时会看到编译和链接错误。

Answer 2

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

nvcc -arch=compute_60 -o out src.cu