Skip to content

Commit a08a8fc

Browse files
spencerpattyspencerpatty
committed
small changes from review
1 parent ea1a9b0 commit a08a8fc

File tree

2 files changed

+29
-32
lines changed

2 files changed

+29
-32
lines changed

‎Libraries/oneMKL/sparse_conjugate_gradient/sparse_cg.cpp

Lines changed: 17 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -333,35 +333,33 @@ int run_sparse_pcg_example(const sycl::device &dev)
333333
const intType nnz = ia_h[n]; // assumes zero indexing
334334

335335
// create arrays for help
336-
intType *ia_d = sycl::malloc_device<intType>(n+1, q);
337-
intType *ja_d = sycl::malloc_device<intType>(nnz, q);
338-
dataType *a_d = sycl::malloc_device<dataType>(nnz, q);
339-
dataType *x_d = sycl::malloc_device<dataType>(n, q);
340-
dataType *b_d = sycl::malloc_device<dataType>(n, q);
341-
dataType *r_d = sycl::malloc_device<dataType>(n, q); // r = b - A*x
342-
dataType *z_d = sycl::malloc_device<dataType>(n, q); // z = M^{-1} * r
343-
dataType *p_d = sycl::malloc_device<dataType>(n, q); // p = ?? z
344-
dataType *t_d = sycl::malloc_device<dataType>(n, q); // helper array
345-
dataType *d_d = sycl::malloc_device<dataType>(n, q); // diag(A)
346-
dataType *invd_d = sycl::malloc_device<dataType>(n, q); // inv(diag(A))
336+
intType *ia_d = sycl::malloc_device<intType>(n+1, q); // matrix rowptr
337+
intType *ja_d = sycl::malloc_device<intType>(nnz, q); // matrix columns
338+
dataType *a_d = sycl::malloc_device<dataType>(nnz, q); // matrix values
339+
dataType *x_d = sycl::malloc_device<dataType>(n, q); // solution
340+
dataType *b_d = sycl::malloc_device<dataType>(n, q); // right hand side
341+
dataType *r_d = sycl::malloc_device<dataType>(n, q); // residual
342+
dataType *z_d = sycl::malloc_device<dataType>(n, q); // preconditioned residualr
343+
dataType *p_d = sycl::malloc_device<dataType>(n, q); // search direction
344+
dataType *t_d = sycl::malloc_device<dataType>(n, q); // helper array
345+
dataType *d_d = sycl::malloc_device<dataType>(n, q); // matrix diagonals
346+
dataType *invd_d = sycl::malloc_device<dataType>(n, q); // matrix reciprocal of diagonals
347347

348348
const intType width = 8; // width * sizeof(dataType) >= cacheline size (64 Bytes)
349-
dataType *temp_d = sycl::malloc_device<dataType>(4*width, q);
350-
dataType *temp_h = sycl::malloc_host<dataType>(4*width, q);
349+
dataType *temp_d = sycl::malloc_device<dataType>(3*width, q);
350+
dataType *temp_h = sycl::malloc_host<dataType>(3*width, q);
351351

352352
if ( !ia_d || !ja_d || !a_d || !x_d || !b_d || !z_d || !p_d || !t_d || !d_d || !invd_d || !temp_d || !temp_h) {
353353
throw std::runtime_error("Failed to allocate device side USM memory");
354354
}
355355

356356
// device side aliases scattered by width elements each
357357
dataType *normr_h = temp_h;
358-
dataType *normr_d = temp_d;
359358
dataType *rtz_h = temp_h+1*width;
359+
dataType *pAp_h = temp_h+2*width;
360+
dataType *normr_d = temp_d;
360361
dataType *rtz_d = temp_d+1*width;
361-
//dataType *oldrtz_h = temp_h+2*width;
362-
//dataType *oldrtz_d = temp_d+2*width;
363-
dataType *pAp_h = temp_h+3*width;
364-
dataType *pAp_d = temp_d+3*width;
362+
dataType *pAp_d = temp_d+2*width;
365363

366364
// copy data from host to device arrays
367365
q.copy(ia_h, ia_d, n+1).wait();
@@ -372,6 +370,7 @@ int run_sparse_pcg_example(const sycl::device &dev)
372370

373371
extract_diagonal<dataType, intType>(q,n, ia_d, ja_d, a_d, d_d, invd_d, {}).wait();
374372

373+
// make the matrix diagonally dominant
375374
modify_diagonal<dataType, intType>(q, dataType(52.0), n, ia_d, ja_d, a_d, d_d, invd_d, {}).wait();
376375

377376
// create and initialize handle for a Sparse Matrix in CSR format

‎Libraries/oneMKL/sparse_conjugate_gradient/sparse_cg2.cpp

Lines changed: 12 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -407,17 +407,17 @@ int run_sparse_pcg_example(const sycl::device &dev)
407407
const intType nnz = ia_h[n]; // assumes zero indexing
408408

409409
// create arrays for help
410-
intType *ia_d = sycl::malloc_device<intType>(n+1, q);
411-
intType *ja_d = sycl::malloc_device<intType>(nnz, q);
412-
dataType *a_d = sycl::malloc_device<dataType>(nnz, q);
413-
dataType *x_d = sycl::malloc_device<dataType>(n, q);
414-
dataType *b_d = sycl::malloc_device<dataType>(n, q);
415-
dataType *r_d = sycl::malloc_device<dataType>(n, q); // r = b - A*x
416-
dataType *z_d = sycl::malloc_device<dataType>(n, q); // z = M^{-1} * r
417-
dataType *p_d = sycl::malloc_device<dataType>(n, q); // p = ?? z
418-
dataType *t_d = sycl::malloc_device<dataType>(n, q); // helper array
419-
dataType *d_d = sycl::malloc_device<dataType>(n, q); // diag(A)
420-
dataType *invd_d = sycl::malloc_device<dataType>(n, q); // inv(diag(A))
410+
intType *ia_d = sycl::malloc_device<intType>(n+1, q); // matrix rowptr
411+
intType *ja_d = sycl::malloc_device<intType>(nnz, q); // matrix columns
412+
dataType *a_d = sycl::malloc_device<dataType>(nnz, q); // matrix values
413+
dataType *x_d = sycl::malloc_device<dataType>(n, q); // solution
414+
dataType *b_d = sycl::malloc_device<dataType>(n, q); // right hand side
415+
dataType *r_d = sycl::malloc_device<dataType>(n, q); // residual
416+
dataType *z_d = sycl::malloc_device<dataType>(n, q); // preconditioned residualr
417+
dataType *p_d = sycl::malloc_device<dataType>(n, q); // search direction
418+
dataType *t_d = sycl::malloc_device<dataType>(n, q); // helper array
419+
dataType *d_d = sycl::malloc_device<dataType>(n, q); // matrix diagonals
420+
dataType *invd_d = sycl::malloc_device<dataType>(n, q); // matrix reciprocal of diagonals
421421

422422
const intType width = 8; // width * sizeof(dataType) >= cacheline size (64 Bytes)
423423
dataType *temp_d = sycl::malloc_device<dataType>(4*width, q);
@@ -430,11 +430,8 @@ int run_sparse_pcg_example(const sycl::device &dev)
430430
// device side aliases scattered by width elements each
431431
dataType *normr_h = temp_h;
432432
dataType *normr_d = temp_d;
433-
//dataType *rtz_h = temp_h+1*width;
434433
dataType *rtz_d = temp_d+1*width;
435-
//dataType *oldrtz_h = temp_h+2*width;
436434
dataType *oldrtz_d = temp_d+2*width;
437-
//dataType *pAp_h = temp_h+3*width;
438435
dataType *pAp_d = temp_d+3*width;
439436

440437
// copy data from host to device arrays
@@ -446,6 +443,7 @@ int run_sparse_pcg_example(const sycl::device &dev)
446443

447444
extract_diagonal<dataType, intType>(q,n, ia_d, ja_d, a_d, d_d, invd_d, {}).wait();
448445

446+
// make the matrix diagonally dominant
449447
modify_diagonal<dataType, intType>(q, dataType(52.0), n, ia_d, ja_d, a_d, d_d, invd_d, {}).wait();
450448

451449
// create and initialize handle for a Sparse Matrix in CSR format

0 commit comments

Comments
 (0)