[Clang][CUDA][HIP] lambda capture of constexpr variable inconsistent between host and device

[mkuron]

Consider the following bit of HIP code:

#include <algorithm>

using std::max;

template<typename F>
static void __global__
kernel(F f)
{
  f(1);
}

void test(float const * fl, float const * A, float * Vf)
{
  float constexpr small(1.0e-25);

  auto f = [=] __device__ __host__ (unsigned int n) {
    float const value = max(small, fl[0]);
    Vf[0] = value * A[0];
  };
  static_assert(sizeof(f) == sizeof(fl) + sizeof(A) + sizeof(Vf));
  kernel<<<1,1>>>(f);
}

The static_assert fails in the host-side compilation but succeeds in the device-side compilation. This means that the layout of the struct synthesized from the lambda is inconsistent between host and device, so if you use any of the captured variables on the device side, they will contain the data of some of the other variables. You can also use -Xclang -fdump-record-layouts to see that. Evidently the constexpr variable is part of the captured variables only on the host side, but not on the device side.
With --cuda-host-only:

*** Dumping AST Record Layout
         0 | class (lambda at <source>:23:12)
         0 |   const float * 
         8 |   const float 
        16 |   float * 
        24 |   const float * 
           | [sizeof=32, dsize=32, align=8,
           |  nvsize=32, nvalign=8]

With --cuda-device-only:

*** Dumping AST Record Layout
         0 | class (lambda at <source>:23:12)
         0 |   const float * 
         8 |   float * 
        16 |   const float * 
           | [sizeof=24, dsize=24, align=8,
           |  nvsize=24, nvalign=8]

Godbolt: https://cuda.godbolt.org/z/KE789sevs.

When you compile the exact same code for CUDA, this does not happen. However, if you add

template <typename T, std::enable_if_t<std::is_arithmetic<T>::value, int> = 0>
__host__ T max(const T a, const T b) {
    return std::max(a, b);
}

after line 3 of the code at the top, you get the exact same layout discrepancy as with HIP. See https://cuda.godbolt.org/z/e3Ybr4hK1.

I can replace the [=] with [fl, A, Vf] and if that __host__ T max overload is present, it tells me that variable 'small' cannot be implicitly captured in a lambda with no capture-default specified, but if I leave out that overload it does not show that error message.

The example code uses std::max (from the GNU libstdc++), but I have no doubt that the same issue can easily be demonstrated without making use of any library headers.

nvcc, for comparison, does not exhibit this issue (https://cuda.godbolt.org/z/33MvMc755). It has a consistent capture size between host and device. Judging by the size, the constexpr variable seems to be part of the capture, though in contrast to Clang, a lambda appears to always occupy 8 bytes more than required for its captured variables.

[Artem-B]

Before we dive deeper, I'm not quite sure that this assertion is valid in principle: static_assert(sizeof(f) == sizeof(fl) + sizeof(A) + sizeof(Vf));

Depending on the types, target-specific alignment requirement and the order of the fields it's not even guaranteed to work with POD types, never mind something as implementation-specific as lambdas.

A more robust comparison would be to assert on offsetof() of captured fields. If we have a discrepancy there, that would indeed be a problem.

[mkuron]

I'm not quite sure that this assertion is valid in principle: static_assert(sizeof(f) == sizeof(fl) + sizeof(A) + sizeof(Vf));

Yeah, it's definitely not valid in principle. I just used this line to demonstrate that for my sample it is valid on the device but not on the host, which implies that the struct size is inconsistent. To get the complete picture, use -Xclang -fdump-record-layouts (see above for output of that), which for my sample shows that there is an extra member in the middle, so the offsets are actually different.

[Artem-B]

In a way your CUDA godbolt example (https://cuda.godbolt.org/z/e3Ybr4hK1) commits an exotic variant of ODR violation -- the lambda with the same signature has different implementations on the host and on the device. If I let both host and device use the same max, the discrepancy disappears.

Another aspect is that [=] leaves it up to the compiler what it wants or needs to capture, and that may be different for the host and GPU compilations. In this particular case, because the host and device lambdas call different max functions, compiler makes different decisions on what needs to be captured. It's very fragile. Making small static or moving it out of test affects whether it gets captured.

In case of HIP reproducer, the ODR violation is more subtle, but it's still there: https://cuda.godbolt.org/z/64sWd6rYj
You can see that the max function it calls on the host and the device are different.

I think the bottom line here is:

• keep lambdas the same on both host and device. Otherwise all bets are off.
• don't use implicit capture, as that may lead to different lambdas on host and GPU, too.

[mkuron]

an exotic variant of ODR violation

That's also the conclusion I reached. Unfortunately, it's not easy to see that there is an ODR violation in my sample code -- there are no preprocessor directives involved, which is the usual way to shoot yourself in the foot with host vs. device discrepancies. In this case, I guess it comes from __device__/__host__ function overloads hidden somewhere in the HIP headers or in __clang_hip_math.h.

Here is a completely standalone sample that uses a subtle difference in function signatures to cause the ODR violation: https://cuda.godbolt.org/z/dPPoqKEW9. nvcc actually gives you a useful error message here and tells you that more than one overloaded function matches the argument list (because unlike Clang, it does not support overloading based on __device__/__host__).

In this particular case, because the host and device lambdas call different max functions, compiler makes different decisions on what needs to be captured. It's very fragile. Making small static or moving it out of test affects whether it gets captured.

When it's static or outside the function, the capture is actually consistent between host and device. My inconsistency only appears to occur when it is constexpr but not static. Seems like the constexpr variable is only captured when the function takes its arguments by value.

don't use implicit capture

I agree, though that can sometimes be impractical. The unfortunate thing is that we have code here that works fine with nvcc but not with HIP (or in the case of https://cuda.godbolt.org/z/dPPoqKEW9, code that nvcc rejects with a useful error message while HIP miscompiles it), and the compiler can't even warn you about it because the issue cannot be detected without actually comparing the device and host compile results with each other.

[Artem-B]

I think, it's an unfixable problem. CUDA/HIP tries to have its cake (separate host/device compilations) and eat it (transparently exchange data across host/GPU boundary). In a plan C++ world it's an equivalent of trying to pass around capturing-lambda-typed objects across TUs, which is an almost canonical way to create an ODR violation. Except that for CUDA/HIP the definition of TU becomes murky. On one hand we are compiling the same file, but the specific sub-compilations of that file are not exactly the same, and as far as C++ machinery is concerned they are separate TUs, with separate lambda types, etc.

Diagnosing it will also be a problem, because host-side compilation has no idea what happens on the GPU side and vice versa, so we can't tell if we end up with a different idea of what the lambda looks like.

We could warn on passing across host/GPU boundary lambdas with non-explicid capture list, but it would only catch some cases, and will create a lot of false positives where [=] does not create a problem.

@yxsamliu Do you have any thought on this?

[yxsamliu]

The issue happened due to that a constexpr variable is captured only on one side in the [=] case.

If the constexpr variable is ODR-used, e.g. passed to a function by reference, then it is captured. If it is passed by value, then it is not captured, since only the literal is used.

In this case, HIP host and device max() functions have different signatures. On one side, it has reference parameters. On the other side, it has value parameters. That's why it is captured on one side only.

We could try making the max()/min() signatures consistent, which may alleviate the issue.

In general, I agree with Artem that the ODR violation on host/device sides is difficult to diagnose since we usually do not have the AST on the other side to compare with, or it is too expensive to build AST for the other side.

However, for this specific case, I may come up with some warning. Basically when we do overloading resolution in a lambda function, if we have a candidate that is not viable only due to wrong-side, and there is a constexpr variable that is passed differently by ref to the chosen candidate and by value to the wrong-sided candidate, and this variable is captured through [=], then we emit a warning that this variable may not be captured on the other side.

[mkuron]

then we emit a warning that this variable may not be captured on the other side.

Thanks, that sounds useful. Do you think that, instead of emitting a warning, you could automatically add the variable to the set of captured variables?

[yxsamliu]

then we emit a warning that this variable may not be captured on the other side.

Thanks, that sounds useful. Do you think that, instead of emitting a warning, you could automatically add the variable to the set of captured variables?

I will give it a try.