Better Enums with CUDA?

[akors]

Hi, I have been trying to use Better Enums with CUDA, in a .cu file.

When supplying --expt-relaxed-constexpr to NVCC, this works alright for the most part. However, when initializing the enum members, NVCC does something very strange.

Here is my testing code:

#include "enum.h"

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

int main() { }

Compile with nvcc constexpr-init.cu -std=c++11 -o constexpr-init

I get the following errors that are not understandable to me:

constexpr-init.cu:3:2362: error: lvalue required as left operand of assignment
 BETTER_ENUM(Channel, char, Red = 1, Green, Blue);
                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                          ^
constexpr-init.cu: In static member function 'static constexpr Channel::_optional Channel::_from_integral_nothrow(Channel::_integral)':
constexpr-init.cu:3:0: error: body of constexpr function 'static constexpr Channel::_optional Channel::_from_integral_nothrow(Channel::_integral)' not a return-statement
 BETTER_ENUM(Channel, char, Red = 1, Green, Blue);
 
constexpr-init.cu: In static member function 'static constexpr Channel::_optional Channel::_from_string_nothrow(const char*)':
constexpr-init.cu:3:0: error: body of constexpr function 'static constexpr Channel::_optional Channel::_from_string_nothrow(const char*)' not a return-statement
constexpr-init.cu: In static member function 'static constexpr Channel::_optional Channel::_from_string_nocase_nothrow(const char*)':
constexpr-init.cu:3:0: error: body of constexpr function 'static constexpr Channel::_optional Channel::_from_string_nocase_nothrow(const char*)' not a return-statement
constexpr-init.cu: In static member function 'static constexpr Channel::_value_iterable Channel::_values()':
constexpr-init.cu:3:0: error: body of constexpr function 'static constexpr Channel::_value_iterable Channel::_values()' not a return-statement

NVCC creates .cpp files that are passed to the host compiler. These files can be inspected by adding the --keep parameter to the command line.

If you do that and inspect the file, you will find the macro expansion somewhere, which looks like this:

namespace better_enums { namespace _data_Channel { }}class Channel { typedef better_enums::optional< Channel>  _optional; typedef better_enums::optional< unsigned long>  _optional_index; public: typedef char _integral; enum _enumerated: char { Red = 1, Green, Blue}; constexpr Channel(_enumerated value) : _value(value) { } constexpr operator _enumerated() const { return (_enumerated)(_value); } constexpr _integral _to_integral() const; static constexpr Channel _from_integral(_integral value); static constexpr Channel _from_integral_unchecked(_integral value); static constexpr _optional _from_integral_nothrow(_integral value); inline const char *_to_string() const; static constexpr Channel _from_string(const char * name); static constexpr _optional _from_string_nothrow(const char * name); static constexpr Channel _from_string_nocase(const char * name); static constexpr _optional _from_string_nocase_nothrow(const char * name); static constexpr bool _is_valid(_integral value); static constexpr bool _is_valid(const char * name); static constexpr bool _is_valid_nocase(const char * name); typedef better_enums::_Iterable< Channel>  _value_iterable; typedef better_enums::_Iterable< const char *>  _name_iterable; typedef better_enums::_Iterable< Channel> ::iterator _value_iterator; typedef better_enums::_Iterable< const char *> ::iterator _name_iterator; static constexpr const std::size_t _size_constant = (3); static constexpr std::size_t _size() { return _size_constant; } static constexpr const char *_name(); static constexpr _value_iterable _values(); static inline _name_iterable _names(); _integral _value; private: Channel() : _value((0)) { } constexpr explicit Channel(const _integral &value) : _value(value) { } static inline int initialize(); static constexpr _optional_index _from_value_loop(_integral value, std::size_t index = 0); static constexpr _optional_index _from_string_loop(const char * name, std::size_t index = 0); static constexpr _optional_index _from_string_nocase_loop(const char * name, std::size_t index = 0); friend struct better_enums::_initialize_at_program_start< Channel> ; }; namespace better_enums { namespace _data_Channel { static _initialize_at_program_start< Channel>  _force_initialization; enum _PutNamesInThisScopeAlso { Red = 1, Green, Blue}; constexpr const Channel _value_array[] = {(Channel::Red = (1)), (Channel::Green), (Channel::Blue)}; constexpr const char *_the_raw_names[] = {("Red = 1"), ("Green"), ("Blue")}; constexpr const char *const *_raw_names() { return _the_raw_names; } inline char *_name_storage() { static char storage[] = "Red = 1,Green,Blue,"; return storage; } inline const char **_name_array() { static const char *value[Channel::_size_constant]; return value; } inline bool &_initialized() { static bool value = false; return value; } }}constexpr const Channel operator+(Channel::_enumerated enumerated) { return static_cast< Channel>(enumerated); } constexpr Channel::_optional_index Channel::_from_value_loop(_integral value, std::size_t index) { return ((index == _size()) ? _optional_index() : ((((((better_enums::_data_Channel::_value_array)[index])._value) == value) ? ((_optional_index)(index)) : (_from_value_loop(value, index + (1)))))); } constexpr Channel::_optional_index Channel::_from_string_loop(const char *name, std::size_t index) { return ((index == _size()) ? _optional_index() : ((::better_enums::_names_match(better_enums::_data_Channel::_raw_names()[index], name) ? ((_optional_index)(index)) : (_from_string_loop(name, index + (1)))))); } constexpr Channel::_optional_index Channel::_from_string_nocase_loop(const char *name, std::size_t index) { return ((index == _size()) ? _optional_index() : ((::better_enums::_names_match_nocase(better_enums::_data_Channel::_raw_names()[index], name) ? ((_optional_index)(index)) : (_from_string_nocase_loop(name, index + (1)))))); } constexpr Channel::_integral Channel::_to_integral() const { return (_integral)(_value); } constexpr Channel Channel::_from_integral_unchecked(_integral value) { return static_cast< _enumerated>(value); } constexpr Channel::_optional Channel::_from_integral_nothrow(_integral value) { return ::better_enums::_map_index< Channel> (better_enums::_data_Channel::_value_array, _from_value_loop(value)); } constexpr Channel Channel::_from_integral(_integral value) { return ::better_enums::_or_throw(_from_integral_nothrow(value), "Channel::_from_integral: invalid argument"); } inline const char *Channel::_to_string() const { return ::better_enums::_or_null(::better_enums::_map_index< const char *> (better_enums::_data_Channel::_name_array(), _from_value_loop(::better_enums::continue_with(initialize(), _value)))); } constexpr Channel::_optional Channel::_from_string_nothrow(const char *name) { return ::better_enums::_map_index< Channel> (better_enums::_data_Channel::_value_array, _from_string_loop(name)); } constexpr Channel Channel::_from_string(const char *name) { return ::better_enums::_or_throw(_from_string_nothrow(name), "Channel::_from_string: invalid argument"); } constexpr Channel::_optional Channel::_from_string_nocase_nothrow(const char *name) { return ::better_enums::_map_index< Channel> (better_enums::_data_Channel::_value_array, _from_string_nocase_loop(name)); } constexpr Channel Channel::_from_string_nocase(const char *name) { return ::better_enums::_or_throw(_from_string_nocase_nothrow(name), "Channel::_from_string_nocase: invalid argument"); } constexpr bool Channel::_is_valid(_integral value) { return _from_value_loop(value); } constexpr bool Channel::_is_valid(const char *name) { return _from_string_loop(name); } constexpr bool Channel::_is_valid_nocase(const char *name) { return _from_string_nocase_loop(name); } constexpr const char *Channel::_name() { return "Channel"; } constexpr Channel::_value_iterable Channel::_values() { return _value_iterable(better_enums::_data_Channel::_value_array, _size()); } inline Channel::_name_iterable Channel::_names() { return _name_iterable(better_enums::_data_Channel::_name_array(), ::better_enums::continue_with(initialize(), _size())); } inline int Channel::initialize() { if (better_enums::_data_Channel::_initialized()) { return 0; }  ::better_enums::_trim_names(better_enums::_data_Channel::_raw_names(), better_enums::_data_Channel::_name_array(), better_enums::_data_Channel::_name_storage(), _size()); better_enums::_data_Channel::_initialized() = true; return 0; } constexpr bool operator==(const Channel &a, const Channel &b) { return (a._to_integral()) == (b._to_integral()); } constexpr bool operator!=(const Channel &a, const Channel &b) { return (a._to_integral()) != (b._to_integral()); } constexpr bool operator<(const Channel &a, const Channel &b) { return (a._to_integral()) < (b._to_integral()); } constexpr bool operator<=(const Channel &a, const Channel &b) { return (a._to_integral()) <= (b._to_integral()); } constexpr bool operator>(const Channel &a, const Channel &b) { return (a._to_integral()) > (b._to_integral()); } constexpr bool operator>=(const Channel &a, const Channel &b) { return (a._to_integral()) >= (b._to_integral()); }

(Sorry for the no spaces, but this is what I get out).
The error occurs at character 2362 which seems to be the expression Channel::Red = (1) in the following statement:

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

This is the end of my investigations, I don't know where to go from here.

I know that it's probably the NVCC that's broken, but can you think of a way to make this initialization work?

[aantron]

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.

[aantron]

That specific code is the instantiation of this macro:

better-enums/enum.h

Lines 494 to 495 in 2fad3f6

	#define BETTER_ENUMS_EAT_ASSIGN_SINGLE(EnumType, index, expression) \
	((::better_enums::_eat_assign<EnumType>)EnumType::expression),

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.

[akors]

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?

[akors]

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.

[aantron]

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:

better-enums/enum.h

Lines 494 to 500 in 2fad3f6

	#define BETTER_ENUMS_EAT_ASSIGN_SINGLE(EnumType, index, expression) \
	((::better_enums::_eat_assign<EnumType>)EnumType::expression),
	#define BETTER_ENUMS_EAT_ASSIGN(EnumType, ...) \
	BETTER_ENUMS_ID( \
	BETTER_ENUMS_PP_MAP( \
	BETTER_ENUMS_EAT_ASSIGN_SINGLE, EnumType, __VA_ARGS__))

and invoked here:

better-enums/enum.h

Lines 678 to 679 in 2fad3f6

	BETTER_ENUMS_CONSTEXPR_ const Enum _value_array[] = \
	{ BETTER_ENUMS_ID(BETTER_ENUMS_EAT_ASSIGN(Enum, __VA_ARGS__)) }; \

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.

[aantron]

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.

[aantron]

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.

[aantron]

...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.