Initial phase of integrating new BL:AS 2 subroutines from merge_reduction… #61

mehdi-goli · 2018-08-07T15:09:38Z

… branch.

…rameters.

AdamBrouwersHarries · 2018-08-13T13:29:17Z

include/interface/blas2_interface.hpp

  if (mA.getAccess()) {
-    auto scalOp1 = make_op<ScalarOp, prdOp2_struct>(_beta, vy);


Could you explain what this change is here? As far as I can tell, it's cleaning up the code to pick one implementation, but I could be very wrong.

Yes, you are right. The main backend code for SYCL BLAS GEMV has changed so now whether it is row-major or column major, there are two separate computation. However, the make_op is the same for both, it has been moved affter if/else.

AdamBrouwersHarries · 2018-08-13T13:56:15Z

include/operations/blas2_trees.hpp

-  value_type eval(IndexType i) {
+  bool valid_thread(cl::sycl::nd_item<1> ndItem) { return true; }
+
+  value_type eval(IndexType i) {  // NOT VERIFIED


What does "NOT VERIFIED" mean in this context?

I have added todo list for @josealiaga to check if this function is not used and extra, so he can let us know and we can remove it from here.

I think the method value_type eval(IndexType i) { .. } has to be defined in all structs.
In this case, it only implements the most simple implementation of the matrix-vector product by rows

AdamBrouwersHarries · 2018-08-13T14:01:16Z

include/queue/queue_base.hpp

+  */
+  inline bool has_local_memory() { return false; }
+  /*
+   @brief This class is used to allocated the a regin of memory on the device


s/the a regin/a region/

AdamBrouwersHarries · 2018-08-13T14:01:52Z

include/queue/queue_base.hpp

+class Queue_Interface<Sequential> {
+  Queue_Interface() {}
+  /*
+  @brief This class is to determine whether or not the underlying device has


This class? (same for other member functions).

Ruyk

There are various code style and documentation concerns. In general I think is great we have now level 2 implementation (thanks @josealiaga !) but we need now to refactor the codebase to make the code a bit clearer. We can do that on other pull request.

Ruyk · 2018-08-15T14:25:26Z

include/executors/executor_sycl.hpp

-                                                   // to the tree not the root
-      //  printf("Index %ld\n", index.get_global_id(0));
+    if (tree.valid_thread(index)) {  // FIXME (Mehdi):: this should move
+                                     // to the tree not the root


Not sure if I understand the FIXME, can you create a gihub issue and describe it more clearly?

Ruyk · 2018-08-15T14:25:43Z

include/executors/executor_sycl.hpp

@@ -263,7 +262,9 @@ class Executor<SYCL> {
   */
  Executor(cl::sycl::queue q) : q_interface(q){};

-  cl::sycl::queue sycl_queue() const { return q_interface.sycl_queue(); }
+  inline Queue_Interface<SYCL> &policy_handler() { return q_interface; }


why do you return a reference to the policy?

Because the queue interface contains the pointer mapper class which deals with allocation and deallocation of buffers and memory. The call to
implicitly-deleted copy constructor of 'Queue_Interface' error will rise in case of not returning the reference

here is the error :
'Queue_Interface' is implicitly deleted because field 'pointer_mapper' has a deleted copy
constructor
mutable cl::sycl::codeplay::PointerMapper pointer_mapper;
^
../../../sycl-blas/include/queue/pointer_mapper.hpp:310:3: note: 'PointerMapper' has been
explicitly marked deleted here
PointerMapper(const PointerMapper&) = delete;

used std::shared_ptr to remove the reference

Ruyk · 2018-08-15T14:26:27Z

include/executors/executor_sycl.hpp

@@ -323,7 +324,7 @@ class Executor<SYCL> {
  */
  template <typename T>
  inline void copy_to_device(T *src, buffer_iterator<T> dst, size_t = 0) {
-    sycl_queue().submit([&](cl::sycl::handler &cgh) {
+    queue().submit([&](cl::sycl::handler &cgh) {


I think is clearer if you rename this to "get_queue" so it is not confused with a constructor

Ruyk · 2018-08-15T14:27:15Z

include/interface/blas1_interface.hpp

-  auto localSize = 256;
-  auto nWG = 512;
+
+  auto localSize = ex.policy_handler().get_work_group_size();


const auto localSize?

Ruyk · 2018-08-15T14:27:29Z

include/interface/blas1_interface.hpp

-  auto nWG = 512;
+
+  auto localSize = ex.policy_handler().get_work_group_size();
+  auto nWG = 2 * localSize;


const auto nWG? Also, why 2? can you ellaborate?

fixed. @josealiaga could you elaborate why this is fixed for 2?

Please, first, read the attached document: SYCL_BLAS1_reduction.txt

If you analyze my original code, you will see that two alternatives are considered, related to the alternatives A and B in the attached document. In the merged code, only the case B appears, whereas the case A is the best one for OpenCL devices that I tested.

Ruyk · 2018-08-15T14:52:44Z

include/operations/blas2_trees.hpp

+        }
+      }
+    }
+    return l.eval(frs_row, idWFC);


there is a lot of duplicated code between the different variants of gemv, cannot this be shared on a separate struct?
Maybe it should be a separate issue that is smaller than this branch

Ruyk · 2018-08-15T14:53:01Z

include/operations/blas2_trees.hpp

-    l.eval(i) = scl * val + r2.eval(i);
-    return val;
-  }
+  IndexType getSize() { return r1.getSize(); }


can be const everywhere

Ruyk · 2018-08-15T14:54:15Z

include/operations/blas2_trees.hpp

+        if (rowid > frs_row)
+          // This barrier is mandatory to be sure the data is on the shared
+          // memory
+          ndItem.barrier(cl::sycl::access::fence_space::local_space);


how do we ensure all threads go through this barrier?
The decision of thread cancelation is not in this part of the code so its difficult to understand if this is correct

This code should be restructured and re-written and the barrier should be out side of the if statement. I have created a issue for it

Ruyk · 2018-08-15T14:56:09Z

include/operations/blas2_trees.hpp

+    // TO SOLVE IT, USE GLOBAL VALUES OF frs_col AND lst_col
+    if ((!Upper && (((idWFC * dimWFC) + ((!Diag) ? 1 : 0)) > (lst_row - 1))) ||
+        (!Lower &&
+         ((frs_row + ((!Diag) ? 1 : 0)) > ((idWFC * dimWFC + dimWFC) - 1)))) {


can we explain what this condition means?
Extract the boolean decisions outside of the if, and use const bool temp variables to explain each step, otherwise its difficult to understand.

@josealiaga could you please add some documentation here.

Some BLAS routines (GER, SYR and SYR2) are computed by using the same nodes (Ger_Row and Ger_col), but changing the Single, Lower, Upper and Diag parameters.
If only one of the triangle of the matrix has to be computed (SYR and SYR2), the function verifies when the computation has to be made.
I don't remember if Diag can be false in some case, but I probably consider the more general case, and I also included it as parameter

Ruyk · 2018-08-15T14:59:05Z

include/queue/queue_sycl.hpp

+    wGSize |= (wGSize >> 4);
+    wGSize |= (wGSize >> 8);
+    wGSize |= (wGSize >> 16);
+#if defined(__x86_64__) || defined(_M_X64) || defined(__amd64) || \


extract that set of macros onto a separate hader and say IS_64BIT or something like that

We are creating a syclblas.hpp file and this macro will be added there, in aseparate PR.

initial phase of integrating new blas 2 function from merge_reduction…

c0073dc

… branch.

mehdi-goli requested review from Ruyk and AdamBrouwersHarries August 7, 2018 15:09

mehdi-goli added the WIP label Aug 7, 2018

mehdi-goli changed the title ~~Initial phase of integrating new BL:AS 2 subroutines from merge_reduction…~~ [WIP] Initial phase of integrating new BL:AS 2 subroutines from merge_reduction… Aug 7, 2018

mehdi-goli added 6 commits August 9, 2018 10:39

Integrating merge_reduction branch with new SYCL interface.

2243be3

Formatting the code,

51dfd12

Adding the valid thread from merge_reduction btanch.

2ce2e0d

Replacing std::cout with std::invalid_arguments for the eErroneous pa…

b634eb5

…rameters.

Replacing std::cout with std::invalid_arguments for the eErroneous pa…

36e5fd9

…rameters.

set the local size to be taken from device.

171f65a

mehdi-goli removed the WIP label Aug 10, 2018

mehdi-goli changed the title ~~[WIP] Initial phase of integrating new BL:AS 2 subroutines from merge_reduction…~~ Initial phase of integrating new BL:AS 2 subroutines from merge_reduction… Aug 10, 2018

AdamBrouwersHarries suggested changes Aug 13, 2018

View reviewed changes

Applying the comments.

32bf6f4

AdamBrouwersHarries approved these changes Aug 14, 2018

View reviewed changes

Ruyk approved these changes Aug 15, 2018

View reviewed changes

mehdi-goli added 5 commits August 16, 2018 14:44

Applying the comments.

0eb80bd

removing commented codes.

b688383

applying the clang format.

c2d8691

replacing the _ from the optimisation parameters.

0617da1

Merge with upstream. resolving the conflict.

b2b9e8f

mehdi-goli merged commit b6329fc into codeplaysoftware:master Aug 17, 2018

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Initial phase of integrating new BL:AS 2 subroutines from merge_reduction… #61

Initial phase of integrating new BL:AS 2 subroutines from merge_reduction… #61

mehdi-goli commented Aug 7, 2018

AdamBrouwersHarries Aug 13, 2018

mehdi-goli Aug 13, 2018

AdamBrouwersHarries Aug 13, 2018

mehdi-goli Aug 13, 2018

josealiaga Aug 30, 2018

AdamBrouwersHarries Aug 13, 2018

mehdi-goli Aug 13, 2018

AdamBrouwersHarries Aug 13, 2018

mehdi-goli Aug 13, 2018

Ruyk left a comment

Ruyk Aug 15, 2018

Ruyk Aug 15, 2018

mehdi-goli Aug 16, 2018

mehdi-goli Aug 16, 2018 •

edited

Loading

mehdi-goli Aug 16, 2018

Ruyk Aug 15, 2018

mehdi-goli Aug 16, 2018

Ruyk Aug 15, 2018

mehdi-goli Aug 16, 2018

Ruyk Aug 15, 2018

mehdi-goli Aug 16, 2018

josealiaga Aug 22, 2018

Ruyk Aug 15, 2018

mehdi-goli Aug 16, 2018

Ruyk Aug 15, 2018

mehdi-goli Aug 16, 2018

Ruyk Aug 15, 2018

mehdi-goli Aug 16, 2018

Ruyk Aug 15, 2018

mehdi-goli Aug 16, 2018

josealiaga Aug 30, 2018

Ruyk Aug 15, 2018

mehdi-goli Aug 16, 2018

		if (mA.getAccess()) {
		auto scalOp1 = make_op<ScalarOp, prdOp2_struct>(_beta, vy);

Initial phase of integrating new BL:AS 2 subroutines from merge_reduction… #61

Initial phase of integrating new BL:AS 2 subroutines from merge_reduction… #61

Conversation

mehdi-goli commented Aug 7, 2018

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Ruyk left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

mehdi-goli Aug 16, 2018 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

mehdi-goli Aug 16, 2018 •

edited

Loading