从 __host__ __device__ 函数调用 __host__ 函数

Question

编译 MWE 时

#include <iostream>
#include "cuda.h"

struct Foo{
///*
    Foo( ){
      std::cout << "Construct" << std::endl;
    }

    Foo( const Foo & that ){
      std::cout << "Copy construct" << std::endl;
    }
//*/
   __host__ __device__
   int bar( ) const {
     return 0;
   }
};

template<typename CopyBody>
__global__ 
void kernel( CopyBody cBody ){
  cBody( );
}

template <typename CopyBody>
void wrapper( CopyBody && cBody ){
  std::cout << "enquing kernel" << std::endl;
  kernel<<<1,32>>>( cBody );
  std::cout << "kernel enqued" << std::endl;
}

int main(int argc, char** argv) {

  Foo foo;

  std::cout << "enquing kernel" << std::endl;
  kernel<<<1,32>>>( [=] __device__ ( ) { foo.bar( ); } );
  std::cout << "kernel enqued" << std::endl;
  cudaDeviceSynchronize( );

  wrapper( [=] __device__ ( ) { foo.bar( ); } );
  cudaDeviceSynchronize( );
  
  return 0;
}

对于 CUDA 10.1 (nvcc --expt-extended-lambda test.cu -o test)，编译器会针对 test.cu(16): warning: calling a __host__ function("Foo::Foo") from a __host__ __device__ function("") is not allowed 发出警告。但是，永远不会在设备上调用复制构造函数。 CUDA 9.1 不会产生此警告。

• 直接调用kernel（不产生警告）和wrapper版本有什么区别？
• 忽略此警告是否安全？
• 将#pragma hd_warning_disable 或#pragma nv_exec_check_disable 放在哪里可以摆脱它？

给定的 MWE 是基于一个更大的项目，其中wrapper 决定是使用__device__ 还是__host__ lambda。构造函数/析构函数不能标记为__host__ __device__，因为它们只需要在 CPU 上调用（（取消）分配 CUDA 内存） - 这或删除构造函数/析构函数（并让编译器创建默认的 __host__ 和 @ 987654333@ 版本）否则会有所帮助。

Answer 1

通过以下修改，我没有得到提及的警告：（我在 Windows 10 上使用了 CUDA 10.1）

#include <stdio.h>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

struct Baz {

   Baz() {
      printf("%s: Construct\n", __FUNCTION__);
   }

   Baz(const Baz & that) {
      printf("%s: Copy Construct\n", __FUNCTION__);
   }

};

struct Foo: 
   public Baz {

   __host__ __device__ 
   int bar() const  {
      return 0;
   } 
};

template<typename CopyBody>
__global__
void kernel(CopyBody cBody) {
   cBody();
}

template <typename CopyBody>
void wrapper(CopyBody && cBody) {
   printf("%s: enquing kernel\n",__FUNCTION__);
   kernel << <1, 32 >> > (cBody);
   printf("%s: kernel enqued\n", __FUNCTION__);
}

int main(int argc, char** argv) {

   Foo foo;

   printf("%s: enquing kernel\n", __FUNCTION__);   
   kernel << <1, 32 >> > ([=] __device__() { foo.bar(); });
   printf("%s: kernel enqued\n", __FUNCTION__);   
   cudaDeviceSynchronize();

   wrapper([=] __device__() { foo.bar(); });
   cudaDeviceSynchronize();

   return 0;
}

以上代码产生以下输出：

Foo::Foo: Construct
main: enquing kernel
Foo::Foo: Copy Construct
Foo::Foo: Copy Construct
main: kernel enqued
Foo::Foo: Copy Construct
Foo::Foo: Copy Construct
wrapper: enquing kernel
Foo::Foo: Copy Construct
wrapper: kernel enqued

为方便起见，我将<iostream> 替换为<stdio.h>。 printf() 在内核中工作。