Error macros now have GPU variant.

[kjetilly]

This will allow us to keep code such as

OPM_THROW(std::exception, "Message");

also in the GPU kernels. The alternative would be to wrap every such throw-statement in the following manner:

#if OPM_IS_INSIDE_DEVICE_FUNCTION
assert(false);
#else
OPM_THROW(std::exception, "Message");
#endif

Note that while it is possible to supply the message to assert as well, this is not really practically feasible in all the places the throw macro is used, and it is probably not needed either.

While not necessary for this to build, I have also included unit tests for the updated macros in the opm-simulators PR 5962

[kjetilly]

jenkins build this opm-simulators=5962 please

[kjetilly]

jenkins build this opm-simulators=5962 please

[akva2]

I like it. Of course those writing gpu code has to be aware of the interface here as it's not all that obvious how things are tied together. (ie you have to re-check on cpu side), but at least those who are writing cpu code and are unaware they are being exploited by the gpu guys don't have to do anything special. One small request which I think makes things more readable.

[kjetilly]

jenkins build this opm-simulators=5962 please

[kjetilly]

Jenkins build this please

[bska]

Just out of curiosity: Does assert() on GPUs have the same semantics as on the CPU–i.e., if someone #defines NDEBUG, then the assert() is just removed?

[kjetilly]

Just out of curiosity: Does assert() on GPUs have the same semantics as on the CPU–i.e., if someone #defines NDEBUG, then the assert() is just removed?

Yes, that is correct. You can see this effect in the following Compiler Explorer example that compiles since the assert statement is just empty:

#define NDEBUG 1
#include <cassert>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cstdio>


// Type your code here, or load an example.
__global__ void square() {
    printf("From the kernel\n");
    assert(variable_that_does_not_exist);
}

int main(int, char**) {
    square<<<1, 1>>>();
    cudaDeviceSynchronize();
}

This is a risk, but we are also talking about code that will mostly be run in an inner loop. I can replace it with __trap() which will always abort, but I think it probably is wise to be able to compile the CUDA code in ways the compiler knows the kernel will never prematurely exit.

[bska]

This is a risk, but we are also talking about code that will mostly be run in an inner loop. I can replace it with __trap() which will always abort, but I think it probably is wise to be able to compile the CUDA code in ways the compiler knows the kernel will never prematurely exit.

Agreed. You're the expert here, so if you're happy with this trade-off, then I have no objection.

[kjetilly]

Now, the current version does not protect against:

OPM_ERROR_IF(functionThatMustBeCalled(), "message");

I can amend this, though. Adding __trap() I am a bit more dubious towards, at least before properly benchmarking it.

[kjetilly]

Jenkins build this please

[kjetilly]

Can this be merged or does it need any other changes?

[kjetilly]

This PR is in principle fine to merge, BUT, we have seen that it does not work as expected on Hipcc, and would like to wait with merging until we have figured out the Hipcc problems

[kjetilly]

jenkins build this opm-simulators=5962 please

[multitalentloes]

jenkins build this opm-simulators=5962 please

[kjetilly]

jenkins build this hipify opm-simulators=5962 please