1 // Copyright (c) 2017-2018, Lawrence Livermore National Security, LLC. 2 // Produced at the Lawrence Livermore National Laboratory. LLNL-CODE-734707. 3 // All Rights reserved. See files LICENSE and NOTICE for details. 4 // 5 // This file is part of CEED, a collection of benchmarks, miniapps, software 6 // libraries and APIs for efficient high-order finite element and spectral 7 // element discretizations for exascale applications. For more information and 8 // source code availability see http://github.com/ceed. 9 // 10 // The CEED research is supported by the Exascale Computing Project 17-SC-20-SC, 11 // a collaborative effort of two U.S. Department of Energy organizations (Office 12 // of Science and the National Nuclear Security Administration) responsible for 13 // the planning and preparation of a capable exascale ecosystem, including 14 // software, applications, hardware, advanced system engineering and early 15 // testbed platforms, in support of the nation's exascale computing imperative. 16 17 #include <ceed/ceed.h> 18 #include <ceed/backend.h> 19 #include <hip/hip_runtime.h> 20 #include <hipblas.h> 21 #include <math.h> 22 #include <string.h> 23 #include "ceed-hip-ref.h" 24 25 //------------------------------------------------------------------------------ 26 // * Bytes used 27 //------------------------------------------------------------------------------ 28 static inline size_t bytes(const CeedVector vec) { 29 int ierr; 30 CeedInt length; 31 ierr = CeedVectorGetLength(vec, &length); CeedChkBackend(ierr); 32 return length * sizeof(CeedScalar); 33 } 34 35 //------------------------------------------------------------------------------ 36 // Sync host to device 37 //------------------------------------------------------------------------------ 38 static inline int CeedVectorSyncH2D_Hip(const CeedVector vec) { 39 int ierr; 40 Ceed ceed; 41 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 42 CeedVector_Hip *impl; 43 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 44 45 if (!impl->h_array) 46 // LCOV_EXCL_START 47 return CeedError(ceed, CEED_ERROR_BACKEND, 48 "No valid host data to sync to device"); 49 // LCOV_EXCL_STOP 50 51 if (impl->d_array_borrowed) { 52 impl->d_array = impl->d_array_borrowed; 53 } else if (impl->d_array_owned) { 54 impl->d_array = impl->d_array_owned; 55 } else { 56 ierr = hipMalloc((void **)&impl->d_array_owned, bytes(vec)); 57 CeedChk_Hip(ceed, ierr); 58 impl->d_array = impl->d_array_owned; 59 } 60 61 ierr = hipMemcpy(impl->d_array, impl->h_array, bytes(vec), 62 hipMemcpyHostToDevice); CeedChk_Hip(ceed, ierr); 63 64 return CEED_ERROR_SUCCESS; 65 } 66 67 //------------------------------------------------------------------------------ 68 // Sync device to host 69 //------------------------------------------------------------------------------ 70 static inline int CeedVectorSyncD2H_Hip(const CeedVector vec) { 71 int ierr; 72 Ceed ceed; 73 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 74 CeedVector_Hip *impl; 75 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 76 77 if (!impl->d_array) 78 // LCOV_EXCL_START 79 return CeedError(ceed, CEED_ERROR_BACKEND, 80 "No valid device data to sync to host"); 81 // LCOV_EXCL_STOP 82 83 if (impl->h_array_borrowed) { 84 impl->h_array = impl->h_array_borrowed; 85 } else if (impl->h_array_owned) { 86 impl->h_array = impl->h_array_owned; 87 } else { 88 CeedInt length; 89 ierr = CeedVectorGetLength(vec, &length); CeedChkBackend(ierr); 90 ierr = CeedCalloc(length, &impl->h_array_owned); CeedChkBackend(ierr); 91 impl->h_array = impl->h_array_owned; 92 } 93 94 ierr = hipMemcpy(impl->h_array, impl->d_array, bytes(vec), 95 hipMemcpyDeviceToHost); CeedChk_Hip(ceed, ierr); 96 97 return CEED_ERROR_SUCCESS; 98 } 99 100 //------------------------------------------------------------------------------ 101 // Sync arrays 102 //------------------------------------------------------------------------------ 103 static inline int CeedVectorSync_Hip(const CeedVector vec, CeedMemType mtype) { 104 switch (mtype) { 105 case CEED_MEM_HOST: return CeedVectorSyncD2H_Hip(vec); 106 case CEED_MEM_DEVICE: return CeedVectorSyncH2D_Hip(vec); 107 } 108 return CEED_ERROR_UNSUPPORTED; 109 } 110 111 //------------------------------------------------------------------------------ 112 // Set all pointers as invalid 113 //------------------------------------------------------------------------------ 114 static inline int CeedVectorSetAllInvalid_Hip(const CeedVector vec) { 115 int ierr; 116 CeedVector_Hip *impl; 117 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 118 119 impl->h_array = NULL; 120 impl->d_array = NULL; 121 122 return CEED_ERROR_SUCCESS; 123 } 124 125 //------------------------------------------------------------------------------ 126 // Check if CeedVector has any valid pointers 127 //------------------------------------------------------------------------------ 128 static inline int CeedVectorHasValidArray_Hip(const CeedVector vec, 129 bool *has_valid_array) { 130 int ierr; 131 CeedVector_Hip *impl; 132 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 133 134 *has_valid_array = !!impl->h_array || !!impl->d_array; 135 136 return CEED_ERROR_SUCCESS; 137 } 138 139 //------------------------------------------------------------------------------ 140 // Check if has any array of given type 141 //------------------------------------------------------------------------------ 142 static inline int CeedVectorHasArrayOfType_Hip(const CeedVector vec, 143 CeedMemType mtype, bool *has_array_of_type) { 144 int ierr; 145 CeedVector_Hip *impl; 146 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 147 148 switch (mtype) { 149 case CEED_MEM_HOST: 150 *has_array_of_type = !!impl->h_array_borrowed || !!impl->h_array_owned; 151 break; 152 case CEED_MEM_DEVICE: 153 *has_array_of_type = !!impl->d_array_borrowed || !!impl->d_array_owned; 154 break; 155 } 156 157 return CEED_ERROR_SUCCESS; 158 } 159 160 //------------------------------------------------------------------------------ 161 // Check if has borrowed array of given type 162 //------------------------------------------------------------------------------ 163 static inline int CeedVectorHasBorrowedArrayOfType_Hip(const CeedVector vec, 164 CeedMemType mtype, bool *has_borrowed_array_of_type) { 165 int ierr; 166 CeedVector_Hip *impl; 167 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 168 169 switch (mtype) { 170 case CEED_MEM_HOST: 171 *has_borrowed_array_of_type = !!impl->h_array_borrowed; 172 break; 173 case CEED_MEM_DEVICE: 174 *has_borrowed_array_of_type = !!impl->d_array_borrowed; 175 break; 176 } 177 178 return CEED_ERROR_SUCCESS; 179 } 180 181 //------------------------------------------------------------------------------ 182 // Sync array of given type 183 //------------------------------------------------------------------------------ 184 static inline int CeedVectorNeedSync_Hip(const CeedVector vec, 185 CeedMemType mtype, bool *need_sync) { 186 int ierr; 187 CeedVector_Hip *impl; 188 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 189 190 bool has_valid_array = false; 191 ierr = CeedVectorHasValidArray(vec, &has_valid_array); CeedChkBackend(ierr); 192 switch (mtype) { 193 case CEED_MEM_HOST: 194 *need_sync = has_valid_array && !impl->h_array; 195 break; 196 case CEED_MEM_DEVICE: 197 *need_sync = has_valid_array && !impl->d_array; 198 break; 199 } 200 201 return CEED_ERROR_SUCCESS; 202 } 203 204 //------------------------------------------------------------------------------ 205 // Set array from host 206 //------------------------------------------------------------------------------ 207 static int CeedVectorSetArrayHost_Hip(const CeedVector vec, 208 const CeedCopyMode cmode, CeedScalar *array) { 209 int ierr; 210 CeedVector_Hip *impl; 211 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 212 213 switch (cmode) { 214 case CEED_COPY_VALUES: { 215 CeedInt length; 216 if (!impl->h_array_owned) { 217 ierr = CeedVectorGetLength(vec, &length); CeedChkBackend(ierr); 218 ierr = CeedMalloc(length, &impl->h_array_owned); CeedChkBackend(ierr); 219 } 220 impl->h_array_borrowed = NULL; 221 impl->h_array = impl->h_array_owned; 222 if (array) 223 memcpy(impl->h_array, array, bytes(vec)); 224 } break; 225 case CEED_OWN_POINTER: 226 ierr = CeedFree(&impl->h_array_owned); CeedChkBackend(ierr); 227 impl->h_array_owned = array; 228 impl->h_array_borrowed = NULL; 229 impl->h_array = array; 230 break; 231 case CEED_USE_POINTER: 232 ierr = CeedFree(&impl->h_array_owned); CeedChkBackend(ierr); 233 impl->h_array_borrowed = array; 234 impl->h_array = array; 235 break; 236 } 237 238 return CEED_ERROR_SUCCESS; 239 } 240 241 //------------------------------------------------------------------------------ 242 // Set array from device 243 //------------------------------------------------------------------------------ 244 static int CeedVectorSetArrayDevice_Hip(const CeedVector vec, 245 const CeedCopyMode cmode, CeedScalar *array) { 246 int ierr; 247 Ceed ceed; 248 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 249 CeedVector_Hip *impl; 250 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 251 252 switch (cmode) { 253 case CEED_COPY_VALUES: 254 if (!impl->d_array_owned) { 255 ierr = hipMalloc((void **)&impl->d_array_owned, bytes(vec)); 256 CeedChk_Hip(ceed, ierr); 257 } 258 impl->d_array_borrowed = NULL; 259 impl->d_array = impl->d_array_owned; 260 if (array) { 261 ierr = hipMemcpy(impl->d_array, array, bytes(vec), 262 hipMemcpyDeviceToDevice); CeedChk_Hip(ceed, ierr); 263 } 264 break; 265 case CEED_OWN_POINTER: 266 ierr = hipFree(impl->d_array_owned); CeedChk_Hip(ceed, ierr); 267 impl->d_array_owned = array; 268 impl->d_array_borrowed = NULL; 269 impl->d_array = array; 270 break; 271 case CEED_USE_POINTER: 272 ierr = hipFree(impl->d_array_owned); CeedChk_Hip(ceed, ierr); 273 impl->d_array_owned = NULL; 274 impl->d_array_borrowed = array; 275 impl->d_array = array; 276 break; 277 } 278 279 return CEED_ERROR_SUCCESS; 280 } 281 282 //------------------------------------------------------------------------------ 283 // Set the array used by a vector, 284 // freeing any previously allocated array if applicable 285 //------------------------------------------------------------------------------ 286 static int CeedVectorSetArray_Hip(const CeedVector vec, const CeedMemType mtype, 287 const CeedCopyMode cmode, CeedScalar *array) { 288 int ierr; 289 Ceed ceed; 290 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 291 CeedVector_Hip *impl; 292 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 293 294 ierr = CeedVectorSetAllInvalid_Hip(vec); CeedChkBackend(ierr); 295 switch (mtype) { 296 case CEED_MEM_HOST: 297 return CeedVectorSetArrayHost_Hip(vec, cmode, array); 298 case CEED_MEM_DEVICE: 299 return CeedVectorSetArrayDevice_Hip(vec, cmode, array); 300 } 301 302 return CEED_ERROR_UNSUPPORTED; 303 } 304 305 //------------------------------------------------------------------------------ 306 // Set host array to value 307 //------------------------------------------------------------------------------ 308 static int CeedHostSetValue_Hip(CeedScalar *h_array, CeedInt length, 309 CeedScalar val) { 310 for (int i = 0; i < length; i++) 311 h_array[i] = val; 312 return CEED_ERROR_SUCCESS; 313 } 314 315 //------------------------------------------------------------------------------ 316 // Set device array to value (impl in .hip file) 317 //------------------------------------------------------------------------------ 318 int CeedDeviceSetValue_Hip(CeedScalar *d_array, CeedInt length, CeedScalar val); 319 320 //------------------------------------------------------------------------------ 321 // Set a vector to a value, 322 //------------------------------------------------------------------------------ 323 static int CeedVectorSetValue_Hip(CeedVector vec, CeedScalar val) { 324 int ierr; 325 Ceed ceed; 326 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 327 CeedVector_Hip *impl; 328 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 329 CeedInt length; 330 ierr = CeedVectorGetLength(vec, &length); CeedChkBackend(ierr); 331 332 // Set value for synced device/host array 333 if (!impl->d_array && !impl->h_array) { 334 if (impl->d_array_borrowed) { 335 impl->d_array = impl->d_array_borrowed; 336 } else if (impl->h_array_borrowed) { 337 impl->h_array = impl->h_array_borrowed; 338 } else if (impl->d_array_owned) { 339 impl->d_array = impl->d_array_owned; 340 } else if (impl->h_array_owned) { 341 impl->h_array = impl->h_array_owned; 342 } else { 343 ierr = CeedVectorSetArray(vec, CEED_MEM_DEVICE, CEED_COPY_VALUES, NULL); 344 CeedChkBackend(ierr); 345 } 346 } 347 if (impl->d_array) { 348 ierr = CeedDeviceSetValue_Hip(impl->d_array, length, val); CeedChkBackend(ierr); 349 } 350 if (impl->h_array) { 351 ierr = CeedHostSetValue_Hip(impl->h_array, length, val); CeedChkBackend(ierr); 352 } 353 354 return CEED_ERROR_SUCCESS; 355 } 356 357 //------------------------------------------------------------------------------ 358 // Vector Take Array 359 //------------------------------------------------------------------------------ 360 static int CeedVectorTakeArray_Hip(CeedVector vec, CeedMemType mtype, 361 CeedScalar **array) { 362 int ierr; 363 Ceed ceed; 364 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 365 CeedVector_Hip *impl; 366 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 367 368 // Sync array to requested memtype 369 bool need_sync = false; 370 ierr = CeedVectorNeedSync_Hip(vec, mtype, &need_sync); CeedChkBackend(ierr); 371 if (need_sync) { 372 ierr = CeedVectorSync_Hip(vec, mtype); CeedChkBackend(ierr); 373 } 374 375 // Update pointer 376 switch (mtype) { 377 case CEED_MEM_HOST: 378 (*array) = impl->h_array_borrowed; 379 impl->h_array_borrowed = NULL; 380 impl->h_array = NULL; 381 break; 382 case CEED_MEM_DEVICE: 383 (*array) = impl->d_array_borrowed; 384 impl->d_array_borrowed = NULL; 385 impl->d_array = NULL; 386 break; 387 } 388 389 return CEED_ERROR_SUCCESS; 390 } 391 392 //------------------------------------------------------------------------------ 393 // Core logic for array syncronization for GetArray. 394 // If a different memory type is most up to date, this will perform a copy 395 //------------------------------------------------------------------------------ 396 static int CeedVectorGetArrayCore_Hip(const CeedVector vec, 397 const CeedMemType mtype, CeedScalar **array) { 398 int ierr; 399 Ceed ceed; 400 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 401 CeedVector_Hip *impl; 402 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 403 404 bool need_sync = false; 405 ierr = CeedVectorNeedSync_Hip(vec, mtype, &need_sync); CeedChkBackend(ierr); 406 CeedChkBackend(ierr); 407 if (need_sync) { 408 // Sync array to requested memtype 409 ierr = CeedVectorSync_Hip(vec, mtype); CeedChkBackend(ierr); 410 } 411 412 // Update pointer 413 switch (mtype) { 414 case CEED_MEM_HOST: 415 *array = impl->h_array; 416 break; 417 case CEED_MEM_DEVICE: 418 *array = impl->d_array; 419 break; 420 } 421 422 return CEED_ERROR_SUCCESS; 423 } 424 425 //------------------------------------------------------------------------------ 426 // Get read-only access to a vector via the specified mtype 427 //------------------------------------------------------------------------------ 428 static int CeedVectorGetArrayRead_Hip(const CeedVector vec, 429 const CeedMemType mtype, const CeedScalar **array) { 430 return CeedVectorGetArrayCore_Hip(vec, mtype, (CeedScalar **)array); 431 } 432 433 //------------------------------------------------------------------------------ 434 // Get read/write access to a vector via the specified mtype 435 //------------------------------------------------------------------------------ 436 static int CeedVectorGetArray_Hip(const CeedVector vec, const CeedMemType mtype, 437 CeedScalar **array) { 438 int ierr; 439 CeedVector_Hip *impl; 440 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 441 442 ierr = CeedVectorGetArrayCore_Hip(vec, mtype, array); CeedChkBackend(ierr); 443 444 ierr = CeedVectorSetAllInvalid_Hip(vec); CeedChkBackend(ierr); 445 switch (mtype) { 446 case CEED_MEM_HOST: 447 impl->h_array = *array; 448 break; 449 case CEED_MEM_DEVICE: 450 impl->d_array = *array; 451 break; 452 } 453 454 return CEED_ERROR_SUCCESS; 455 } 456 457 //------------------------------------------------------------------------------ 458 // Get write access to a vector via the specified mtype 459 //------------------------------------------------------------------------------ 460 static int CeedVectorGetArrayWrite_Hip(const CeedVector vec, 461 const CeedMemType mtype, CeedScalar **array) { 462 int ierr; 463 CeedVector_Hip *impl; 464 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 465 466 bool has_array_of_type = true; 467 ierr = CeedVectorHasArrayOfType_Hip(vec, mtype, &has_array_of_type); 468 CeedChkBackend(ierr); 469 if (!has_array_of_type) { 470 // Allocate if array is not yet allocated 471 ierr = CeedVectorSetArray(vec, mtype, CEED_COPY_VALUES, NULL); 472 CeedChkBackend(ierr); 473 } else { 474 // Select dirty array 475 switch (mtype) { 476 case CEED_MEM_HOST: 477 if (impl->h_array_borrowed) 478 impl->h_array = impl->h_array_borrowed; 479 else 480 impl->h_array = impl->h_array_owned; 481 break; 482 case CEED_MEM_DEVICE: 483 if (impl->d_array_borrowed) 484 impl->d_array = impl->d_array_borrowed; 485 else 486 impl->d_array = impl->d_array_owned; 487 } 488 } 489 490 return CeedVectorGetArray_Hip(vec, mtype, array); 491 } 492 493 //------------------------------------------------------------------------------ 494 // Restore an array obtained using CeedVectorGetArrayRead() 495 //------------------------------------------------------------------------------ 496 static int CeedVectorRestoreArrayRead_Hip(const CeedVector vec) { 497 return CEED_ERROR_SUCCESS; 498 } 499 500 //------------------------------------------------------------------------------ 501 // Restore an array obtained using CeedVectorGetArray() 502 //------------------------------------------------------------------------------ 503 static int CeedVectorRestoreArray_Hip(const CeedVector vec) { 504 return CEED_ERROR_SUCCESS; 505 } 506 507 //------------------------------------------------------------------------------ 508 // Get the norm of a CeedVector 509 //------------------------------------------------------------------------------ 510 static int CeedVectorNorm_Hip(CeedVector vec, CeedNormType type, 511 CeedScalar *norm) { 512 int ierr; 513 Ceed ceed; 514 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 515 CeedVector_Hip *impl; 516 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 517 CeedInt length; 518 ierr = CeedVectorGetLength(vec, &length); CeedChkBackend(ierr); 519 hipblasHandle_t handle; 520 ierr = CeedHipGetHipblasHandle(ceed, &handle); CeedChkBackend(ierr); 521 522 // Compute norm 523 const CeedScalar *d_array; 524 ierr = CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &d_array); 525 CeedChkBackend(ierr); 526 switch (type) { 527 case CEED_NORM_1: { 528 if (CEED_SCALAR_TYPE == CEED_SCALAR_FP32) { 529 ierr = hipblasSasum(handle, length, (float *) d_array, 1, (float *) norm); 530 } else { 531 ierr = hipblasDasum(handle, length, (double *) d_array, 1, (double *) norm); 532 } 533 CeedChk_Hipblas(ceed, ierr); 534 break; 535 } 536 case CEED_NORM_2: { 537 if (CEED_SCALAR_TYPE == CEED_SCALAR_FP32) { 538 ierr = hipblasSnrm2(handle, length, (float *) d_array, 1, (float *) norm); 539 } else { 540 ierr = hipblasDnrm2(handle, length, (double *) d_array, 1, (double *) norm); 541 } 542 CeedChk_Hipblas(ceed, ierr); 543 break; 544 } 545 case CEED_NORM_MAX: { 546 CeedInt indx; 547 if (CEED_SCALAR_TYPE == CEED_SCALAR_FP32) { 548 ierr = hipblasIsamax(handle, length, (float *) d_array, 1, &indx); 549 } else { 550 ierr = hipblasIdamax(handle, length, (double *) d_array, 1, &indx); 551 } 552 CeedChk_Hipblas(ceed, ierr); 553 CeedScalar normNoAbs; 554 ierr = hipMemcpy(&normNoAbs, impl->d_array+indx-1, sizeof(CeedScalar), 555 hipMemcpyDeviceToHost); CeedChk_Hip(ceed, ierr); 556 *norm = fabs(normNoAbs); 557 break; 558 } 559 } 560 ierr = CeedVectorRestoreArrayRead(vec, &d_array); CeedChkBackend(ierr); 561 562 return CEED_ERROR_SUCCESS; 563 } 564 565 //------------------------------------------------------------------------------ 566 // Take reciprocal of a vector on host 567 //------------------------------------------------------------------------------ 568 static int CeedHostReciprocal_Hip(CeedScalar *h_array, CeedInt length) { 569 for (int i = 0; i < length; i++) 570 if (fabs(h_array[i]) > CEED_EPSILON) 571 h_array[i] = 1./h_array[i]; 572 return CEED_ERROR_SUCCESS; 573 } 574 575 //------------------------------------------------------------------------------ 576 // Take reciprocal of a vector on device (impl in .cu file) 577 //------------------------------------------------------------------------------ 578 int CeedDeviceReciprocal_Hip(CeedScalar *d_array, CeedInt length); 579 580 //------------------------------------------------------------------------------ 581 // Take reciprocal of a vector 582 //------------------------------------------------------------------------------ 583 static int CeedVectorReciprocal_Hip(CeedVector vec) { 584 int ierr; 585 Ceed ceed; 586 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 587 CeedVector_Hip *impl; 588 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 589 CeedInt length; 590 ierr = CeedVectorGetLength(vec, &length); CeedChkBackend(ierr); 591 592 // Set value for synced device/host array 593 if (impl->d_array) { 594 ierr = CeedDeviceReciprocal_Hip(impl->d_array, length); CeedChkBackend(ierr); 595 } 596 if (impl->h_array) { 597 ierr = CeedHostReciprocal_Hip(impl->h_array, length); CeedChkBackend(ierr); 598 } 599 600 return CEED_ERROR_SUCCESS; 601 } 602 603 //------------------------------------------------------------------------------ 604 // Compute x = alpha x on the host 605 //------------------------------------------------------------------------------ 606 static int CeedHostScale_Hip(CeedScalar *x_array, CeedScalar alpha, 607 CeedInt length) { 608 for (int i = 0; i < length; i++) 609 x_array[i] *= alpha; 610 return CEED_ERROR_SUCCESS; 611 } 612 613 //------------------------------------------------------------------------------ 614 // Compute x = alpha x on device (impl in .cu file) 615 //------------------------------------------------------------------------------ 616 int CeedDeviceScale_Hip(CeedScalar *x_array, CeedScalar alpha, 617 CeedInt length); 618 619 //------------------------------------------------------------------------------ 620 // Compute x = alpha x 621 //------------------------------------------------------------------------------ 622 static int CeedVectorScale_Hip(CeedVector x, CeedScalar alpha) { 623 int ierr; 624 Ceed ceed; 625 ierr = CeedVectorGetCeed(x, &ceed); CeedChkBackend(ierr); 626 CeedVector_Hip *x_impl; 627 ierr = CeedVectorGetData(x, &x_impl); CeedChkBackend(ierr); 628 CeedInt length; 629 ierr = CeedVectorGetLength(x, &length); CeedChkBackend(ierr); 630 631 // Set value for synced device/host array 632 if (x_impl->d_array) { 633 ierr = CeedDeviceScale_Hip(x_impl->d_array, alpha, length); 634 CeedChkBackend(ierr); 635 } 636 if (x_impl->h_array) { 637 ierr = CeedHostScale_Hip(x_impl->h_array, alpha, length); CeedChkBackend(ierr); 638 } 639 640 return CEED_ERROR_SUCCESS; 641 } 642 643 //------------------------------------------------------------------------------ 644 // Compute y = alpha x + y on the host 645 //------------------------------------------------------------------------------ 646 static int CeedHostAXPY_Hip(CeedScalar *y_array, CeedScalar alpha, 647 CeedScalar *x_array, CeedInt length) { 648 for (int i = 0; i < length; i++) 649 y_array[i] += alpha * x_array[i]; 650 return CEED_ERROR_SUCCESS; 651 } 652 653 //------------------------------------------------------------------------------ 654 // Compute y = alpha x + y on device (impl in .cu file) 655 //------------------------------------------------------------------------------ 656 int CeedDeviceAXPY_Hip(CeedScalar *y_array, CeedScalar alpha, 657 CeedScalar *x_array, CeedInt length); 658 659 //------------------------------------------------------------------------------ 660 // Compute y = alpha x + y 661 //------------------------------------------------------------------------------ 662 static int CeedVectorAXPY_Hip(CeedVector y, CeedScalar alpha, CeedVector x) { 663 int ierr; 664 Ceed ceed; 665 ierr = CeedVectorGetCeed(y, &ceed); CeedChkBackend(ierr); 666 CeedVector_Hip *y_impl, *x_impl; 667 ierr = CeedVectorGetData(y, &y_impl); CeedChkBackend(ierr); 668 ierr = CeedVectorGetData(x, &x_impl); CeedChkBackend(ierr); 669 CeedInt length; 670 ierr = CeedVectorGetLength(y, &length); CeedChkBackend(ierr); 671 672 // Set value for synced device/host array 673 if (y_impl->d_array) { 674 ierr = CeedVectorSyncArray(x, CEED_MEM_DEVICE); CeedChkBackend(ierr); 675 ierr = CeedDeviceAXPY_Hip(y_impl->d_array, alpha, x_impl->d_array, length); 676 CeedChkBackend(ierr); 677 } 678 if (y_impl->h_array) { 679 ierr = CeedVectorSyncArray(x, CEED_MEM_HOST); CeedChkBackend(ierr); 680 ierr = CeedHostAXPY_Hip(y_impl->h_array, alpha, x_impl->h_array, length); 681 CeedChkBackend(ierr); 682 } 683 684 return CEED_ERROR_SUCCESS; 685 } 686 687 //------------------------------------------------------------------------------ 688 // Compute the pointwise multiplication w = x .* y on the host 689 //------------------------------------------------------------------------------ 690 static int CeedHostPointwiseMult_Hip(CeedScalar *w_array, CeedScalar *x_array, 691 CeedScalar *y_array, CeedInt length) { 692 for (int i = 0; i < length; i++) 693 w_array[i] = x_array[i] * y_array[i]; 694 return CEED_ERROR_SUCCESS; 695 } 696 697 //------------------------------------------------------------------------------ 698 // Compute the pointwise multiplication w = x .* y on device (impl in .cu file) 699 //------------------------------------------------------------------------------ 700 int CeedDevicePointwiseMult_Hip(CeedScalar *w_array, CeedScalar *x_array, 701 CeedScalar *y_array, CeedInt length); 702 703 //------------------------------------------------------------------------------ 704 // Compute the pointwise multiplication w = x .* y 705 //------------------------------------------------------------------------------ 706 static int CeedVectorPointwiseMult_Hip(CeedVector w, CeedVector x, 707 CeedVector y) { 708 int ierr; 709 Ceed ceed; 710 ierr = CeedVectorGetCeed(w, &ceed); CeedChkBackend(ierr); 711 CeedVector_Hip *w_impl, *x_impl, *y_impl; 712 ierr = CeedVectorGetData(w, &w_impl); CeedChkBackend(ierr); 713 ierr = CeedVectorGetData(x, &x_impl); CeedChkBackend(ierr); 714 ierr = CeedVectorGetData(y, &y_impl); CeedChkBackend(ierr); 715 CeedInt length; 716 ierr = CeedVectorGetLength(w, &length); CeedChkBackend(ierr); 717 718 // Set value for synced device/host array 719 if (!w_impl->d_array && !w_impl->h_array) { 720 ierr = CeedVectorSetValue(w, 0.0); CeedChkBackend(ierr); 721 } 722 if (w_impl->d_array) { 723 ierr = CeedVectorSyncArray(x, CEED_MEM_DEVICE); CeedChkBackend(ierr); 724 ierr = CeedVectorSyncArray(y, CEED_MEM_DEVICE); CeedChkBackend(ierr); 725 ierr = CeedDevicePointwiseMult_Hip(w_impl->d_array, x_impl->d_array, 726 y_impl->d_array, length); 727 CeedChkBackend(ierr); 728 } 729 if (w_impl->h_array) { 730 ierr = CeedVectorSyncArray(x, CEED_MEM_HOST); CeedChkBackend(ierr); 731 ierr = CeedVectorSyncArray(y, CEED_MEM_HOST); CeedChkBackend(ierr); 732 ierr = CeedHostPointwiseMult_Hip(w_impl->h_array, x_impl->h_array, 733 y_impl->h_array, length); 734 CeedChkBackend(ierr); 735 } 736 737 return CEED_ERROR_SUCCESS; 738 } 739 740 //------------------------------------------------------------------------------ 741 // Destroy the vector 742 //------------------------------------------------------------------------------ 743 static int CeedVectorDestroy_Hip(const CeedVector vec) { 744 int ierr; 745 Ceed ceed; 746 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 747 CeedVector_Hip *impl; 748 ierr = CeedVectorGetData(vec, &impl); CeedChkBackend(ierr); 749 750 ierr = hipFree(impl->d_array_owned); CeedChk_Hip(ceed, ierr); 751 ierr = CeedFree(&impl->h_array_owned); CeedChkBackend(ierr); 752 ierr = CeedFree(&impl); CeedChkBackend(ierr); 753 754 return CEED_ERROR_SUCCESS; 755 } 756 757 //------------------------------------------------------------------------------ 758 // Create a vector of the specified length (does not allocate memory) 759 //------------------------------------------------------------------------------ 760 int CeedVectorCreate_Hip(CeedInt n, CeedVector vec) { 761 CeedVector_Hip *impl; 762 int ierr; 763 Ceed ceed; 764 ierr = CeedVectorGetCeed(vec, &ceed); CeedChkBackend(ierr); 765 766 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "HasValidArray", 767 CeedVectorHasValidArray_Hip); CeedChkBackend(ierr); 768 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "HasBorrowedArrayOfType", 769 CeedVectorHasBorrowedArrayOfType_Hip); 770 CeedChkBackend(ierr); 771 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "SetArray", 772 CeedVectorSetArray_Hip); CeedChkBackend(ierr); 773 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "TakeArray", 774 CeedVectorTakeArray_Hip); CeedChkBackend(ierr); 775 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "SetValue", 776 (int (*)())(CeedVectorSetValue_Hip)); CeedChkBackend(ierr); 777 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "GetArray", 778 CeedVectorGetArray_Hip); CeedChkBackend(ierr); 779 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "GetArrayRead", 780 CeedVectorGetArrayRead_Hip); CeedChkBackend(ierr); 781 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "GetArrayWrite", 782 CeedVectorGetArrayWrite_Hip); CeedChkBackend(ierr); 783 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "RestoreArray", 784 CeedVectorRestoreArray_Hip); CeedChkBackend(ierr); 785 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "RestoreArrayRead", 786 CeedVectorRestoreArrayRead_Hip); CeedChkBackend(ierr); 787 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "Norm", 788 CeedVectorNorm_Hip); CeedChkBackend(ierr); 789 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "Reciprocal", 790 CeedVectorReciprocal_Hip); CeedChkBackend(ierr); 791 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "Scale", 792 (int (*)())(CeedVectorScale_Hip)); CeedChkBackend(ierr); 793 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "AXPY", 794 (int (*)())(CeedVectorAXPY_Hip)); CeedChkBackend(ierr); 795 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "PointwiseMult", 796 CeedVectorPointwiseMult_Hip); CeedChkBackend(ierr); 797 ierr = CeedSetBackendFunction(ceed, "Vector", vec, "Destroy", 798 CeedVectorDestroy_Hip); CeedChkBackend(ierr); 799 800 ierr = CeedCalloc(1, &impl); CeedChkBackend(ierr); 801 ierr = CeedVectorSetData(vec, impl); CeedChkBackend(ierr); 802 803 return CEED_ERROR_SUCCESS; 804 } 805