7

Switching from CUDA 8.0 to CUDA 9.0 RC, I get a warning about:

__host__ __device__ ~Foo() = default;

The warning is:

path/to/Foo.cuh(69): warning: __host__ annotation on a defaulted function("~Foo") is ignored

which I didn't use to get before. Should I really be getting this warning? What's wrong with indicating you want the default destructor on both the device and the host side?

einpoklum
  • 118,144
  • 57
  • 340
  • 684
  • I think you probably should get a warning. Defaulting a constructor or destructor is telling the compiler to generate its own trivial default implementation automagically. Adding an annotation is irrelevant in this case. Both compilers will generate a default, specifying that the default from either host or device compiler should exist on both host and device is wrong in this case. – talonmies Sep 28 '17 at 15:44
  • @talonmies: So, you're saying that even if, without this line, no default destructor is generated on the host side nor on the device side, I should still have it without any annotation? – einpoklum Sep 28 '17 at 15:58
  • 1
    NVIDIA claim that the device toolchain supports [N2346](http://www.open-std.org/jtc1/sc22/wg21/docs/papers/2007/n2346.htm#trivial) . If you want that behaviour (and actually understand what it entails), then by all means use defaulted constructors or destructors. But in that case specifying `__host__ __device__` for a defaulted function doesn't make sense to me, and I think the warning is valid. But what do I know.... – talonmies Sep 28 '17 at 16:08
  • @talonmies: What I want is to use the default-generated dtor despite it not being generated implicitly. I _think_ that's what N2346 is about though I'm not entirely sure. I could just avoid it with an empty destructor, but there's other code with the same issue, like equality operators etc. – einpoklum Sep 28 '17 at 16:12
  • Be aware that an empty destructor and a default destructor are not the same thing. This is really a language lawyer question and I am definitely not one of those, so I'm not even going to try and answer this. – talonmies Sep 28 '17 at 16:14
  • @talonmies The same warning occurs for defaulted constructors and assignment operators. If you remove `__host__ __device__` , how would the nvcc know that objects of a class/struct will be created in device memory? – Matthias Oct 06 '17 at 17:30
  • 1
    @Matthias: Very late follow up, but remember nvcc isn't a compiler. It runs all code through two parallel compiler passes (host and device), and each compiler will emit its own default implementation. The warning comes for exactly this reason -- the device code compiler sees the `__host__` decorator applied to a device default and warns it is irrelevant *at that point in the compilation trajectory*. – talonmies Aug 09 '18 at 08:32

1 Answers1

7

What's wrong with indicating you want the default destructor on both the device and the host side?

But that isn't what the code says. It says you want the same trivial default destructor in both host and device, and that is why there is a warning, because neither compiler (host and device) will potentially emit the same default destructor (and because the the way the compilation trajectory works that can't happen).

NVIDIA claim that the recent device toolchains support N2346. If you want that behaviour (and actually understand what it entails), then the correct code should just be:

~Foo() = default;

Both compilers will automagically emit their own default destructors in both codes and everything will just work.

If you want a hacky workaround for an existing code base, adding 

-Xcudafe="--diag_suppress=2886"

to your nvcc build commands should eliminate the warning, although I counsel against suppressing warnings.

[Answer added as a summary of comments discussion to get question off the unanswered list for the CUDA tag. ]

talonmies
  • 70,661
  • 34
  • 192
  • 269
  • What would be the portable solution for "host-device" explicitly defaulted functions in nvcc as well as native [clang++](https://www.llvm.org/docs/CompileCudaWithLLVM.html)? Is the warning 2886 related only to this problem alone or actually more things? For the record, the warning number seems to have changed to 2929 in CUDA 10 (or 10.1)... – Jakub Klinkovský Aug 24 '19 at 12:24
  • Do you claim that `~Foo() = default` implies that the destructor is a no-op? This is not a case in C++: defaulted destructor calls members' destructors. Ok, I've read your comments on the questing, and it seems like you know that. Still, the answer is a little misleading – Nikita Petrenko Sep 07 '19 at 17:23
  • "Both compilers will automagically emit their own default destructors in both codes and everything will just work." <- Are you referring to future versions of CUDA, or the current one (10.x)? – einpoklum Dec 29 '19 at 20:29
  • I was referring to CUDA 9 which was current when this answer was written – talonmies Dec 29 '19 at 20:33