add the test for ConfigSet and review update

Co-authored-by: Terry Cojean <terry.cojean@kit.edu>
ginkgo-project · May 19, 2021 · a26bee0 · a26bee0
1 parent ad4c587
commit a26bee0
Show file tree

Hide file tree

Showing 8 changed files with 154 additions and 54 deletions.
diff --git a/core/test/base/types.cpp b/core/test/base/types.cpp
@@ -33,6 +33,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include <ginkgo/core/base/types.hpp>
 
 
+#include <array>
+#include <cstdint>
+#include <type_traits>
+
+
 #include <gtest/gtest.h>
 
 
@@ -100,4 +105,109 @@ TEST(PrecisionReduction, ComputesCommonEncoding)
 }
 
 
+TEST(ConfigSet, MaskCorrectly)
+{
+    constexpr auto mask3_u = gko::detail::mask<3>();
+    constexpr auto fullmask_u = gko::detail::mask<32>();
+    constexpr auto mask3_u64 = gko::detail::mask<3, std::uint64_t>();
+    constexpr auto fullmask_u64 = gko::detail::mask<64, std::uint64_t>();
+
+    ASSERT_EQ(mask3_u, 7u);
+    ASSERT_EQ(fullmask_u, 0xffffffffu);
+    ASSERT_EQ((std::is_same<decltype(mask3_u), const unsigned int>::value),
+              true);
+    ASSERT_EQ((std::is_same<decltype(fullmask_u), const unsigned int>::value),
+              true);
+    ASSERT_EQ(mask3_u64, 7ull);
+    ASSERT_EQ(fullmask_u64, 0xffffffffffffffffull);
+    ASSERT_EQ((std::is_same<decltype(mask3_u64), const std::uint64_t>::value),
+              true);
+    ASSERT_EQ(
+        (std::is_same<decltype(fullmask_u64), const std::uint64_t>::value),
+        true);
+}
+
+
+TEST(ConfigSet, ShiftCorrectly)
+{
+    constexpr std::array<char, 3> bits{3, 5, 7};
+
+    constexpr auto shift0 = gko::detail::shift<3, 0>(bits);
+    constexpr auto shift1 = gko::detail::shift<3, 1>(bits);
+    constexpr auto shift2 = gko::detail::shift<3, 2>(bits);
+
+    ASSERT_EQ(shift0, 12);
+    ASSERT_EQ(shift1, 7);
+    ASSERT_EQ(shift2, 0);
+}
+
+
+TEST(ConfigSet, ConfigSet1Correctly)
+{
+    using Cfg = gko::ConfigSet<3>;
+
+    constexpr auto encoded = Cfg::encode(2);
+    constexpr auto decoded = Cfg::decode<0>(encoded);
+
+    ASSERT_EQ(encoded, 2);
+    ASSERT_EQ(decoded, 2);
+}
+
+
+TEST(ConfigSet, ConfigSet1FullCorrectly)
+{
+    using Cfg = gko::ConfigSet<32>;
+
+    constexpr auto encoded = Cfg::encode(0xffffffff);
+    constexpr auto decoded = Cfg::decode<0>(encoded);
+
+    ASSERT_EQ(encoded, 0xffffffff);
+    ASSERT_EQ(decoded, 0xffffffff);
+}
+
+
+TEST(ConfigSet, ConfigSet2FullCorrectly)
+{
+    using Cfg = gko::ConfigSet<1, 31>;
+
+    constexpr auto encoded = Cfg::encode(1, 33);
+
+    ASSERT_EQ(encoded, (1u << 31) + 33);
+}
+
+
+TEST(ConfigSet, ConfigSetSomeCorrectly)
+{
+    using Cfg = gko::ConfigSet<3, 5, 7>;
+
+    constexpr auto encoded = Cfg::encode(2, 11, 13);
+    constexpr auto decoded_0 = Cfg::decode<0>(encoded);
+    constexpr auto decoded_1 = Cfg::decode<1>(encoded);
+    constexpr auto decoded_2 = Cfg::decode<2>(encoded);
+
+    ASSERT_EQ(encoded, (2 << 12) + (11 << 7) + 13);
+    ASSERT_EQ(decoded_0, 2);
+    ASSERT_EQ(decoded_1, 11);
+    ASSERT_EQ(decoded_2, 13);
+}
+
+
+TEST(ConfigSet, ConfigSetSomeFullCorrectly)
+{
+    using Cfg = gko::ConfigSet<2, 6, 7, 17>;
+
+    constexpr auto encoded = Cfg::encode(2, 11, 13, 19);
+    constexpr auto decoded_0 = Cfg::decode<0>(encoded);
+    constexpr auto decoded_1 = Cfg::decode<1>(encoded);
+    constexpr auto decoded_2 = Cfg::decode<2>(encoded);
+    constexpr auto decoded_3 = Cfg::decode<3>(encoded);
+
+    ASSERT_EQ(encoded, (2 << 30) + (11 << 24) + (13 << 17) + 19);
+    ASSERT_EQ(decoded_0, 2);
+    ASSERT_EQ(decoded_1, 11);
+    ASSERT_EQ(decoded_2, 13);
+    ASSERT_EQ(decoded_3, 19);
+}
+
+
 }  // namespace
diff --git a/dev_tools/scripts/format_header.sh b/dev_tools/scripts/format_header.sh
@@ -155,8 +155,6 @@ CONSIDER_REGEX="${START_BLOCK_REX}|${END_BLOCK_REX}|${COMMENT_REGEX}|${INCLUDE_R
 
 # This part capture the main header and give the possible fail arrangement information
 while IFS='' read -r line || [ -n "$line" ]; do
-    # if [ "${line}" = '#include "hip/hip_runtime.h"' ] && [ "${SKIP}" = "true" ]; then
-    #     HAS_HIP_RUNTIME="true"
     if [ "${line}" = "/*${GINKGO_LICENSE_BEACON}" ] || [ "${DURING_LICENSE}" = "true" ]; then
         DURING_LICENSE="true"
         if [ "${line}" = "${GINKGO_LICENSE_BEACON}*/" ]; then
@@ -268,9 +266,6 @@ fi
 # Arrange the remain files and give
 if [ -f "${CONTENT}" ]; then
     add_regroup
-    # if [ "${HAS_HIP_RUNTIME}" = "true" ]; then
-    #     echo "#include <hip/hip_runtime.h>" > temp
-    # fi
     head -n -${KEEP_LINES} ${CONTENT} >> temp
     if [ ! -z "${IFNDEF}" ] && [ ! -z "${DEFINE}" ]; then
         # Ignore the last line #endif

diff --git a/dpcpp/base/dim3.dp.hpp b/dpcpp/base/dim3.dp.hpp
@@ -63,12 +63,12 @@ struct dim3 {
     {}
 
     /**
-     * reverse returns the range for sycl with correct ordering (reverse of
+     * get_range returns the range for sycl with correct ordering (reverse of
      * cuda)
      *
      * @return sycl::range<3>
      */
-    sycl::range<3> reverse() { return sycl::range<3>(z, y, x); }
+    sycl::range<3> get_range() { return sycl::range<3>(z, y, x); }
 };
 
 
@@ -82,8 +82,8 @@ struct dim3 {
  */
 inline sycl::nd_range<3> sycl_nd_range(dim3 grid, dim3 block)
 {
-    auto local_range = block.reverse();
-    auto global_range = grid.reverse() * local_range;
+    auto local_range = block.get_range();
+    auto global_range = grid.get_range() * local_range;
     return sycl::nd_range<3>(global_range, local_range);
 }
 

diff --git a/dpcpp/base/helper.hpp b/dpcpp/base/helper.hpp
@@ -45,18 +45,21 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 /**
  * GKO_ENABLE_DEFAULT_HOST_CONFIG gives a default host implementation for those
- * kernels which require config but do not need explicit template parameter and
- * share memory
+ * kernels which require encoded config but do not need explicit template
+ * parameter and share memory
+ *
+ * @param name_  the name of the host function with config
+ * @param kernel_  the kernel name
  */
 #define GKO_ENABLE_DEFAULT_HOST_CONFIG(name_, kernel_)                     \
-    template <int config, typename... InferredArgs>                        \
+    template <int encoded, typename... InferredArgs>                       \
     inline void name_(dim3 grid, dim3 block, size_t dynamic_shared_memory, \
-                      sycl::queue *stream, InferredArgs... args)           \
+                      sycl::queue *queue, InferredArgs... args)            \
     {                                                                      \
-        stream->submit([&](sycl::handler &cgh) {                           \
+        queue->submit([&](sycl::handler &cgh) {                            \
             cgh.parallel_for(sycl_nd_range(grid, block),                   \
                              [=](sycl::nd_item<3> item_ct1) {              \
-                                 kernel_<config>(args..., item_ct1);       \
+                                 kernel_<encoded>(args..., item_ct1);      \
                              });                                           \
         });                                                                \
     }
@@ -65,11 +68,18 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
  * GKO_ENABLE_DEFAULT_CONFIG_CALL gives a default config selection call
  * implementation for those kernels which require config selection but do not
  * need explicit template parameter
+ *
+ * @param name_  the name of the calling function
+ * @param callable_  the host function with selection
+ * @param cfg_  the ConfigSet for encode/decode method
+ * @param list_  the list for encoded config selection, whose value should be
+ *               available to decode<0> for blocksize and decode<1> for
+ *               subgroup_size by cfg_
  */
 #define GKO_ENABLE_DEFAULT_CONFIG_CALL(name_, callable_, cfg_, list_)      \
     template <typename... InferredArgs>                                    \
     void name_(dim3 grid, dim3 block, size_t dynamic_shared_memory,        \
-               sycl::queue *stream,                                        \
+               sycl::queue *queue,                                         \
                std::shared_ptr<const gko::DpcppExecutor> exec,             \
                InferredArgs... args)                                       \
     {                                                                      \
@@ -84,7 +94,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
             ::gko::syn::value_list<bool>(), ::gko::syn::value_list<int>(), \
             ::gko::syn::value_list<gko::size_type>(),                      \
             ::gko::syn::type_list<>(), grid, block, dynamic_shared_memory, \
-            stream, std::forward<InferredArgs>(args)...);                  \
+            queue, std::forward<InferredArgs>(args)...);                   \
     }
 
 // __WG_BOUND__ gives the cuda-like launch bound in cuda ordering

diff --git a/dpcpp/components/cooperative_groups.dp.hpp b/dpcpp/components/cooperative_groups.dp.hpp
@@ -348,7 +348,7 @@ class thread_block_tile<1> {
 using detail::thread_block_tile;
 
 
-// Only support tile_partition with 8, 16, 32.
+// Only support tile_partition with 2, 4, 8, 16, 32, 64.
 template <unsigned Size, typename Group>
 __dpct_inline__
     std::enable_if_t<(Size > 1) && Size <= 64 && !(Size & (Size - 1)),

diff --git a/dpcpp/test/base/dim3.dp.cpp b/dpcpp/test/base/dim3.dp.cpp
@@ -45,20 +45,10 @@ namespace {
 using namespace gko::kernels::dpcpp;
 
 
-class DpcppDim3 : public ::testing::Test {
-protected:
-    DpcppDim3() {}
-
-    void SetUp() {}
-
-    void TearDown() {}
-};
-
-
-TEST_F(DpcppDim3, CanGenerate1DRange)
+TEST(DpcppDim3, CanGenerate1DRange)
 {
     dim3 block(3);
-    auto sycl_block = block.reverse();
+    auto sycl_block = block.get_range();
 
     ASSERT_EQ(block.x, 3);
     ASSERT_EQ(block.y, 1);
@@ -69,10 +59,10 @@ TEST_F(DpcppDim3, CanGenerate1DRange)
 }
 
 
-TEST_F(DpcppDim3, CanGenerate2DRange)
+TEST(DpcppDim3, CanGenerate2DRange)
 {
     dim3 block(3, 5);
-    auto sycl_block = block.reverse();
+    auto sycl_block = block.get_range();
 
     ASSERT_EQ(block.x, 3);
     ASSERT_EQ(block.y, 5);
@@ -83,10 +73,10 @@ TEST_F(DpcppDim3, CanGenerate2DRange)
 }
 
 
-TEST_F(DpcppDim3, CanGenerate3DRange)
+TEST(DpcppDim3, CanGenerate3DRange)
 {
     dim3 block(3, 5, 7);
-    auto sycl_block = block.reverse();
+    auto sycl_block = block.get_range();
 
     ASSERT_EQ(block.x, 3);
     ASSERT_EQ(block.y, 5);
@@ -97,7 +87,7 @@ TEST_F(DpcppDim3, CanGenerate3DRange)
 }
 
 
-TEST_F(DpcppDim3, CanGenerateNDRange)
+TEST(DpcppDim3, CanGenerateNDRange)
 {
     dim3 block(3, 5, 7);
     dim3 grid(17, 13, 11);

diff --git a/dpcpp/test/components/cooperative_groups_kernels.dp.cpp b/dpcpp/test/components/cooperative_groups_kernels.dp.cpp
@@ -70,8 +70,8 @@ class CooperativeGroups : public testing::TestWithParam<int> {
     CooperativeGroups()
         : ref(gko::ReferenceExecutor::create()),
           dpcpp(gko::DpcppExecutor::create(0, ref)),
-          max_test_case(3),
-          max_num(max_test_case * 64),
+          test_case(3),
+          max_num(test_case * 64),
           result(ref, max_num),
           dresult(dpcpp)
     {
@@ -87,7 +87,7 @@ class CooperativeGroups : public testing::TestWithParam<int> {
         auto subgroup_size = GetParam();
         auto exec_info = dpcpp->get_const_exec_info();
         if (exec_info.validate(subgroup_size, subgroup_size)) {
-            for (int i = 0; i < max_test_case * subgroup_size; i++) {
+            for (int i = 0; i < test_case * subgroup_size; i++) {
                 result.get_data()[i] = true;
             }
 
@@ -102,7 +102,7 @@ class CooperativeGroups : public testing::TestWithParam<int> {
         }
     }
 
-    int max_test_case;
+    int test_case;
     int max_num;
     std::shared_ptr<gko::ReferenceExecutor> ref;
     std::shared_ptr<gko::DpcppExecutor> dpcpp;
@@ -111,13 +111,6 @@ class CooperativeGroups : public testing::TestWithParam<int> {
 };
 
 
-void test_assert(bool *success, bool partial)
-{
-    if (!partial) {
-        *success = false;
-    }
-}
-
 // kernel implementation
 template <int config>
 __WG_BOUND__(KCfg::decode<0>(config))

diff --git a/include/ginkgo/core/base/types.hpp b/include/ginkgo/core/base/types.hpp
@@ -268,6 +268,8 @@ constexpr std::enable_if_t<(num_groups > current_shift + 1), int> shift(
  * 1/2/k position and r is for rest of unused bits.
  *
  * @tparam num_bits...  the number of bits for each position.
+ *
+ * @note the num_bit is required at least $log_2(maxval) + 1$
  */
 template <int... num_bits>
 class ConfigSet {
@@ -285,7 +287,7 @@ class ConfigSet {
      * @return the decoded information at position
      */
     template <int position>
-    static constexpr int decode(int encoded)
+    static constexpr unsigned decode(unsigned encoded)
     {
         static_assert(position < num_groups,
                       "This position is over the bounds.");
@@ -300,7 +302,7 @@ class ConfigSet {
      * @note the last case of nested template.
      */
     template <size_type current_iter>
-    static constexpr std::enable_if_t<(current_iter == num_groups), int>
+    static constexpr std::enable_if_t<(current_iter == num_groups), unsigned>
     encode()
     {
         return 0;
@@ -310,16 +312,16 @@ class ConfigSet {
      * Encodes the information with given bit set to encoded integer.
      *
      * @tparam current_iter  the encoded place
-     * @tparam First  the current_iter type
-     * @tparam Rest...  the rest information
+     * @tparam Rest...  the rest type
      *
-     * @param informations... the information will be encoded
+     * @param first  the current encoded information
+     * @param rest...  the rest of others information waits for encoding
      *
      * @return the encoded integer
      */
-    template <size_type current_iter = 0, typename First, typename... Rest>
-    static constexpr std::enable_if_t<(current_iter < num_groups), int> encode(
-        First first, Rest &&... rest)
+    template <size_type current_iter = 0, typename... Rest>
+    static constexpr std::enable_if_t<(current_iter < num_groups), unsigned>
+    encode(unsigned first, Rest &&... rest)
     {
         constexpr int shift = detail::shift<num_groups, current_iter>(bits);
         return (first << shift) |