From 8f4e7df1b386990491bfe9d6594839b11a6fe433 Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Mon, 11 Nov 2024 16:50:44 +0100 Subject: [PATCH 01/10] Added a pow specialization for integer. The new specialization works for integers. So if all arguments are integer the return value is guaranteed to be integer as well. This should also solve the issue with the fft test mentioned by Tal. https://github.com/spcl/dace/pull/1742/files#diff-a81c4594b19da98f6516283a50f6958b04976ae2913f097358141b2681868f3eR158 --- dace/runtime/include/dace/math.h | 36 +++++++++++++++++++------------- 1 file changed, 22 insertions(+), 14 deletions(-) diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index 4dae494a8a..3b9defb176 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -512,32 +512,40 @@ namespace dace return (thrust::complex)thrust::pow(a, b); } #endif - template + 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) { return std::pow(a, b); } #ifndef DACE_XILINX - static DACE_CONSTEXPR DACE_HDFI int pow(const int& a, const int& b) - { - if (b < 0) return 0; - int result = 1; - for (int i = 0; i < b; ++i) - result *= a; - return result; - } + template + using IntPowReturnType_t = std::conditional_t::value, unsigned long long int, long long int>; - static DACE_CONSTEXPR DACE_HDFI unsigned int pow(const unsigned int& a, - const unsigned int& b) + template + static DACE_CONSTEXPR DACE_HDFI + std::enable_if_t::value && std::is_integral::value, IntPowReturnType_t > + pow(const T1& a, const T2& b) { - unsigned int result = 1; - for (unsigned int i = 0; i < b; ++i) + /* 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. */ + if(std::is_signed::value && (b < 0)) + return 0; + using IterationBound_t = typename std::make_unsigned::type; + IntPowReturnType_t result = 1; + const IterationBound_t stop = b; + for (IterationBound_t i = 0; i < stop; ++i) result *= a; return result; - } + }; #endif + template DACE_HDFI T ipow(const T& a, const unsigned int& b) { T result = a; From eb887d917332fb5ff0f23c9705cf1469e371d994 Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Mon, 11 Nov 2024 16:54:36 +0100 Subject: [PATCH 02/10] Enabled the FFT test again. However, to work we need https://github.com/spcl/dace/pull/1747. --- tests/npbench/misc/stockham_fft_test.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/tests/npbench/misc/stockham_fft_test.py b/tests/npbench/misc/stockham_fft_test.py index 8fc5e88203..240cc1dc9f 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) From e7d7990e7ce7712e3f35eb0c758f606d0ac61f9e Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Tue, 12 Nov 2024 08:04:44 +0100 Subject: [PATCH 03/10] Fixed some errors in the power implementation. --- dace/runtime/include/dace/math.h | 45 +++++++++++++++++++++++++------- 1 file changed, 35 insertions(+), 10 deletions(-) diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index 3b9defb176..d8df71074a 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -512,6 +512,23 @@ namespace dace return (thrust::complex)thrust::pow(a, b); } #endif + + +#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); + } + +#else + /* If `DACE_XILINX` is not defined, there is a specialization for `pow` + * which only operates on integers. This allows it to use in `new[]` + * expressions. */ template< typename T, typename U, @@ -522,30 +539,38 @@ namespace dace return std::pow(a, b); } -#ifndef DACE_XILINX + + //TODO: Should this always be the largest integer? template using IntPowReturnType_t = std::conditional_t::value, unsigned long long int, long long int>; - template + + /* 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 - std::enable_if_t::value && std::is_integral::value, IntPowReturnType_t > + IntPowReturnType_t pow(const T1& a, const T2& b) { - /* 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. */ - if(std::is_signed::value && (b < 0)) - return 0; - using IterationBound_t = typename std::make_unsigned::type; + 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 - template DACE_HDFI T ipow(const T& a, const unsigned int& b) { T result = a; From db4b411b874fd0e6d41bb7b7191de6701f8e8a0e Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Tue, 12 Nov 2024 08:12:09 +0100 Subject: [PATCH 04/10] Some better types. --- dace/runtime/include/dace/math.h | 10 ++++++---- dace/runtime/include/dace/types.h | 2 ++ 2 files changed, 8 insertions(+), 4 deletions(-) diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index d8df71074a..4486a6381c 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -514,6 +514,10 @@ namespace dace #endif +/* 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< @@ -526,9 +530,7 @@ namespace dace } #else - /* If `DACE_XILINX` is not defined, there is a specialization for `pow` - * which only operates on integers. This allows it to use in `new[]` - * expressions. */ + template< typename T, typename U, @@ -542,7 +544,7 @@ namespace dace //TODO: Should this always be the largest integer? template - using IntPowReturnType_t = std::conditional_t::value, unsigned long long int, long long int>; + 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 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; From 238f15d7d26f00cdc381bf1b8433cade61ecf1de Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Tue, 12 Nov 2024 08:15:56 +0100 Subject: [PATCH 05/10] Fixed a bug in `ipow`. --- dace/runtime/include/dace/math.h | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index 4486a6381c..2df6f43093 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -574,7 +574,10 @@ namespace dace #endif template - DACE_HDFI T ipow(const T& a, const unsigned int& b) { + DACE_CONSTEXPR DACE_HDFI T ipow(const T& a, const unsigned int& b) { + if(b == 0) { + return T(1); + }; T result = a; for (unsigned int i = 1; i < b; ++i) result *= a; From d355c12250cb2793e8acbe2387ccdda2aef66cf1 Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Tue, 12 Nov 2024 08:20:42 +0100 Subject: [PATCH 06/10] Made `ipow` also more general. --- dace/runtime/include/dace/math.h | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index 2df6f43093..747be8bf93 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -573,13 +573,18 @@ namespace dace }; #endif - template - DACE_CONSTEXPR DACE_HDFI T ipow(const T& a, const unsigned int& b) { + template< + typename T, + typename U, + typename = std::enable_if_t<(std::is_integral::value && std::is_unsigned::value)> + > + DACE_CONSTEXPR DACE_HDFI T ipow(const T& a, const U b) + { if(b == 0) { return T(1); }; T result = a; - for (unsigned int i = 1; i < b; ++i) + for (U i = 1; i < b; ++i) result *= a; return result; } From def4d4291697ad43c96d747fd30ba659fe7d779e Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Tue, 12 Nov 2024 08:37:58 +0100 Subject: [PATCH 07/10] Added a specialization of `ipow` for signed integers. The original signature of `ipow` only accepted an unsigned integer as exponent. However, it seems that it was called using integers as exponents. Thus there is now a specialization for this case. For the sake of argument let's consider the following situation `ipow(a, b)`: - If `b` was zero then the previous implementation would return `a`, the new implementation returns 1. - If `b` was negative then it would perform a lot of iterations, because a signed integer was converted to an unsigned one. In the new implementation, the function would return 0, which is in agreement with what `pow` does. --- dace/runtime/include/dace/math.h | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index 747be8bf93..61c329ced3 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -589,6 +589,22 @@ namespace dace return result; } + + template< + typename T, + typename U, + typename = std::enable_if_t<(std::is_integral::value && std::is_signed::value)> + > + DACE_CONSTEXPR DACE_HDFI T ipow(const T& a, const U b) + { + if(b < 0) { + return T(0); + }; + using UnsignedU = std::make_unsigned_t; + return ipow(a, UnsignedU{b}); + } + + template::value>::type* = nullptr> DACE_CONSTEXPR DACE_HDFI T ifloor(const T& a) { From 28feda5fa5e7ae1054477dbdf74df0f0ea4dad3f Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Tue, 12 Nov 2024 08:58:34 +0100 Subject: [PATCH 08/10] Again fixed `ipow` now it was a redeclaration. I now combined both versions into one, this makes it similar to what `pow` is. However another version would be to put the enable if on the arguments, but I thought it would be too ugly. --- dace/runtime/include/dace/math.h | 25 ++++++++----------------- 1 file changed, 8 insertions(+), 17 deletions(-) diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index 61c329ced3..61e7c0a640 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -576,35 +576,26 @@ namespace dace template< typename T, typename U, - typename = std::enable_if_t<(std::is_integral::value && std::is_unsigned::value)> + typename = std::enable_if_t::value> > DACE_CONSTEXPR DACE_HDFI T ipow(const T& a, const U b) { + if(std::is_signed::value) { + if(b < 0) + return 0; + } if(b == 0) { return T(1); }; + using IterationBound_t = std::make_unsigned_t; T result = a; - for (U i = 1; i < b; ++i) + const IterationBound_t stop{b}; + for (IterationBound_t i = 1; i < stop; ++i) result *= a; return result; } - template< - typename T, - typename U, - typename = std::enable_if_t<(std::is_integral::value && std::is_signed::value)> - > - DACE_CONSTEXPR DACE_HDFI T ipow(const T& a, const U b) - { - if(b < 0) { - return T(0); - }; - using UnsignedU = std::make_unsigned_t; - return ipow(a, UnsignedU{b}); - } - - template::value>::type* = nullptr> DACE_CONSTEXPR DACE_HDFI T ifloor(const T& a) { From 28efc8c65d37e8807ceb6d08e0b59e3ce9feec22 Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Tue, 12 Nov 2024 15:07:13 +0100 Subject: [PATCH 09/10] Let's try this. --- dace/runtime/include/dace/math.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index 61e7c0a640..0fde72fd4a 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -578,7 +578,7 @@ namespace dace typename U, typename = std::enable_if_t::value> > - DACE_CONSTEXPR DACE_HDFI T ipow(const T& a, const U b) + DACE_HDFI T ipow(const T& a, const U b) { if(std::is_signed::value) { if(b < 0) From 9773ee819b2da91528bb4eec1746d26eef753051 Mon Sep 17 00:00:00 2001 From: Philip Mueller Date: Wed, 13 Nov 2024 08:00:47 +0100 Subject: [PATCH 10/10] Unfixed `ipow`. After some digging I found out that `ipow` is only called for positive integers that are known at compile time. And I know now whym because the `half` vectors do not play nice. --- dace/runtime/include/dace/math.h | 28 +++++++++++----------------- 1 file changed, 11 insertions(+), 17 deletions(-) diff --git a/dace/runtime/include/dace/math.h b/dace/runtime/include/dace/math.h index 0fde72fd4a..1ed6fb45d2 100644 --- a/dace/runtime/include/dace/math.h +++ b/dace/runtime/include/dace/math.h @@ -570,27 +570,21 @@ namespace dace for (IterationBound_t i = 0; i < stop; ++i) result *= a; return result; - }; + } #endif - template< - typename T, - typename U, - typename = std::enable_if_t::value> - > - DACE_HDFI T ipow(const T& a, const U b) + /* 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) { - if(std::is_signed::value) { - if(b < 0) - return 0; - } - if(b == 0) { - return T(1); - }; - using IterationBound_t = std::make_unsigned_t; T result = a; - const IterationBound_t stop{b}; - for (IterationBound_t i = 1; i < stop; ++i) + for (unsigned int i = 1; i < b; ++i) result *= a; return result; }