Add ability for KERNEL_FLOAT_TILING_FOR_IMPL to accept two or three arguments

Author: stijnh

include/kernel_float/macros.h

@@ -44,4 +44,9 @@
     } while (0)
 #define KERNEL_FLOAT_UNREACHABLE __builtin_unreachable()
 
+// Somet utility macros
+#define KERNEL_FLOAT_CONCAT_IMPL(A, B) A##B
+#define KERNEL_FLOAT_CONCAT(A, B)      KERNEL_FLOAT_CONCAT_IMPL(A, B)
+#define KERNEL_FLOAT_CALL(F, ...)      F(__VA_ARGS__)
+
 #endif  //KERNEL_FLOAT_MACROS_H

include/kernel_float/meta.h

@@ -77,7 +77,7 @@ template<typename T>
 using decay_t = typename detail::decay_impl<T>::type;
 
 template<typename A, typename B>
-struct promote_type;
+struct promote_type {};
 
 template<typename T>
 struct promote_type<T, T> {

include/kernel_float/tiling.h

@@ -504,11 +504,20 @@ template<size_t TileDim, size_t BlockDim, typename D = dist::cyclic, typename In
 using tiling_1d = tiling<tile_size<TileDim>, block_size<BlockDim>, distributions<D>, IndexType>;
 
 // clang-format off
-#define KERNEL_FLOAT_TILING_FOR(TILING_VARIABLE__, INDEX_VARIABLE__, POINT_VARIABLE__)                                      \
-        _Pragma("unroll")                                                                                                   \
-        for (::std::size_t INDEX_VARIABLE__ = 0; INDEX_VARIABLE__ < TILING_VARIABLE__.size(); INDEX_VARIABLE__++)           \
-            if (typename decltype(TILING_VARIABLE__)::point_type POINT_VARIABLE__ = TILING_VARIABLE__.at(INDEX_VARIABLE__); \
-                TILING_VARIABLE__.is_present(INDEX_VARIABLE__))
+#define KERNEL_FLOAT_TILING_FOR_IMPL1(ITER_VAR, TILING, POINT_VAR, _)      \
+        _Pragma("unroll")                                                         \
+        for (size_t ITER_VAR = 0; ITER_VAR < (TILING).size(); ITER_VAR++)         \
+            if (POINT_VAR = (TILING).at(ITER_VAR); (TILING).is_present(ITER_VAR)) \
+
+#define KERNEL_FLOAT_TILING_FOR_IMPL2(ITER_VAR, TILING, INDEX_VAR, POINT_VAR)      \
+        KERNEL_FLOAT_TILING_FOR_IMPL1(ITER_VAR, TILING, POINT_VAR, _) \
+            if (INDEX_VAR = ITER_VAR; true)
+
+#define KERNEL_FLOAT_TILING_FOR_IMPL(ITER_VAR, TILING, A, B, N, ...) \
+    KERNEL_FLOAT_CALL(KERNEL_FLOAT_CONCAT(KERNEL_FLOAT_TILING_FOR_IMPL, N), ITER_VAR, TILING, A, B)
+
+#define KERNEL_FLOAT_TILING_FOR(...) \
+    KERNEL_FLOAT_TILING_FOR_IMPL(KERNEL_FLOAT_CONCAT(__tiling_index_variable__, __LINE__), __VA_ARGS__, 2, 1)
 // clang-format on
 
 }  // namespace kernel_float

tests/tiling.cu

@@ -4,12 +4,14 @@
 struct basic_tiling_test {
     template<typename T>
     __host__ __device__ void operator()(generator<T> gen) {
-        auto tiling = kf::tiling<
+        using TestTiling = kf::tiling<
             kf::tile_size<8, 8>,
             kf::block_size<2, 4>,
-            kf::distributions<kf::dist::cyclic, kf::dist::blocked>>(dim3(1, 2, 0));
+            kf::distributions<kf::dist::cyclic, kf::dist::blocked>  //
+            >;
+        auto tiling = TestTiling(dim3(1, 2, 0));
 
-        ASSERT_EQ(tiling.size(), size_t(8));
+        ASSERT_EQ(TestTiling::size(), size_t(8));
 
         ASSERT_EQ(
             tiling.local_points(),
@@ -24,7 +26,6 @@ struct basic_tiling_test {
                 kf::make_vec(7, 6)));
 
         ASSERT_EQ(tiling.local_points(0), kf::make_vec(1, 3, 5, 7, 1, 3, 5, 7));
-
         ASSERT_EQ(tiling.local_points(1), kf::make_vec(2, 2, 2, 2, 6, 6, 6, 6));
 
         ASSERT_EQ(tiling.at(0), kf::make_vec(1, 2));
@@ -44,6 +45,7 @@ struct basic_tiling_test {
         ASSERT_EQ(
             tiling.local_mask(),
             kf::make_vec(true, true, true, true, true, true, true, true));
+        ASSERT_EQ(TestTiling::all_present(), true);
         ASSERT_EQ(tiling.is_present(0), true);
         ASSERT_EQ(tiling.is_present(1), true);
         ASSERT_EQ(tiling.is_present(2), true);
@@ -54,19 +56,16 @@ struct basic_tiling_test {
         ASSERT_EQ(tiling.thread_index(2), 0);
         ASSERT_EQ(tiling.thread_index(), kf::make_vec(1, 2));
 
-        ASSERT_EQ(tiling.block_size(0), 2);
-        ASSERT_EQ(tiling.block_size(1), 4);
-        ASSERT_EQ(tiling.block_size(2), 1);
-        ASSERT_EQ(tiling.block_size(), kf::make_vec(2, 4));
-
-        ASSERT_EQ(tiling.tile_size(0), 8);
-        ASSERT_EQ(tiling.tile_size(1), 8);
-        ASSERT_EQ(tiling.tile_size(2), 1);
-        ASSERT_EQ(tiling.tile_size(), kf::make_vec(8, 8));
+        ASSERT_EQ(TestTiling::block_size(0), 2);
+        ASSERT_EQ(TestTiling::block_size(1), 4);
+        ASSERT_EQ(TestTiling::block_size(2), 1);
+        ASSERT_EQ(TestTiling::block_size(), kf::make_vec(2, 4));
 
-        ASSERT_EQ(tiling.size(), size_t(8));
+        ASSERT_EQ(TestTiling::tile_size(0), 8);
+        ASSERT_EQ(TestTiling::tile_size(1), 8);
+        ASSERT_EQ(TestTiling::tile_size(2), 1);
+        ASSERT_EQ(TestTiling::tile_size(), kf::make_vec(8, 8));
 
-        size_t counter = 0;
         const int points[8][2] = {
             {1, 2},
             {3, 2},
@@ -78,14 +77,26 @@ struct basic_tiling_test {
             {7, 6},
         };
 
-        KERNEL_FLOAT_TILING_FOR(tiling, i, point) {
-            ASSERT_EQ(counter, i);
+        size_t counter = 0;
+        KERNEL_FLOAT_TILING_FOR(tiling, auto point) {
+            ASSERT(counter < 8);
+            ASSERT_EQ(point[0], points[counter][0]);
+            ASSERT_EQ(point[1], points[counter][1]);
+            counter++;
+        }
+
+        ASSERT(counter == 8);
+
+        counter = 0;
+        KERNEL_FLOAT_TILING_FOR(tiling, int i, auto point) {
+            ASSERT(counter < 8);
+            ASSERT_EQ(counter, size_t(i));
             ASSERT_EQ(point[0], points[i][0]);
             ASSERT_EQ(point[1], points[i][1]);
             counter++;
         }
 
-        ASSERT_EQ(counter, size_t(8));
+        ASSERT(counter == 8);
     }
 };