Macro expansion during JIT compilation?

[NamjaeChoi]

Hello,

I'm wondering if JIT compilation can properly expand macros. I guess it is somehow treating macros since CEED_QFUNCTION is also a macro, but I'm not sure if you have a special rule to specifically handle it or it is being handled by the compiler. The reason for this question is because we are defining QFunctions using macros like:

#define CEED_KERNEL(name, parent)                                                                  \
  class name final : public parent<name>                                                           \
  {                                                                                                \
  public:                                                                                          \
    static CEEDQFunctionPointer _qf;                                                               \
    auto getQFunction() { return _qf; }                                                            \
    std::string getQFunctionPath() { return std::string(__FILE__) + ":" + #name + "_qf"; }
#define CEED_KERNEL_END(name)                                                                      \
  }                                                                                                \
  ;                                                                                                \
                                                                                                   \
  CEED_QFUNCTION(name##_qf)                                                                        \
  (void * ctx, const CeedInt Q, const CeedScalar * const * in, CeedScalar * const * out)           \
  {                                                                                                \
    auto kernel = static_cast<name *>(ctx);                                                        \
    kernel->computeResidualInternal(Q, in, out);                                                   \
    return 0;                                                                                      \
  }                                                                                                \
                                                                                                   \
  CEEDQFunctionPointer name::_qf = name##_qf;

which is used in the user code like:

CEED_KERNEL(CEEDDiffusionKernel, CEEDDiffusion)
public:
static InputParameters validParams();

CEEDDiffusionKernel(const InputParameters & parameters);
CEED_KERNEL_END(CEEDDiffusionKernel)

But this leads to a JIT error that the QFunction was not defined:

/data/choin/moose/libCEED/backends/cuda/ceed-cuda-compile.cpp:125 in CeedCompile_Cuda(): NVRTC_ERROR_COMPILATION
default_program(63222): error: identifier "CEEDDiffusionKernel_qf" is undefined
      CEEDDiffusionKernel_qf(ctx, 1, inputs, outputs);

[jeremylt]

I'll try to make a running example of this to reproduce and understand what's happening, but it will take me time to do so.

[jeremylt]

If you aren't familiar, a minimal reproducible example greatly helps library maintainers in seeing the issue you are seeing.

[NamjaeChoi]

Let me try. This might be easy to make an example.

[jeremylt]

For CUDA JiT, the CEED_QFUNCTION macro is defined as

#define CEED_QFUNCTION(name) inline __device__ int name

I would guess you are getting a mixup with macro composition?

[jeremylt]

You can set the environment variable CEED_DEBUG=1 to see the code being JiTed, though that will potentially go away in the future as a result of #1695

[NamjaeChoi]

Hmm that just give me the code before expanding macro.

[jeremylt]

Ah, that's right. I'm still suspicious of chaining string manipulation macros. That's very touchy in my experience

[NamjaeChoi]

I just tried chaining macros in ex1-volume but seems to work fine. Let look into this more.

[NamjaeChoi]

Ah I think I know the reason. All our CEED source files are protected with a preprocessor, which is defined by our static compiler. But I guess JIT compiler doesn't know that? Is there a way to provide a preprocessor to the JIT compiler?

[jrwrigh]

The code itself is already in my fork. It's just that compiling and running MOOSE will be a pain for you.

Just being able to read the code is quite helpful, even if we don't attempt to compile it.

Edit: Just saw that Jeremy said more-or-less the same thing.

[jeremylt]

Alright, give the updated branch a try

[NamjaeChoi]

Works now. The next issue is now on our side. Thanks for the quick fix!

[lindsayad]

The particular branch @NamjaeChoi is working in is https://github.com/NamjaeChoi/moose/tree/libceed

[lindsayad]

We're probably going to put our ceed exploration on the backburner for a bit as we haven't figured out a good way yet to incorporate it into MOOSE's object oriented design