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