Skip to content

Commit 8c975f1

Browse files
stijnhstijnh
committed
Fix bug where fma(vec, constant, constant) does not become type of vec
1 parent 55045c0 commit 8c975f1

4 files changed

Lines changed: 55 additions & 37 deletions

File tree

‎include/kernel_float/constant.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -76,12 +76,12 @@ struct promote_type<constant<L>, constant<R>> {
7676

7777
template<typename L, typename R>
7878
struct promote_type<constant<L>, R> {
79-
using type = typename promote_type<L, R>::type;
79+
using type = R;
8080
};
8181

8282
template<typename L, typename R>
8383
struct promote_type<L, constant<R>> {
84-
using type = typename promote_type<L, R>::type;
84+
using type = L;
8585
};
8686

8787
namespace ops {

‎single_include/kernel_float.h‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -16,8 +16,8 @@
1616

1717
//================================================================================
1818
// this file has been auto-generated, do not modify its contents!
19-
// date: 2026-06-18 10:19:56.606511
20-
// git hash: 3810ae656e2763c17aa31fa5d86fb24b7c3200ef
19+
// date: 2026-06-18 14:56:21.693021
20+
// git hash: 044163ec361ae1b0757b505bccc15fa894543cb1
2121
//================================================================================
2222

2323
#ifndef KERNEL_FLOAT_MACROS_H
@@ -1406,7 +1406,7 @@ KERNEL_FLOAT_DEFINE_UNARY_MATH(tgamma)
14061406
KERNEL_FLOAT_DEFINE_UNARY_MATH(sqrt)
14071407
KERNEL_FLOAT_DEFINE_UNARY_MATH(cbrt)
14081408

1409-
KERNEL_FLOAT_DEFINE_UNARY_MATH(abs)
1409+
KERNEL_FLOAT_DEFINE_UNARY_MATH_FUN(abs, fabs, fabsf)
14101410
KERNEL_FLOAT_DEFINE_UNARY_MATH(floor)
14111411
KERNEL_FLOAT_DEFINE_UNARY_MATH(round)
14121412
KERNEL_FLOAT_DEFINE_UNARY_MATH(ceil)
@@ -2212,12 +2212,12 @@ struct promote_type<constant<L>, constant<R>> {
22122212

22132213
template<typename L, typename R>
22142214
struct promote_type<constant<L>, R> {
2215-
using type = typename promote_type<L, R>::type;
2215+
using type = R;
22162216
};
22172217

22182218
template<typename L, typename R>
22192219
struct promote_type<L, constant<R>> {
2220-
using type = typename promote_type<L, R>::type;
2220+
using type = L;
22212221
};
22222222

22232223
namespace ops {

‎tests/common.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -121,8 +121,8 @@ struct equals_helper<kf::vec<T, N>> {
121121

122122
} // namespace detail
123123

124-
template<typename T>
125-
__host__ __device__ bool equals(const T& left, const T& right) {
124+
template<typename L, typename R, typename T = std::enable_if_t<std::is_same_v<L, R>, L>>
125+
__host__ __device__ bool equals(const L& left, const R& right) {
126126
return detail::equals_helper<T>::call(left, right);
127127
}
128128

‎tests/constant.cu‎

Lines changed: 46 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -6,39 +6,49 @@ struct constant_ops_tests {
66
T value = gen.next();
77
kf::vec<T, 2> vector = {gen.next(), gen.next()};
88

9-
ASSERT_EQ(kf::make_constant(5.0) + value, T(5.0) + value);
10-
ASSERT_EQ(value + kf::make_constant(5.0), value + T(5.0));
11-
ASSERT_EQ(kf::make_constant(5.0) + vector, T(5.0) + vector);
12-
ASSERT_EQ(vector + kf::make_constant(5.0), vector + T(5.0));
9+
auto C = [](auto x) { return kf::make_constant(x); };
1310

14-
ASSERT_EQ(kf::make_constant(5.0) - value, T(5.0) - value);
15-
ASSERT_EQ(value - kf::make_constant(5.0), value - T(5.0));
16-
ASSERT_EQ(kf::make_constant(5.0) - vector, T(5.0) - vector);
17-
ASSERT_EQ(vector - kf::make_constant(5.0), vector - T(5.0));
11+
ASSERT_EQ(C(5.0) + value, T(5.0) + value);
12+
ASSERT_EQ(value + C(5.0), value + T(5.0));
13+
ASSERT_EQ(C(5.0) + vector, T(5.0) + vector);
14+
ASSERT_EQ(vector + C(5.0), vector + T(5.0));
1815

19-
ASSERT_EQ(kf::make_constant(5.0) * value, T(5.0) * value);
20-
ASSERT_EQ(value * kf::make_constant(5.0), value * T(5.0));
21-
ASSERT_EQ(kf::make_constant(5.0) * vector, T(5.0) * vector);
22-
ASSERT_EQ(vector * kf::make_constant(5.0), vector * T(5.0));
16+
ASSERT_EQ(C(5.0) - value, T(5.0) - value);
17+
ASSERT_EQ(value - C(5.0), value - T(5.0));
18+
ASSERT_EQ(C(5.0) - vector, T(5.0) - vector);
19+
ASSERT_EQ(vector - C(5.0), vector - T(5.0));
20+
21+
ASSERT_EQ(C(5.0) * value, T(5.0) * value);
22+
ASSERT_EQ(value * C(5.0), value * T(5.0));
23+
ASSERT_EQ(C(5.0) * vector, T(5.0) * vector);
24+
ASSERT_EQ(vector * C(5.0), vector * T(5.0));
2325

2426
// These results in division by zero for integers
25-
// ASSERT_EQ(kf::make_constant(5.0) / value, T(5) / value);
26-
// ASSERT_EQ(value / kf::make_constant(5.0), value / T(5));
27-
// ASSERT_EQ(kf::make_constant(5.0) / vector, T(5) / vector);
28-
// ASSERT_EQ(vector / kf::make_constant(5.0), vector / T(5));
27+
// ASSERT_EQ(C(5.0) / value, T(5) / value);
28+
// ASSERT_EQ(value / C(5.0), value / T(5));
29+
// ASSERT_EQ(C(5.0) / vector, T(5) / vector);
30+
// ASSERT_EQ(vector / C(5.0), vector / T(5));
2931
//
30-
// ASSERT_EQ(kf::make_constant(5.0) % value, T(5) % value);
31-
// ASSERT_EQ(value % kf::make_constant(5.0), value % T(5));
32-
// ASSERT_EQ(kf::make_constant(5.0) % vector, T(5) % vector);
33-
// ASSERT_EQ(vector % kf::make_constant(5.0), vector % T(5));
34-
35-
ASSERT_EQ(kf::cast<double>(kf::make_constant(T(5.0))), kf::make_vec(5.0));
36-
ASSERT_EQ(kf::cast<float>(kf::make_constant(T(5.0))), kf::make_vec(5.0f));
37-
ASSERT_EQ(kf::cast<int>(kf::make_constant(T(5.0))), kf::make_vec(5));
38-
39-
ASSERT_EQ(kf::cast<T>(kf::make_constant(double(5.0))), kf::make_vec(T(5.0)));
40-
ASSERT_EQ(kf::cast<T>(kf::make_constant(float(5.0))), kf::make_vec(T(5.0)));
41-
ASSERT_EQ(kf::cast<T>(kf::make_constant(int(5.0))), kf::make_vec(T(5.0)));
32+
// ASSERT_EQ(C(5.0) % value, T(5) % value);
33+
// ASSERT_EQ(value % C(5.0), value % T(5));
34+
// ASSERT_EQ(C(5.0) % vector, T(5) % vector);
35+
// ASSERT_EQ(vector % C(5.0), vector % T(5));
36+
37+
ASSERT_EQ(kf::cast<double>(C(T(5.0))), kf::make_vec(5.0));
38+
ASSERT_EQ(kf::cast<float>(C(T(5.0))), kf::make_vec(5.0f));
39+
ASSERT_EQ(kf::cast<int>(C(T(5.0))), kf::make_vec(5));
40+
41+
ASSERT_EQ(kf::cast<T>(C(5.0)), kf::make_vec(T(5.0)));
42+
ASSERT_EQ(kf::cast<T>(C(5.0f)), kf::make_vec(T(5.0)));
43+
ASSERT_EQ(kf::cast<T>(C(5)), kf::make_vec(T(5.0)));
44+
45+
auto expected = kf::make_vec(T(30.0));
46+
ASSERT_EQ(kf::fma(T(5), C(5.0), C(5.0)), expected);
47+
ASSERT_EQ(kf::fma(T(5), C(5.0f), C(5.0f)), expected);
48+
ASSERT_EQ(kf::fma(T(5), C(5), C(5)), expected);
49+
ASSERT_EQ(kf::fma(T(5), C(5.0f), C(5)), expected);
50+
ASSERT_EQ(kf::fma(T(5), C(5.0), C(5)), expected);
51+
ASSERT_EQ(kf::fma(T(5), C(5.0f), C(5.0)), expected);
4252
}
4353
};
4454

@@ -55,6 +65,14 @@ struct constant_eq_tests {
5565
ASSERT(kf::make_constant(double(5.0)) == T(5.0));
5666
ASSERT(kf::make_constant(float(5.0)) == T(5.0));
5767
ASSERT(kf::make_constant(int(5.0)) == T(5.0));
68+
69+
ASSERT(kf::make_constant(T(5.0)) != double(6.0));
70+
ASSERT(kf::make_constant(T(5.0)) != float(6.0));
71+
ASSERT(kf::make_constant(T(5.0)) != int(6.0));
72+
73+
ASSERT(kf::make_constant(double(5.0)) != T(6.0));
74+
ASSERT(kf::make_constant(float(5.0)) != T(6.0));
75+
ASSERT(kf::make_constant(int(5.0)) != T(6.0));
5876
}
5977
};
6078

0 commit comments

Comments
 (0)