Review updates.

+ Add some environment variable output to `get_info.cmake`.
+ Fix formatting of `GKO_NOT_IMPLEMENTED` declarations.
+ Allow lowercase in the `DPC++` strings passed to executors.
+ Fix kernels storing SYCL `id` into `int` types.
+ Add some extra comments to tricky parts of the code.
+ Improve some other code formatting issues.

Co-authored-by: Tobias Ribizel <ribizel@kit.edu>
Co-authored-by: Yuhsiang M. Tsai <yhmtsai@gmail.com>

cmake/get_info.cmake

@@ -72,6 +72,21 @@ function(ginkgo_print_variable log_type var_name)
     FILE(APPEND ${log_type} "${upd_string}")
 endfunction()
 
+
+function(ginkgo_print_env_variable log_type var_name)
+    string(SUBSTRING
+        "
+--        ${var_name}:                                                          " 0 55 upd_string)
+    if(NOT ENV{${var_name}})
+        set(str_value "<empty>")
+    else()
+        set(str_value "$ENV{${var_name}}")
+    endif()
+    string(APPEND upd_string "${str_value}")
+    FILE(APPEND ${log_type} "${upd_string}")
+endfunction()
+
+
 macro(ginkgo_print_foreach_variable variables)
     foreach(var ${variables})
         ginkgo_print_variable(${${log_type}} ${var} )

core/test/reorder/rcm.cpp

@@ -1,5 +1,5 @@
 /*******************************<GINKGO LICENSE>******************************
-Copyright (c) 2017-2019, the Ginkgo authors
+Copyright (c) 2017-2020, the Ginkgo authors
 All rights reserved.
 
 Redistribution and use in source and binary forms, with or without

dpcpp/base/executor.dp.cpp

@@ -33,6 +33,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include <ginkgo/core/base/executor.hpp>
 
 
+#include <algorithm>
+#include <cctype>
 #include <iostream>
 #include <map>
 #include <string>
@@ -56,6 +58,8 @@ const std::vector<sycl::device> get_devices(std::string device_type)
         {"all", sycl::info::device_type::all},
         {"cpu", sycl::info::device_type::cpu},
         {"gpu", sycl::info::device_type::gpu}};
+    std::for_each(device_type.begin(), device_type.end(),
+                  [](char &c) { c = std::tolower(c); });
     return sycl::device::get_devices(device_type_map.at(device_type));
 }
 
@@ -123,7 +127,6 @@ void DpcppExecutor::raw_copy_to(const DpcppExecutor *dest, size_type num_bytes,
                                 const void *src_ptr, void *dest_ptr) const
 {
     if (num_bytes > 0) {
-        // TODO: does this work? Or is it needed to go through host?
         dest->get_queue()->memcpy(dest_ptr, src_ptr, num_bytes).wait();
     }
 }
@@ -147,13 +150,19 @@ int DpcppExecutor::get_num_devices(std::string device_type)
 }
 
 
+namespace detail {
+
+
 void delete_queue(sycl::queue *queue)
 {
     queue->wait();
     delete queue;
 }
 
 
+}  // namespace detail
+
+
 void DpcppExecutor::set_device_property()
 {
     assert(device_id_ < DpcppExecutor::get_num_devices(device_type_));
@@ -174,8 +183,13 @@ void DpcppExecutor::set_device_property()
     }
     max_workgroup_size_ =
         device.get_info<sycl::info::device::max_work_group_size>();
+    // Here we declare the queue with the property `in_order` which ensures the
+    // kernels are executed in the submission order. Otherwise, calls to
+    // `wait()` would be needed after every call to a DPC++ function or kernel.
+    // For example, without `in_order`, doing a copy, a kernel, and a copy, will
+    // not necessarily happen in that order by default, which we need to avoid.
     auto *queue = new sycl::queue{device, sycl::property::queue::in_order{}};
-    queue_ = std::move(queue_manager<sycl::queue>{queue, delete_queue});
+    queue_ = std::move(queue_manager<sycl::queue>{queue, detail::delete_queue});
 }
 
 

dpcpp/components/absolute_array.dp.cpp

@@ -51,7 +51,7 @@ void inplace_absolute_array(std::shared_ptr<const DefaultExecutor> exec,
 {
     exec->get_queue()->submit([&](sycl::handler &cgh) {
         cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx_id) {
-            const int idx = idx_id[0];
+            const auto idx = idx_id[0];
             data[idx] = dpcpp::abs(data[idx]);
         });
     });
@@ -67,7 +67,7 @@ void outplace_absolute_array(std::shared_ptr<const DefaultExecutor> exec,
 {
     exec->get_queue()->submit([&](sycl::handler &cgh) {
         cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx_id) {
-            const int idx = idx_id[0];
+            const auto idx = idx_id[0];
             out[idx] = dpcpp::abs(in[idx]);
         });
     });

dpcpp/components/fill_array.dp.cpp

@@ -48,7 +48,7 @@ void fill_array(std::shared_ptr<const DefaultExecutor> exec, ValueType *array,
 {
     exec->get_queue()->submit([&](sycl::handler &cgh) {
         cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx_id) {
-            const int idx = idx_id[0];
+            const auto idx = idx_id[0];
             array[idx] = val;
         });
     });

dpcpp/components/format_conversion.dp.hpp

@@ -56,10 +56,8 @@ namespace dpcpp {
 template <typename IndexType>
 inline void convert_unsorted_idxs_to_ptrs(const IndexType *idxs,
                                           size_type num_nonzeros,
-                                          IndexType *ptrs, size_type length)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+                                          IndexType *ptrs,
+                                          size_type length) GKO_NOT_IMPLEMENTED;
 
 
 /**
@@ -72,18 +70,12 @@ inline void convert_unsorted_idxs_to_ptrs(const IndexType *idxs,
 template <typename IndexType>
 inline void convert_sorted_idxs_to_ptrs(const IndexType *idxs,
                                         size_type num_nonzeros, IndexType *ptrs,
-                                        size_type length)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+                                        size_type length) GKO_NOT_IMPLEMENTED;
 
 
 template <typename IndexType>
 inline void convert_ptrs_to_idxs(const IndexType *ptrs, size_type num_rows,
-                                 IndexType *idxs)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+                                 IndexType *idxs) GKO_NOT_IMPLEMENTED;
 
 
 }  // namespace dpcpp

dpcpp/components/matrix_operations.dp.hpp

@@ -48,13 +48,9 @@ namespace dpcpp {
  * Computes the infinity norm of a column-major matrix.
  */
 template <typename ValueType>
-remove_complex<ValueType> compute_inf_norm(size_type num_rows,
-                                           size_type num_cols,
-                                           const ValueType *matrix,
-                                           size_type stride)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+remove_complex<ValueType> compute_inf_norm(
+    size_type num_rows, size_type num_cols, const ValueType *matrix,
+    size_type stride) GKO_NOT_IMPLEMENTED;
 
 
 }  // namespace dpcpp

dpcpp/components/precision_conversion.dp.cpp

@@ -48,7 +48,7 @@ void convert_precision(std::shared_ptr<const DefaultExecutor> exec,
 {
     exec->get_queue()->submit([&](sycl::handler &cgh) {
         cgh.parallel_for(sycl::range<1>{size}, [=](sycl::id<1> idx_id) {
-            const int idx = idx_id[0];
+            const auto idx = idx_id[0];
             out[idx] = in[idx];
         });
     });

dpcpp/components/prefix_sum.dp.cpp

@@ -47,13 +47,9 @@ namespace components {
 
 template <typename IndexType>
 void prefix_sum(std::shared_ptr<const DefaultExecutor> exec, IndexType *counts,
-                size_type num_entries)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+                size_type num_entries) GKO_NOT_IMPLEMENTED;
 
 GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE(GKO_DECLARE_PREFIX_SUM_KERNEL);
-
 // instantiate for size_type as well, as this is used in the Sellp format
 template GKO_DECLARE_PREFIX_SUM_KERNEL(size_type);
 

dpcpp/factorization/factorization_kernels.dp.cpp

@@ -59,29 +59,21 @@ namespace factorization {
 template <bool IsSorted, typename ValueType, typename IndexType>
 void find_missing_diagonal_elements(
     const matrix::Csr<ValueType, IndexType> *mtx,
-    IndexType *elements_to_add_per_row, bool *changes_required)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+    IndexType *elements_to_add_per_row,
+    bool *changes_required) GKO_NOT_IMPLEMENTED;
 
 
 template <typename ValueType, typename IndexType>
-void add_missing_diagonal_elements(const matrix::Csr<ValueType, IndexType> *mtx,
-                                   ValueType *new_values,
-                                   IndexType *new_col_idxs,
-                                   const IndexType *row_ptrs_addition)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+void add_missing_diagonal_elements(
+    const matrix::Csr<ValueType, IndexType> *mtx, ValueType *new_values,
+    IndexType *new_col_idxs,
+    const IndexType *row_ptrs_addition) GKO_NOT_IMPLEMENTED;
 
 
 template <typename ValueType, typename IndexType>
 void add_diagonal_elements(std::shared_ptr<const DpcppExecutor> exec,
                            matrix::Csr<ValueType, IndexType> *mtx,
-                           bool is_sorted)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+                           bool is_sorted) GKO_NOT_IMPLEMENTED;
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
     GKO_DECLARE_FACTORIZATION_ADD_DIAGONAL_ELEMENTS_KERNEL);
@@ -91,10 +83,7 @@ template <typename ValueType, typename IndexType>
 void initialize_row_ptrs_l_u(
     std::shared_ptr<const DpcppExecutor> exec,
     const matrix::Csr<ValueType, IndexType> *system_matrix,
-    IndexType *l_row_ptrs, IndexType *u_row_ptrs)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+    IndexType *l_row_ptrs, IndexType *u_row_ptrs) GKO_NOT_IMPLEMENTED;
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
     GKO_DECLARE_FACTORIZATION_INITIALIZE_ROW_PTRS_L_U_KERNEL);
@@ -105,9 +94,7 @@ void initialize_l_u(std::shared_ptr<const DpcppExecutor> exec,
                     const matrix::Csr<ValueType, IndexType> *system_matrix,
                     matrix::Csr<ValueType, IndexType> *csr_l,
                     matrix::Csr<ValueType, IndexType> *csr_u)
-{
     GKO_NOT_IMPLEMENTED;
-}
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
     GKO_DECLARE_FACTORIZATION_INITIALIZE_L_U_KERNEL);
@@ -117,10 +104,7 @@ template <typename ValueType, typename IndexType>
 void initialize_row_ptrs_l(
     std::shared_ptr<const DpcppExecutor> exec,
     const matrix::Csr<ValueType, IndexType> *system_matrix,
-    IndexType *l_row_ptrs)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+    IndexType *l_row_ptrs) GKO_NOT_IMPLEMENTED;
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
     GKO_DECLARE_FACTORIZATION_INITIALIZE_ROW_PTRS_L_KERNEL);
@@ -129,10 +113,8 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
 template <typename ValueType, typename IndexType>
 void initialize_l(std::shared_ptr<const DpcppExecutor> exec,
                   const matrix::Csr<ValueType, IndexType> *system_matrix,
-                  matrix::Csr<ValueType, IndexType> *csr_l, bool diag_sqrt)
-{
-    GKO_NOT_IMPLEMENTED;
-}
+                  matrix::Csr<ValueType, IndexType> *csr_l,
+                  bool diag_sqrt) GKO_NOT_IMPLEMENTED;
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
     GKO_DECLARE_FACTORIZATION_INITIALIZE_L_KERNEL);

dpcpp/factorization/par_ict_kernels.dp.cpp

@@ -48,7 +48,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "core/base/utils.hpp"
 #include "core/components/prefix_sum.hpp"
 #include "core/matrix/csr_builder.hpp"
-#include "dpcpp/components/csr_spgeam.dp.hpp"
 
 
 namespace gko {
@@ -67,9 +66,7 @@ void compute_factor(std::shared_ptr<const DefaultExecutor> exec,
                     const matrix::Csr<ValueType, IndexType> *a,
                     matrix::Csr<ValueType, IndexType> *l,
                     const matrix::Coo<ValueType, IndexType> *)
-{
     GKO_NOT_IMPLEMENTED;
-}
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
     GKO_DECLARE_PAR_ICT_COMPUTE_FACTOR_KERNEL);
@@ -81,9 +78,7 @@ void add_candidates(std::shared_ptr<const DefaultExecutor> exec,
                     const matrix::Csr<ValueType, IndexType> *a,
                     const matrix::Csr<ValueType, IndexType> *l,
                     matrix::Csr<ValueType, IndexType> *l_new)
-{
     GKO_NOT_IMPLEMENTED;
-}
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
     GKO_DECLARE_PAR_ICT_ADD_CANDIDATES_KERNEL);