cuda - 从 __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 内存) - 否则会有所帮助。

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()从内核工作。