Alpaka: MSVC and g++ host compiler fail to deduce return type of min/max

[bernhardmgruber]

When I use the alpaka::math::min or alpaka::math::max function, nvcc compiles fine, but MSVC or g++ as host compiler fails to compile the remaining code after nvcc.

Given, from mallocMC,
`` const uint32 fullsegments = alpaka::math::min(acc, 32u, pagesize / segmentsize); ```` g++ complains with:error: void value not ignored as it ought to beat the assignment tofullsegments. The error message indicates, that the host compiler sees a return type ofvoid` for the alpaka min/max functions, probably in the code it is given my nvcc.

The problem disappears, if the return type for alpaka::math::traits::Min<MinUniformCudaHipBuiltIn, ...>::min(...) is explicitely specified as e.g. Tx or decltype(::min(x, y)).

[BenjaminW3]

Can you reproduce this in an alpaka unit test so that it can be checked in CI?

[psychocoderHPC]

@sbastrakov said that it could be an issue when min is a function which can not be inlined. ::min() is in the CUDA headers defined as extern __device__ function so it could be the reason.

[psychocoderHPC]

Another problem could be that a,b in min(acc, a,b) has a different type.

[bernhardmgruber]

So while reproducing the issue in a small unit test, I ran into a different issue. But maybe this is caused by the same problem.

The following code:

#include <alpaka/alpaka.hpp>
void test() {
    using Idx = std::size_t;
    using Dim = alpaka::dim::DimInt<1>;
    using Acc = alpaka::acc::AccGpuCudaRt<Dim, Idx>;
    auto f = [] ALPAKA_FN_ACC(const Acc & acc) {
        const auto l = 4;
        const auto u = 5;
        const auto r = alpaka::math::min(acc, l, u);
    };
}

fails to compile with the host hompiler (MSVC):

1>C:/dev/mallocMC/alpaka/include\alpaka/math/min/Traits.hpp(57): error #2970: calling a __device__ function("min") from a __host__ __device__ function("min") is not allowed
1>          detected during instantiation of "auto alpaka::math::min(const T &, const Tx &, const Ty &) [with T=alpaka::acc::AccGpuCudaRt<alpaka::dim::DimInt<1Ui64>, size_t>, Tx=int, Ty=int]".

nvcc compiles fine.

The problem is that the CUDA ::min() function is marked as __device__, whereas the alpaka::math::min() is marked as __host__ __device__. Since the return type of alpaka::math::min() is deduced, the host compiler seems to fail to deduce this type from the __device__ function and assumes it's void.

[bernhardmgruber]

As a sidenote: the alpaka unit tests seem to only test the math functions on float and double, but not on integral arguments (which does not make sense for all functions).

[psychocoderHPC]

I think it is only solvable by adding the return type. I got the same issue in cupla during I implemented math functions there https://github.com/alpaka-group/cupla/blob/7e7509463b48d393edc7596d36b20de2b0fef2f4/include/cupla/device/math/Common.hpp#L89-L103

See my comment in the cupla code

return type is required for the compiler to detect host, device  function qualifier correctly                                             

[psychocoderHPC]

@bernhardmgruber could you please test if you workaround decltype(::min(x, y)) helps in you unit test.

[bernhardmgruber]

@psychocoderHPC yes, adding decltype(::min(x, y)) solves it. This is what i added a few days ago on top of the alpaka subtree in mallocMC.

[psychocoderHPC]

Is this issue already solved?

[bernhardmgruber]

Well you committed the fix here: https://github.com/alpaka-group/alpaka/commit/2bf8dad8059c0200968d800c2d7d87b92242d4d8
That should do it!

[psychocoderHPC]

solved with #1085