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