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