split __host__ __device__, include thrust in test

ginkgo-project · Nov 15, 2019 · 8ca796f · 8ca796f
1 parent 70809ce
commit 8ca796f
Show file tree

Hide file tree

Showing 6 changed files with 93 additions and 76 deletions.
diff --git a/cmake/create_test.cmake b/cmake/create_test.cmake
@@ -53,8 +53,10 @@ function(ginkgo_create_hip_test test_name)
  target_include_directories("${TEST_TARGET_NAME}"
  PRIVATE
  "$<BUILD_INTERFACE:${Ginkgo_BINARY_DIR}>"
-
- # Only `exception_helpers` requires these so far, but it's much easier
+ # Only `math` requires it so far, but it's much easier
+ # to put these this way.
+ ${GINKGO_HIP_THRUST_PATH}
+ # Only `exception_helpers` requires thess so far, but it's much easier
  # to put these this way.
  ${HIPBLAS_INCLUDE_DIRS}
  ${HIPSPARSE_INCLUDE_DIRS}

diff --git a/core/factorization/par_ilu.cpp b/core/factorization/par_ilu.cpp
@@ -78,7 +78,7 @@ ParIlu<ValueType, IndexType>::generate_l_u(
  const auto host_exec = exec->get_master();
 
  // If required, it is also possible to make this a Factory parameter
- auto csr_strategy = std::make_shared<typename CsrMatrix::cusparse>();
+ auto csr_strategy = std::make_shared<typename CsrMatrix::sparselib>();
 
  // Only copies the matrix if it is not on the same executor or was not in
  // the right format. Throws an exception if it is not convertable.

diff --git a/hip/base/hipblas_bindings.hip.hpp b/hip/base/hipblas_bindings.hip.hpp
@@ -153,10 +153,6 @@ GKO_BIND_HIPBLAS_GEAM(ValueType, detail::not_implemented);
 
 GKO_BIND_HIPBLAS_SCAL(float, hipblasSscal);
 GKO_BIND_HIPBLAS_SCAL(double, hipblasDscal);
-/* not implemented
-GKO_BIND_HIPBLAS_SCAL(std::complex<float>, hipblasCscal);
-GKO_BIND_HIPBLAS_SCAL(std::complex<double>, hipblasZscal);
-*/
 template <typename ValueType>
 GKO_BIND_HIPBLAS_SCAL(ValueType, detail::not_implemented);
 
@@ -177,10 +173,6 @@ GKO_BIND_HIPBLAS_SCAL(ValueType, detail::not_implemented);
 
 GKO_BIND_HIPBLAS_AXPY(float, hipblasSaxpy);
 GKO_BIND_HIPBLAS_AXPY(double, hipblasDaxpy);
-/* not implemented
-GKO_BIND_HIPBLAS_AXPY(std::complex<float>, hipblasCaxpy);
-GKO_BIND_HIPBLAS_AXPY(std::complex<double>, hipblasZaxpy);
-*/
 template <typename ValueType>
 GKO_BIND_HIPBLAS_AXPY(ValueType, detail::not_implemented);
 
@@ -201,10 +193,6 @@ GKO_BIND_HIPBLAS_AXPY(ValueType, detail::not_implemented);
 
 GKO_BIND_HIPBLAS_DOT(float, hipblasSdot);
 GKO_BIND_HIPBLAS_DOT(double, hipblasDdot);
-/* not implemented
-GKO_BIND_HIPBLAS_DOT(std::complex<float>, hipblasCdotc);
-GKO_BIND_HIPBLAS_DOT(std::complex<double>, hipblasZdotc);
-*/
 template <typename ValueType>
 GKO_BIND_HIPBLAS_DOT(ValueType, detail::not_implemented);
 

diff --git a/hip/matrix/csr_kernels.hip.cpp b/hip/matrix/csr_kernels.hip.cpp
@@ -938,7 +938,7 @@ __global__
 
  if (sliceid * slice_size + tid_in_warp < num_rows) {
  size_type thread_result = 0;
- for (auto i = tid_in_warp; i < slice_size; i += warp_size) {
+ for (int i = tid_in_warp; i < slice_size; i += warp_size) {
  thread_result =
  (i + slice_size * sliceid < num_rows)
  ? max(thread_result, nnz_per_row[sliceid * slice_size + i])

diff --git a/hip/matrix/dense_kernels.hip.cpp b/hip/matrix/dense_kernels.hip.cpp
@@ -632,7 +632,7 @@ __global__
 
  if (sliceid * slice_size + tid_in_warp < num_rows) {
  size_type thread_result = 0;
- for (auto i = tid_in_warp; i < slice_size; i += warp_size) {
+ for (int i = tid_in_warp; i < slice_size; i += warp_size) {
  thread_result =
  (i + slice_size * sliceid < num_rows)
  ? max(thread_result, nnz_per_row[sliceid * slice_size + i])
@@ -1019,11 +1019,11 @@ void transpose(std::shared_ptr<const HipExecutor> exec,
  hipblas::pointer_mode_guard pm_guard(handle);
  auto alpha = one<ValueType>();
  auto beta = zero<ValueType>();
- hipblas::geam(
- handle, HIPBLAS_OP_T, HIPBLAS_OP_N, orig->get_size()[0],
- orig->get_size()[1], &alpha, orig->get_const_values(),
- orig->get_stride(), &beta, static_cast<ValueType *>(nullptr),
- trans->get_size()[1], trans->get_values(), trans->get_stride());
+ hipblas::geam(handle, HIPBLAS_OP_T, HIPBLAS_OP_N,
+  orig->get_size()[0], orig->get_size()[1], &alpha,
+  orig->get_const_values(), orig->get_stride(), &beta,
+  orig->get_const_values(), trans->get_size()[1],
+   trans->get_values(), trans->get_stride());
  }
  } else {
  GKO_NOT_IMPLEMENTED;

diff --git a/include/ginkgo/core/base/math.hpp b/include/ginkgo/core/base/math.hpp
@@ -351,58 +351,85 @@ GKO_INLINE GKO_ATTRIBUTES constexpr int64 ceildiv(int64 num, int64 den)
 }
 
 
-/**
- * Returns the additive identity for T.
- *
- * @return additive identity for T
- */
-template <typename T>
-GKO_INLINE GKO_ATTRIBUTES constexpr T zero()
-{
- return T(0);
-}
-
-
-/**
- * Returns the additive identity for T.
- *
- * @return additive identity for T
- *
- * @note This version takes an unused reference argument to avoid complicated
- * calls like `zero<decltype(x)>()`. Instead, it allows `zero(x)`.
- */
-template <typename T>
-GKO_INLINE GKO_ATTRIBUTES constexpr T zero(const T &)
-{
- return zero<T>();
-}
-
-
-/**
- * Returns the multiplicative identity for T.
- *
- * @return the multiplicative identity for T
- */
-template <typename T>
-GKO_INLINE GKO_ATTRIBUTES constexpr T one()
-{
- return T(1);
-}
-
-
-/**
- * Returns the multiplicative identity for T.
- *
- * @return the multiplicative identity for T
- *
- * @note This version takes an unused reference argument to avoid complicated
- * calls like `one<decltype(x)>()`. Instead, it allows `one(x)`.
- */
-template <typename T>
-GKO_INLINE GKO_ATTRIBUTES constexpr T one(const T &)
-{
- return one<T>();
-}
+#define GKO_BIND_ZERO_ONE(_attribute) \
+ { \
+ /** \
+ * Returns the additive identity for T. \
+ * \
+ * @return additive identity for T \
+ */ \
+ template <typename T> \
+ GKO_INLINE _attribute constexpr T zero() \
+ { \
+ return T(0); \
+ } \
+ \
+ \
+ /** \
+ * Returns the additive identity for T. \
+ * \
+ * @return additive identity for T \
+ * \
+ * @note This version takes an unused reference argument to avoid \
+ * complicated calls like `zero<decltype(x)>()`. Instead, it allows \
+ * `zero(x)`. \
+ */ \
+ template <typename T> \
+ GKO_INLINE _attribute constexpr T zero(const T &) \
+ { \
+ return zero<T>(); \
+ } \
+ \
+ \
+ /** \
+ * Returns the multiplicative identity for T. \
+ * \
+ * @return the multiplicative identity for T \
+ */ \
+ template <typename T> \
+ GKO_INLINE _attribute constexpr T one() \
+ { \
+ return T(1); \
+ } \
+ \
+ \
+ /** \
+ * Returns the multiplicative identity for T. \
+ * \
+ * @return the multiplicative identity for T \
+ * \
+ * @note This version takes an unused reference argument to avoid \
+ * complicated calls like `one<decltype(x)>()`. Instead, it allows \
+ * `one(x)`. \
+ */ \
+ template <typename T> \
+ GKO_INLINE __host__ constexpr T one(const T &) \
+ { \
+ return one<T>(); \
+ } \
+ } \
+ static_assert(true, \
+ "This assert is used to counter the false positive extra " \
+ "semi-colon warnings")
+
+
+#if defined(__HIPCC__) && GINKGO_HIP_PLATFORM_HCC
+
+
+GKO_BIND_ZERO_ONE(__host__);
+GKO_BIND_ZERO_ONE(__device__);
+
+
+#else
+
+
+GKO_BIND_ZERO_ONE(GKO_ATTRIBUTES);
+
+
+#endif // defined(__HIPCC__) && GINKGO_HIP_PLATFORM_HCC
+
+
+#undef GKO_BIND_ZERO_ONE
 
 
 /**
@@ -561,7 +588,7 @@ GKO_INLINE GKO_ATTRIBUTES constexpr T get_superior_power(
 }
 
 
-#if !(defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__))
+#if !(defined(__CUDA_ARCH__))
 
 
 // Since a lot of compiler in combination with CUDA seem to have difficulties
@@ -575,7 +602,7 @@ isfinite(const T &value)
  return std::isfinite(value);
 }
 
-#endif // defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__))
+#endif // defined(__CUDA_ARCH__)
 
 
 /**