diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index 4dae494a8a..1ed6fb45d2 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -512,40 +512,84 @@ namespace dace return (thrust::complex)thrust::pow(a, b); } #endif - template + + +/* If `DACE_XILINX` is defined then `pow` only returns floats, even if their arguments + * are integer. If it is not defined, then if all arguments of `pow` are integers + * the return value will also be an integer, currently it is always the largest supported by the platform. + * This is the original behaviour prior to [PR#1748](https://github.com/spcl/dace/pull/1748). */ +#if defined(DACE_XILINX) + + template< + typename T, + typename U + > DACE_CONSTEXPR DACE_HDFI auto pow(const T& a, const U& b) { return std::pow(a, b); } -#ifndef DACE_XILINX - static DACE_CONSTEXPR DACE_HDFI int pow(const int& a, const int& b) +#else + + template< + typename T, + typename U, + typename = std::enable_if_t::value && std::is_integral::value)> + > + DACE_CONSTEXPR DACE_HDFI auto pow(const T& a, const U& b) { - if (b < 0) return 0; - int result = 1; - for (int i = 0; i < b; ++i) - result *= a; - return result; + return std::pow(a, b); } - static DACE_CONSTEXPR DACE_HDFI unsigned int pow(const unsigned int& a, - const unsigned int& b) + + //TODO: Should this always be the largest integer? + template + using IntPowReturnType_t = std::conditional_t::value, uintmax, intmax>; + + + /* TODO: The return value is always an integer, this is different from the behaviour of `std::pow` that + * always return a float. We should probably patch the code generator to generate `ipow` calls if all + * arguments are integers. */ + template< + typename T1, + typename T2, + typename = std::enable_if_t::value && std::is_integral::value> + > + static DACE_CONSTEXPR DACE_HDFI + IntPowReturnType_t + pow(const T1& a, const T2& b) { - unsigned int result = 1; - for (unsigned int i = 0; i < b; ++i) + if(std::is_signed::value) { + if(b < 0) + return 0; + } + using IterationBound_t = std::make_unsigned_t; + IntPowReturnType_t result = 1; + const IterationBound_t stop = b; + //TODO: Implement logarithmic version. + for (IterationBound_t i = 0; i < stop; ++i) result *= a; return result; - } + } #endif + /* Implements the power where the exponent is known as code generation time. + * + * Furthermore, as in accordance with `CPPUnparse._BinOp` the function assumes + * that the exponent is a positive, i.e. $b > 0` integer. + * + * TODO: Find out whyt the function can not be `constexpr`. + */ template - DACE_HDFI T ipow(const T& a, const unsigned int& b) { + DACE_HDFI T ipow(const T& a, const unsigned int b) + { T result = a; for (unsigned int i = 1; i < b; ++i) result *= a; return result; } + template::value>::type* = nullptr> DACE_CONSTEXPR DACE_HDFI T ifloor(const T& a) { diff --git a/dace/runtime/include/dace/types.h b/dace/runtime/include/dace/types.h index e5eed1e35e..da13b89680 100644 --- a/dace/runtime/include/dace/types.h +++ b/dace/runtime/include/dace/types.h @@ -74,6 +74,8 @@ namespace dace typedef uint16_t uint16; typedef uint32_t uint32; typedef uint64_t uint64; + typedef intmax_t intmax; + typedef uintmax_t uintmax; typedef unsigned int uint; typedef float float32; typedef double float64; diff --git a/tests/npbench/misc/stockham_fft_test.py b/tests/npbench/misc/stockham_fft_test.py index 5878cf621a..eeedaf2034 100644 --- a/tests/npbench/misc/stockham_fft_test.py +++ b/tests/npbench/misc/stockham_fft_test.py @@ -155,12 +155,10 @@ def run_stockham_fft(device_type: dace.dtypes.DeviceType): return sdfg -@pytest.mark.skip(reason="Assertion error in read_and_write_sets") def test_cpu(): run_stockham_fft(dace.dtypes.DeviceType.CPU) -@pytest.mark.skip(reason="Assertion error in read_and_write_sets") @pytest.mark.gpu def test_gpu(): run_stockham_fft(dace.dtypes.DeviceType.GPU)