| 476 | , rColIdx_(rColIdx) {} |
| 477 | |
| 478 | void operator()(sycl::nd_item<1> it) const { |
| 479 | common::Binary<T, op> binOP; |
| 480 | |
| 481 | const uint row = it.get_global_id(0); |
| 482 | |
| 483 | const bool valid = row < M_; |
| 484 | const uint lEnd = (valid ? lRowIdx_[row + 1] : 0); |
| 485 | const uint rEnd = (valid ? rRowIdx_[row + 1] : 0); |
| 486 | const uint offset = (valid ? oRowIdx_[row] : 0); |
| 487 | |
| 488 | T *ovPtr = oVals_.get_pointer() + offset; |
| 489 | int *ocPtr = oColIdx_.get_pointer() + offset; |
| 490 | |
| 491 | uint l = (valid ? lRowIdx_[row] : 0); |
| 492 | uint r = (valid ? rRowIdx_[row] : 0); |
| 493 | |
| 494 | uint nnz = 0; |
| 495 | while (l < lEnd && r < rEnd) { |
| 496 | uint lci = lColIdx_[l]; |
| 497 | uint rci = rColIdx_[r]; |
| 498 | |
| 499 | T lhs = (lci <= rci ? lVals_[l] : common::Binary<T, op>::init()); |
| 500 | T rhs = (lci >= rci ? rVals_[r] : common::Binary<T, op>::init()); |
| 501 | |
| 502 | ovPtr[nnz] = binOP(lhs, rhs); |
| 503 | ocPtr[nnz] = (lci <= rci) ? lci : rci; |
| 504 | |
| 505 | l += (lci <= rci); |
| 506 | r += (lci >= rci); |
| 507 | nnz++; |
| 508 | } |
| 509 | while (l < lEnd) { |
| 510 | ovPtr[nnz] = binOP(lVals_[l], common::Binary<T, op>::init()); |
| 511 | ocPtr[nnz] = lColIdx_[l]; |
| 512 | l++; |
| 513 | nnz++; |
| 514 | } |
| 515 | while (r < rEnd) { |
| 516 | ovPtr[nnz] = binOP(common::Binary<T, op>::init(), rVals_[r]); |
| 517 | ocPtr[nnz] = rColIdx_[r]; |
| 518 | r++; |
| 519 | nnz++; |
| 520 | } |
| 521 | } |
| 522 | |
| 523 | private: |
| 524 | write_accessor<T> oVals_; |