add some note to indicate the porting TODO

need to revisit these TODO when we are close to fully porting ginkgo Co-authored-by: Terry Cojean <terry.cojean@kit.edu> Co-authored-by: Thomas Grützmacher <thomas.gruetzmacher@kit.edu> Co-authored-by: Tobias Ribizel <ribizel@kit.edu>
ginkgo-project · Jul 20, 2021 · 097f9ce · 097f9ce
1 parent 6169c2f
commit 097f9ce
Show file tree

Hide file tree

Showing 3 changed files with 32 additions and 24 deletions.
diff --git a/dpcpp/base/helper.hpp b/dpcpp/base/helper.hpp
@@ -56,17 +56,17 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
  * @param name_ the name of the host function with config
  * @param kernel_ the kernel name
  */
-#define GKO_ENABLE_DEFAULT_HOST(name_, kernel_) \
- template <typename... InferredArgs> \
- void name_(dim3 grid, dim3 block, size_t, sycl::queue *queue, \
- InferredArgs... args) \
- { \
- queue->submit([&](sycl::handler &cgh) { \
- cgh.parallel_for(sycl_nd_range(grid, block), \
- [=](sycl::nd_item<3> item_ct1) { \
- kernel_(args..., item_ct1); \
- }); \
- }); \
+#define GKO_ENABLE_DEFAULT_HOST(name_, kernel_)  \
+ template <typename... InferredArgs>  \
+ void name_(dim3 grid, dim3 block, size_type, sycl::queue *queue, \
+ InferredArgs... args)  \
+ {  \
+ queue->submit([&](sycl::handler &cgh) {  \
+ cgh.parallel_for(sycl_nd_range(grid, block),  \
+ [=](sycl::nd_item<3> item_ct1) {  \
+ kernel_(args..., item_ct1);  \
+ });  \
+ });  \
  }
 
 
@@ -78,17 +78,17 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
  * @param name_ the name of the host function with config
  * @param kernel_ the kernel name
  */
-#define GKO_ENABLE_DEFAULT_HOST_CONFIG(name_, kernel_) \
- template <std::uint32_t encoded, typename... InferredArgs> \
- inline void name_(dim3 grid, dim3 block, size_t, sycl::queue *queue, \
- InferredArgs... args) \
- { \
- queue->submit([&](sycl::handler &cgh) { \
- cgh.parallel_for(sycl_nd_range(grid, block), \
- [=](sycl::nd_item<3> item_ct1) { \
- kernel_<encoded>(args..., item_ct1); \
- }); \
- }); \
+#define GKO_ENABLE_DEFAULT_HOST_CONFIG(name_, kernel_)  \
+ template <std::uint32_t encoded, typename... InferredArgs>  \
+ inline void name_(dim3 grid, dim3 block, size_type, sycl::queue *queue, \
+ InferredArgs... args)  \
+ {  \
+ queue->submit([&](sycl::handler &cgh) {  \
+ cgh.parallel_for(sycl_nd_range(grid, block),  \
+ [=](sycl::nd_item<3> item_ct1) {  \
+ kernel_<encoded>(args..., item_ct1);  \
+ });  \
+ });  \
  }
 
 /**
@@ -106,7 +106,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #define GKO_ENABLE_DEFAULT_CONFIG_CALL(name_, callable_, list_) \
  template <typename... InferredArgs> \
  void name_(std::uint32_t desired_cfg, dim3 grid, dim3 block, \
- size_t dynamic_shared_memory, sycl::queue *queue,  \
+ size_type dynamic_shared_memory, sycl::queue *queue, \
  InferredArgs... args) \
  { \
  callable_( \
@@ -174,7 +174,6 @@ std::uint32_t get_first_cfg(IterArr &arr, Validate verify)
  }
  }
  GKO_NOT_SUPPORTED(arr);
- return 0;
 }
 
 

diff --git a/dpcpp/components/thread_ids.dp.hpp b/dpcpp/components/thread_ids.dp.hpp
@@ -52,6 +52,12 @@ namespace dpcpp {
 namespace thread {
 
 
+// TODO: porting - need to refine functions and their name in this file
+// the grid/block description uses the cuda dim3 to represent. i.e. using dim3
+// to launch dpcpp kernel, the kernel will reverse the ordering to keep the same
+// linear memory usage as cuda.
+
+
 /**
  * @internal
  *

diff --git a/dpcpp/components/uninitialized_array.hpp b/dpcpp/components/uninitialized_array.hpp
@@ -45,6 +45,9 @@ namespace kernels {
 namespace dpcpp {
 
 
+// TODO: porting - consider directly use the array as shared memory
+
+
 /**
  * Stores an array with uninitialized contents.
  *