Merge Dpcpp port isai (row_limit=32)

This pr ports the isai to dpcpp executor under row_limit=32 limitation and also ports some essential components like ffs.

Related PR: #924

.github/workflows/windows-cygwin.yml

@@ -53,7 +53,7 @@ jobs:
         mkdir build
         cd build
         bash -c "cmake -DBUILD_SHARED_LIBS=${{ matrix.config.shared }} -DCMAKE_BUILD_TYPE=${{ matrix.config.build_type }} -DGINKGO_COMPILER_FLAGS=${{ matrix.config.cflags }} .."
-        bash -c "make -j4"
+        bash -c "make -j2"
         bash -c "ctest . --output-on-failure"
       shell: cmd
 

common/cuda_hip/components/searching.hpp.inc

@@ -66,12 +66,12 @@ __forceinline__ __device__ IndexType binary_search(IndexType offset,
  * @internal
  * Generic implementation of a fixed-size binary search.
  * The implementation makes sure that the number of predicate evaluations only
- * depends on `length` and not on the actual position of the partition point.
- * It assumes that the predicate partitions the range [offset, offset + length)
- * into two subranges [offset, middle), [middle, offset + length) such that
- * the predicate is `false` for all elements in the first range and `true` for
- * all elements in the second range. `middle` is called the partition point.
- * If the predicate is `false` everywhere, `middle` equals `offset + length`.
+ * depends on `size` and not on the actual position of the partition point.
+ * It assumes that the predicate partitions the range [0, size) into two
+ * subranges [0, middle), [middle, size) such that the predicate is `false` for
+ * all elements in the first range and `true` for all elements in the second
+ * range. `middle` is called the partition point.
+ * If the predicate is `false` everywhere, `middle` equals `size`.
  *
  * @tparam size  the length of the partitioned range - must be a power of two
  * @param p  the predicate to be evaluated on the range - it should not have
@@ -103,12 +103,12 @@ __forceinline__ __device__ int synchronous_fixed_binary_search(Predicate p)
  * @internal
  * Generic implementation of a synchronous binary search.
  * The implementation makes sure that the number of predicate evaluations only
- * depends on `length` and not on the actual position of the partition point.
- * It assumes that the predicate partitions the range [offset, offset + length)
- * into two subranges [offset, middle), [middle, offset + length) such that
- * the predicate is `false` for all elements in the first range and `true` for
- * all elements in the second range. `middle` is called the partition point.
- * If the predicate is `false` everywhere, `middle` equals `offset + length`.
+ * depends on `size` and not on the actual position of the partition point.
+ * It assumes that the predicate partitions the range [0, size) into two
+ * subranges [0, middle), [middle, size) such that the predicate is `false` for
+ * all elements in the first range and `true` for all elements in the second
+ * range. `middle` is called the partition point.
+ * If the predicate is `false` everywhere, `middle` equals `size`.
  *
  * @param size  the length of the partitioned range - must be a power of two
  * @param p  the predicate to be evaluated on the range - it should not have

cuda/preconditioner/isai_kernels.cu

@@ -48,6 +48,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "cuda/components/reduction.cuh"
 #include "cuda/components/thread_ids.cuh"
 #include "cuda/components/uninitialized_array.hpp"
+#include "cuda/components/warp_blas.cuh"
 
 
 namespace gko {
@@ -66,8 +67,6 @@ constexpr int subwarps_per_block{2};
 constexpr int default_block_size{subwarps_per_block * subwarp_size};
 
 
-#include "common/cuda_hip/components/atomic.hpp.inc"
-#include "common/cuda_hip/components/warp_blas.hpp.inc"
 #include "common/cuda_hip/preconditioner/isai_kernels.hpp.inc"
 
 

dev_tools/oneapi/convert_source.sh

@@ -16,8 +16,9 @@
 #                    If GTEST_HEADER_DIR is available elsewhere GINKGO_BUILD_TESTS is not required.
 #                    CMake's step is not required if copying the ginkgo config.hpp from another ginkgo build into "${ROOT_DIR}/include/ginkgo/".
 #   ROOT_BUILD_DIR: the complete path for build folder. The default is "${ROOT_DIR}/${BUILD_DIR}"
-#   GTEST_HEADER_DIR: the gtest header folder. The default is "${ROOT_BUILD_DIR}/third_party_gtest/src/googletest/include"
+#   GTEST_HEADER_DIR: the gtest header folder. The default is "${ROOT_BUILD_DIR}/_deps/googletest-src/googletest/include"
 #   CLANG_FORMAT: the clang-format exec. The default is "clang-format"
+#   VERBOSE: if it is set as 1, script will ouput the path information
 CURRENT_DIR="$( pwd )"
 cd "$( dirname "${BASH_SOURCE[0]}" )"
 SCRIPT_DIR="$( pwd )"
@@ -300,6 +301,8 @@ replace_regex="${replace_regex};s|trick/thread_ids.hpp|dpcpp/components/thread_i
 replace_regex="${replace_regex};s|trick/cooperative_groups\.hpp|dpcpp/components/cooperative_groups.dp.hpp|g"
 replace_regex="${replace_regex};s|trick/sorting\.hpp|dpcpp/components/sorting.dp.hpp|g"
 replace_regex="${replace_regex};s|trick/reduction\.hpp|dpcpp/components/reduction.dp.hpp|g"
+replace_regex="${replace_regex};s|cuda/base/math\.hpp|limits|g"
+replace_regex="${replace_regex};s|device_numeric_limits(<.*>)::(.*)(,|;)|std::numeric_limits\1::\2()\3|g"
 replace_regex="${replace_regex};s/#define GET_QUEUE 0//g"
 replace_regex="${replace_regex};s/GET_QUEUE/exec->get_queue()/g"
 replace_regex="${replace_regex};s/cuda/dpcpp/g"

dpcpp/components/intrinsics.dp.hpp

@@ -0,0 +1,96 @@
+/*******************************<GINKGO LICENSE>******************************
+Copyright (c) 2017-2022, the Ginkgo authors
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions
+are met:
+
+1. Redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
+IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
+PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
+LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+******************************<GINKGO LICENSE>*******************************/
+
+#ifndef GKO_DPCPP_COMPONENTS_INTRINSICS_DP_HPP_
+#define GKO_DPCPP_COMPONENTS_INTRINSICS_DP_HPP_
+
+
+#include <CL/sycl.hpp>
+
+
+#include <ginkgo/core/base/types.hpp>
+
+
+#include "dpcpp/base/dpct.hpp"
+
+
+namespace gko {
+namespace kernels {
+namespace dpcpp {
+
+
+/**
+ * @internal
+ * Returns the number of set bits in the given mask.
+ */
+__dpct_inline__ int popcnt(uint32 mask) { return sycl::popcount(mask); }
+
+/** @copydoc popcnt */
+__dpct_inline__ int popcnt(uint64 mask) { return sycl::popcount(mask); }
+
+
+/**
+ * @internal
+ * Returns the (1-based!) index of the first set bit in the given mask,
+ * starting from the least significant bit.
+ *
+ * @note can not use length - clz because subgroup size is not always 32 or 64
+ */
+__dpct_inline__ int ffs(uint32 mask)
+{
+    return (mask == 0) ? 0 : (sycl::ext::intel::ctz(mask) + 1);
+}
+
+/** @copydoc ffs */
+__dpct_inline__ int ffs(uint64 mask)
+{
+    return (mask == 0) ? 0 : (sycl::ext::intel::ctz(mask) + 1);
+}
+
+
+/**
+ * @internal
+ * Returns the number of zero bits before the first set bit in the given mask,
+ * starting from the most significant bit.
+ */
+__dpct_inline__ int clz(uint32 mask) { return sycl::clz(mask); }
+
+/** @copydoc clz */
+__dpct_inline__ int clz(uint64 mask) { return sycl::clz(mask); }
+
+
+}  // namespace dpcpp
+}  // namespace kernels
+}  // namespace gko
+
+
+#endif  // GKO_DPCPP_COMPONENTS_INTRINSICS_DP_HPP_