第二次内核调用后cudaMemcpyAsync不同步 问题版本

编程问答 2022-06-09

问题描述

我的目标是设置通过引用传递到cuda内核的主机变量：

// nvcc test_cudamemcpyAsync.cu -rdc=true
#include <iostream>

__global__ void setHostvar(double& host_var) {
  double const var = 2.0;
  cudamemcpyAsync(&host_var,&var,sizeof(double),cudamemcpyDevicetoHost);
  // identifier "cudamemcpy" is undefined in device code
  // cudamemcpy(&host_var,cudamemcpyDevicetoHost);
}

int main() {
  double host_var = 1.0;

  setHostvar<<<1,1>>>(host_var);
  cudaDeviceSynchronize();
  std::cout << "host_var = " << host_var << std::endl;

  setHostvar<<<1,1>>>(host_var);
  cudaDeviceSynchronize();
  std::cout << "host_var = " << host_var << std::endl;

  return 0;
}

编译并运行：

$ nvcc test_cudamemcpyAsync.cu -rdc=true
$ ./a.out

输出：

host_var = 1
host_var = 1

除了对host_var = 1的异步调用之外，我还可以理解给定的异步内核调用的第一条输出线cudamemcpyAsync()。但是，我本以为第二个内核调用是在先前的异步调用完成之后执行的，但是host_var保持不变。

问题

  1. 我的期望有什么不对？
  2. 将引用/指针传递的主机变量设置到内核中的最佳/更好方法是什么？

版本

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
copyright (c) 2005-2020 NVIDIA Corporation
Built on Wed_Jul_22_19:09:09_PDT_2020
Cuda compilation tools,release 11.0,V11.0.221
Build cuda_11.0_bu.TC445_37.28845127_0

解决方法

我的期望有什么不对？

如果我们忽略托管内存和主机固定内存（即，如果我们专注于典型的主机内存，例如您在此处使用的内存），则CUDA中的一项基本原则是设备代码无法触摸/修改/访问主机内存（在Power9处理器平台上除外）。对此的直接扩展是，您（不能附带这些附带条件）不能传递对CUDA内核的引用，而希望对它进行任何有用的操作。

如果您真的想通过引用传递变量，则必须使用托管内存或主机固定的内存。这些需要特定的分配器，因此取决于指针的使用方式。

无论如何，除非您在Power9平台上，否则无法将对基于主机的堆栈内存的引用传递给CUDA内核并明智地使用它。

如果您想查看主机与设备之间的内存合理使用情况，请研究任何CUDA示例代码。

将引用/指针传递的主机变量设置到内核中的最佳/更好方法是什么？

我建议您与此处显示的内容最接近的内容是这样的（使用主机固定的分配器）：

$ cat t14.cu
#include <iostream>

__global__ void setHostVar(double *host_var) {
  double const var = 2.0;
  *host_var = var;
}

int main() {
  double *host_var_ptr;
  cudaHostAlloc(&host_var_ptr,sizeof(double),cudaHostAllocDefault);
  *host_var_ptr = 1.0;

  setHostVar<<<1,1>>>(host_var_ptr);
  cudaDeviceSynchronize();
  std::cout << "host_var = " << *host_var_ptr << std::endl;

  setHostVar<<<1,1>>>(host_var_ptr);
  cudaDeviceSynchronize();
  std::cout << "host_var = " << *host_var_ptr << std::endl;

  return 0;
}
$ nvcc -o t14 t14.cu
$ cuda-memcheck ./t14
========= CUDA-MEMCHECK
host_var = 2
host_var = 2
========= ERROR SUMMARY: 0 errors
$

尽管这可能并不完全符合您的要求。

您可能还对CUDA中如何使用异步感到困惑。在不试图涵盖该主题的各个方面的情况下，CUDA内核是异步启动的，这意味着CPU线程在继续操作之前不会等待CUDA内核完成。但是cudaDeviceSynchronize()强制在允许CPU线程继续执行之前，对该设备的所有先前发布的工作必须完成。其中包括内核以及与内核相关的任何事物，例如从内核/设备代码发出的数据复制（无论如何）。因此，我们希望在调用之后内核活动将是完整/一致的。

c++cudacuda