Abstract the use of constexpr aggregate host/device constants

[fwyzard]

The use of host-side consexpr constants in device code is limited to

• scalars (other than long double)
• scalar elements of aggregates used inside constexpr functions

In particular, it's not possible to use constexpr scalars by pointer or reference (e.g. std::min() takes arguments by reference), or pass constexpr arrays as pointers, or access elements of constexpr arrays outside of constexpr functions.

The workaround that we found, is to write portable code that declare two copies of the constants, one on the host and one on the device:

#ifdef __CUDA_ARCH__
__device__
#endif
constexpr uint32_t values[N] = { ... };

We should probably abstract it with something like

#ifdef __CUDA_ARCH__
#define HOST_DEVICE_CONSTANT __device__ constexpr
#else
#define HOST_DEVICE_CONSTANT constexpr
#endif

to be used as

HOST_DEVICE_CONSTANT uint32_t values[N] = { ... };

Suggestions for a better name are welcome :-)

[cmsbuild]

A new Issue was created by @fwyzard Andrea Bocci.

@Dr15Jones, @perrotta, @dpiparo, @makortel, @smuzaffar, @qliphy can you please review it and eventually sign/assign? Thanks.

cms-bot commands are listed here

[fwyzard]

assign heterogeneous

[cmsbuild]

New categories assigned: heterogeneous

@fwyzard,@makortel you have been requested to review this Pull request/Issue and eventually sign? Thanks

[fwyzard]

Duplicate of #35370

[fwyzard]

Implemented in https://github.com/cms-sw/cmssw/pull/37159

[fwyzard]

+heterogeneous

[cmsbuild]

This issue is fully signed and ready to be closed.

[fwyzard]

-heterogeneous

[thomreis]

The method mentioned above does not work if the non-scalar element is a static member of a class like, e.g. in https://github.com/cms-sw/cmssw/blob/master/DataFormats/EcalDigi/interface/EcalConstants.h#L9

Replacing constexpr with HOST_DEVICE_CONSTANT as defined in #37159 leads to the compile error EcalConstants.h(15): error: memory qualifier on data member is not allowed

I have also tried to apply it to the whole class with #define HOST_DEVICE_CONSTANT __device__ class and then use it as HOST_DEVICE_CONSTANT ecalPh2 {... but that does not work neither. The compiler generates a warning EcalConstants.h(11): warning #1866-D: attribute does not apply to any entity and the identifier remains undefined on the device.

[VinInn]

@thomreis a possible solution is to declare/define the constexpr in a namespace instead of the class scope (and eventually "use" the namespace in the methods of that class)

[fwyzard]

@VinInn yes, definitely. The downside is that a class can be passed as a template argument to another type (e.g. to implement Phase-1 vs Phase-2 parameters), while a namespace cannot :-(

[thomreis]

The ability to pass the class as a template is exactly why it was set up like this, yes. The Phase 2 ECAL reconstruction we are currently working on relies on it.

[VinInn]

what about

namespace one {
   constexpr int n = 2;
   __device__ constexpr float g[n] = {1.,2.};
}

namespace two {
   constexpr int n = 4;
   __device__ constexpr float g[n] = {1.,2.,3.,4.};
}

struct One {
   static constexpr int n = one::n;
   static constexpr float const * g = one::g;

};

struct Two {
   static constexpr int n = two::n;
   static constexpr float const * g = two::g;
};

[thomreis]

That compiles, yes, and runs as well.

[fwyzard]

With

namespace one {
   __device__ constexpr float g[n] = {1.,2.};
}
struct One {
   static constexpr float const * g = one::g;
};

is the access to One::g[0] still constexpr ?

[thomreis]

Yes it seems as if it is.

[fwyzard]

To do: document the use of the HOST_DEVICE_CONSTANT macro and its limitations.