OpenCL Add a DECLSPEC or INLINE macro

[magnumripper]

We've had this in opencl_misc.h for ages:

/*
 * Some runtimes/drivers breaks on using inline, others breaks on lack of it,
 * yet others require use of static as well.
 */
#if __MESA__
#define inline	// empty!
#elif __POCL__
// Do nothing (POCL complains if we redefine)
#elif gpu_amd(DEVICE_INFO) // We really target ROCM here
#define inline	static inline
#else
// Do nothing
#endif

This is suboptimal because 1) it's not obvious in affected functions that inline may be changed, and 2) it would change static inline to static static inline. Also, the whole thing might be a long gone problem - we do have code that use neither or just static and there's no reports of it breaking anything. On the other hand, just maybe #5456 is affected by what syntax (or none) is used - I recall a few formats do not emit those warnings, while most do.

FWIW hashcat has this:

/**
 * function declarations can have a large influence depending on the opencl runtime
 * fast but pure kernels on rocm is a good example
 */

#ifdef NO_INLINE
#define HC_INLINE
#else
#define HC_INLINE inline static
#endif

#if defined IS_AMD && defined IS_GPU
#define DECLSPEC HC_INLINE
#elif defined IS_HIP
#define DECLSPEC __device__ HC_INLINE
#else
#define DECLSPEC
#endif

...then all their functions start with DECLSPEC unless they're already ifdef'ed for a specific platform.

[magnumripper]

it's not obvious in affected functions that inline may be changed

I guess the very least we should do about that is rename our inline macro to MAYBE_INLINE like in host code. I like the DECLSPEC though.

@solardiz I'm not touching this until I hear your opinion, as I suspect you may have one. Except I'll investigate at some point in time whether static/inline/both/none (or the case of not using any functions) is affecting #5456.

[magnumripper]

FWIW the "correct" way of writing it is either static inline or nothing at all (and most runtimes would inline anyway). No underscores or attribute should be needed.

For NOT inlining, __attribute__((noinline)) might be the best bet but I would expect it to be ignored more often than not.

[magnumripper]

I think I'll just go with this:

#define INLINE      static inline
#define NOINLINE    __attribute__((noinline))

[magnumripper]

I'm also considering dropping most uses of inline, relying on it mostly happening anyway, but using the INLINE macro for things like our LUT3 function and stuff like that. Then benchmark all formats before and after, comparing ptxas output and binary size and see if/what changes. Imported subdir code such as ed25519 should mostly be kept as-is.

[magnumripper]

Then benchmark all formats before and after, comparing ptxas output and binary size and see if/what changes.

Tried this (with nvidia). Nothing changed, ptxas output identical.

[magnumripper]

Next is to experiment with what affects #5456 and what doesn't. This has to wait a week or so, I'm stuck abroad with a laptop.

[solardiz]

@solardiz I'm not touching this until I hear your opinion, as I suspect you may have one.

Just like you, I'm unhappy we were redefining inline, so let's move away from that, like you do in your PR now.

As to the rest, I'm not sure. You experimented with it just recently and continue to do so. So I'll defer to you on it.

[magnumripper]

I'm merging the nearly cosmetical #5638 for now. #5456 is more interesting and being able to spot where the macro is used helps a lot with that.