Better Enums with CUDA? about better-enums HOT 8 CLOSED

I suspect something is broken with macro expansion in NVCC, but I am not sure yet. I'll point to the most likely offending line.

The NVCC output you've pasted includes this:

constexpr const Channel _value_array[] = {(Channel::Red = (1)), (Channel::Green), (Channel::Blue)};

This is definitely wrong, looking at the assignment of (1). It's supposed to be this:

const Channel _value_array[] = { ((::better_enums::_eat_assign<Channel>)Channel::Red = 1), ((::better_enums::_eat_assign<Channel>)Channel::Green), ((::better_enums::_eat_assign<Channel>)Channel::Blue), };

Clang++ generates that on my system, with c++ -E foo.cc.

from better-enums.

That specific code is the instantiation of this macro:

Perhaps NVCC doesn't like the nested parentheses?

Also, don't mind the missing constexpr in the "correct" output, that is because I forgot to pass -std=c++11 or higher.

from better-enums.

Turns out that it actually works, when I define BETTER_ENUMS_NO_CONSTEXPR before including it.

It would still be nice to use BETTER_ENUM with constexpr, I believe NVCC should theoretically support it.

I'm pretty sure that this is actually an NVCC problem, but I would request your help in finding the cause of this issue so we can report this to NVIDIA.

So here is a document that goes through the NVCC compilation phases, page 22 shows a nice overview over the process.

You can watch what NVCC is doing by adding the --verbose flag to compilation.

Both of the C++ intermediate files constexpr-init.cpp1.ii and constexpr-init.cpp4.ii that are created by using gcc -E contain the proper code:

constexpr const Channel _value_array[] = {
    ((::better_enums::_eat_assign<Channel>) Channel::Red = 1),
    ((::better_enums::_eat_assign<Channel>) Channel::Green),
    ((::better_enums::_eat_assign<Channel>) Channel::Blue), };

Then however, the output file of cudafe++ produces garbage:

constexpr const Channel _value_array[] = {(Channel::Red = (1)), (Channel::Green), (Channel::Blue)};

If I understand correctly, cudafe++ does some sort of template parsing, which I gather from the fact that NVCC runs it with the parameter --parse_templates. Is it possible that something goes wrong here with your _eat_assign template?

What is it actually supposed to do?

from better-enums.

Turns out it doesn't work. With the following example:

#include <iostream>

#define BETTER_ENUMS_NO_CONSTEXPR
#include "enum.h"

BETTER_ENUM(Channel, char, Red /* = 1*/, Green, Blue);

__device__ void calcSomething(int* result, int value, Channel mode) {
    switch(mode) {
    case Channel::Blue:
        *result = 2*value;
        break;
    default:
        *result = 10*value;
    }
}

int main() { }

I get the following error:

constexpr-init.cu(10): error: calling a __host__ function("Channel::operator  ::Channel::_enumerated const") from a __device__ function("calcSomething") is not allowed

constexpr-init.cu(10): error: identifier "Channel::operator  ::Channel::_enumerated const" is undefined in device code

Without the #define BETTER_ENUMS_NO_CONSTEXPR the code compiles, however it breaks again when I initialize Red = 1.

As it looks now, BETTER_ENUMS is simply not compatible with NVCC.

from better-enums.

It's almost certainly the case that NVCC is broken.

The point of _eat_assign is to convert the = 1 syntax into something that is acceptable inside an array initializer. Normally, you can't have an array initializer {A = 1, B = 2}. The way Better Enums turns that into valid syntax, is by prepending a cast to A and B that turns them into objects that have an assignment operator. Then, the (Cast)A = 1, (Cast)B = 2 become valid expressions, that are evaluated as part of the initializer. As a last step, the casted A and B have to convert back to the type of the array element. _eat_assign does all that (it "eats" the assignments, maybe I should rename it :p).

I don't have the necessary expertise to look into NVCC, unfortunately. But you should be able to copy out the _eat_assign template, make minor modifications to it if needed, and use it to narrow down the bug. It could also be an interaction between _eat_assign and the BETTER_ENUMS_PP_MAP macro. The job of that macro is to apply another macro to each one of its arguments. Maybe tokens are being somehow lost there by NVCC – it certainly would explain some of the output in your second-to-last post. Try copying that macro out as well, and seeing if NVCC is able to process it properly.

The macro is defined here:

and invoked here:

You should be able to extract all this to a file, for reproducing, pretty easily.

Working backwards, to mimic the last bit of code, that creates the array, create a static array somewhere, with the same definition as _value_array above. I don't know if you will need to nest it in a macro or a namespace to trigger the bug in NVCC, but if so, wrap it:

#define BETTER_ENUMS_TYPE_REPRO(Enum, ...) \
namespace better_enums_data_ ## Enum { \
BETTER_ENUMS_CONSTEXPR_ const Enum      _value_array[] =                       \
    { BETTER_ENUMS_ID(BETTER_ENUMS_EAT_ASSIGN(Enum, __VA_ARGS__)) };           \
}

As you can see, you now need to extract only:

• BETTER_ENUMS_CONSTEXPR_: just define it as constexpr.

• BETTER_ENUMS_ID: included in BETTER_ENUMS_PP_MAP.

• BETTER_ENUMS_EAT_ASSIGN: copy it from above. That will also pull in the BETTER_ENUMS_PP_MAP macro. Its definition starts here (but it's quite large, because it is a manual expansion of a loop):

  better-enums/enum.h

  Line 103 in 2fad3f6

  #define BETTER_ENUMS_PP_MAP(macro, data, ...) \

• The eat_assign template itself:

  better-enums/enum.h

  Lines 359 to 374 in 2fad3f6

  	// Values array declaration helper.
  	template <typename EnumType>
  	struct _eat_assign {
  	explicit BETTER_ENUMS_CONSTEXPR_ _eat_assign(EnumType value) : _value(value)
  	{ }
  	template <typename Any>
  	BETTER_ENUMS_CONSTEXPR_ const _eat_assign&
  	operator =(Any) const { return *this; }
  	BETTER_ENUMS_CONSTEXPR_ operator EnumType () const { return _value; }
  	private:
  	EnumType _value;
  	};

  Which has no further dependencies beyond what you will already have extracted.

That should be self-contained, and enough to find the bug by applying BETTER_ENUMS_TYPE_REPRO(Channel, Red = 1, Green, Blue), and checking what tokens are in the output.

If you want to try building up to the bug instead, try following the explanation of how a simplified Better Enums works in this article: https://stackoverflow.com/questions/28828957/enum-to-string-in-modern-c11-c14-and-future-c17-c20/31362042#31362042.

from better-enums.

Ah, maybe the macro is not involved, as I see the expansion was correct at an earlier step of the NVCC process. But the above steps should get you a relatively simple, self-contained file to work with. Also, I don't know if NVCC is starting over with macros at some point or not.

from better-enums.

This error seems reasonable with BETTER_ENUMS_NO_CONSTEXPR:

constexpr-init.cu(10): error: calling a __host__ function("Channel::operator  ::Channel::_enumerated const") from a __device__ function("calcSomething") is not allowed

constexpr-init.cu(10): error: identifier "Channel::operator  ::Channel::_enumerated const" is undefined in device code

without knowing too much about the semantics of CUDA, I don't know if it makes sense to try adding __device__ in places inside the macro definition.

As for the error with constexpr enabled and Red = 1, it is because of the missing cast in that output you showed, and Channel::Red not being an assignable expression without that cast, as described in my previous comment.

from better-enums.

...and to give yet more detail about the error you are seeing when BETTER_ENUMS_NO_CONSTEXPR is defined:

Better Enums are objects, and to each Better Enums type Foo, corresponds an actual C++ enum or enum class Foo::_enumerated. So that Better Enums can be used in switch statements, with exhaustiveness checking by the compiler, they include an implicit conversion from Foo to Foo::_enumerated. Writing switch(e) where e is a Better Enum causes the compiler to insert this conversion, and it seems NVCC is rightly complaining that when the cast is not constexpr, it looks like an ordinary host function.

from better-enums.