Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Fixed dace::math::pow for Integer case #1748

Open
wants to merge 11 commits into
base: main
Choose a base branch
from
Open

Fixed dace::math::pow for Integer case #1748

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
8f4e7df
Added a pow specialization for integer.
philip-paul-mueller Nov 11, 2024
eb887d9
Enabled the FFT test again.
philip-paul-mueller Nov 11, 2024
4357ee5
Merge branch 'main' into better_pow
tbennun Nov 12, 2024
e7d7990
Fixed some errors in the power implementation.
philip-paul-mueller Nov 12, 2024
db4b411
Some better types.
philip-paul-mueller Nov 12, 2024
238f15d
Fixed a bug in `ipow`.
philip-paul-mueller Nov 12, 2024
d355c12
Made `ipow` also more general.
philip-paul-mueller Nov 12, 2024
def4d42
Added a specialization of `ipow` for signed integers.
philip-paul-mueller Nov 12, 2024
28feda5
Again fixed `ipow` now it was a redeclaration.
philip-paul-mueller Nov 12, 2024
28efc8c
Let's try this.
philip-paul-mueller Nov 12, 2024
9773ee8
Unfixed `ipow`.
philip-paul-mueller Nov 13, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
72 changes: 58 additions & 14 deletions dace/runtime/include/dace/math.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -512,40 +512,84 @@ namespace dace
return (thrust::complex<T>)thrust::pow(a, b);
}
#endif
template<typename T, typename U>


/* 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<!(std::is_integral<T>::value && std::is_integral<U>::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<typename T>
using IntPowReturnType_t = std::conditional_t<std::is_unsigned<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<std::is_integral<T1>::value && std::is_integral<T2>::value>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

So, if both a and b are integers (signed or unsigned), this path will be followed? In other words, if we want to compute, e.g., 3^-5, we need to convert one of them to float. Are we sure this is the behavior we want?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is correct.
But we kind of need this behaviour, to enable the FFT test this commit enables.
There the problem is that a size of an array is given as a = np.ones(2**3).
Which the code generator translates to

double* a;
a = new double[pow(2, 3)];

If pow would return a double, then it is a compilation error.

There are technically three things we can do:

  • Leave it as it is.
  • Adding the specialization that if the exponent is an unsigned int then we return a double and hope that expressions such as above are never called where the exponent is unsigned, i.e.
unsigned int b = 4;
//...
double* a;
a = new double[pow(2, b)];
  • Patch the code generator that sizes in new[] expressions are always casted to std::size_t.
    However, this would be another PR.

What is your opinion on this @tbennun ?

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree with @alexnick83. The reason we didn't have these pow implementations in the first place was because of negative exponents. I think we should use ipow when we know something or the equivalent version only in the case of an unsigned exponent and integer base.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I looked it up in the code generator, ipow is only used if the exponent is known at code generation time.
This is not the case in the FFT test, there the exponent is not known and thus pow is selected.

The reason why it lead to the compilation error in the FFT test was because the arguments of pow where int64 and since there is only an overload for int and unsigned int, the template was selected.
However, even before the expression 5^-3 would have resulted in the value 0, given that both were integer.

For the reasons outlined above, I do not think that this PR changes anything big.
But how should I proceed @alexnick83 @tbennun

PS: I also tried to implement ipow for negative exponents and that failed with some compilation bug related to half, which I did not understood.

>
static DACE_CONSTEXPR DACE_HDFI
IntPowReturnType_t<T1>
pow(const T1& a, const T2& b)
{
unsigned int result = 1;
for (unsigned int i = 0; i < b; ++i)
if(std::is_signed<T2>::value) {
if(b < 0)
return 0;
}
using IterationBound_t = std::make_unsigned_t<T2>;
IntPowReturnType_t<T1> 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<typename T>
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<typename T, typename std::enable_if<std::is_integral<T>::value>::type* = nullptr>
DACE_CONSTEXPR DACE_HDFI T ifloor(const T& a)
{
Expand Down
2 changes: 2 additions & 0 deletions dace/runtime/include/dace/types.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 0 additions & 2 deletions tests/npbench/misc/stockham_fft_test.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
Loading