Lines Matching full:-
1 // Copyright (c) 2017-2026, Lawrence Livermore National Security, LLC and other CEED contributors.
2 // All Rights Reserved. See the top-level LICENSE and NOTICE files for details.
4 // SPDX-License-Identifier: BSD-2-Clause
15 #include "../hip/ceed-hip-common.h"
16 #include "ceed-hip-ref.h"
18 //------------------------------------------------------------------------------
20 //------------------------------------------------------------------------------
29 *need_sync = has_valid_array && !impl->h_array; in CeedVectorNeedSync_Hip()
32 *need_sync = has_valid_array && !impl->d_array; in CeedVectorNeedSync_Hip()
38 //------------------------------------------------------------------------------
40 //------------------------------------------------------------------------------
48 …CeedCheck(impl->h_array, CeedVectorReturnCeed(vec), CEED_ERROR_BACKEND, "No valid host data to syn… in CeedVectorSyncH2D_Hip()
52 if (impl->d_array_borrowed) { in CeedVectorSyncH2D_Hip()
53 impl->d_array = impl->d_array_borrowed; in CeedVectorSyncH2D_Hip()
54 } else if (impl->d_array_owned) { in CeedVectorSyncH2D_Hip()
55 impl->d_array = impl->d_array_owned; in CeedVectorSyncH2D_Hip()
57 CeedCallHip(CeedVectorReturnCeed(vec), hipMalloc((void **)&impl->d_array_owned, bytes)); in CeedVectorSyncH2D_Hip()
58 impl->d_array = impl->d_array_owned; in CeedVectorSyncH2D_Hip()
60 …CeedCallHip(CeedVectorReturnCeed(vec), hipMemcpy(impl->d_array, impl->h_array, bytes, hipMemcpyHos… in CeedVectorSyncH2D_Hip()
64 //------------------------------------------------------------------------------
66 //------------------------------------------------------------------------------
74 …CeedCheck(impl->d_array, CeedVectorReturnCeed(vec), CEED_ERROR_BACKEND, "No valid device data to s… in CeedVectorSyncD2H_Hip()
76 if (impl->h_array_borrowed) { in CeedVectorSyncD2H_Hip()
77 impl->h_array = impl->h_array_borrowed; in CeedVectorSyncD2H_Hip()
78 } else if (impl->h_array_owned) { in CeedVectorSyncD2H_Hip()
79 impl->h_array = impl->h_array_owned; in CeedVectorSyncD2H_Hip()
84 CeedCallBackend(CeedCalloc(length, &impl->h_array_owned)); in CeedVectorSyncD2H_Hip()
85 impl->h_array = impl->h_array_owned; in CeedVectorSyncD2H_Hip()
90 …CeedCallHip(CeedVectorReturnCeed(vec), hipMemcpy(impl->h_array, impl->d_array, bytes, hipMemcpyDev… in CeedVectorSyncD2H_Hip()
94 //------------------------------------------------------------------------------
96 //------------------------------------------------------------------------------
103 if (impl->has_unified_addressing && !impl->h_array_borrowed) { in CeedVectorSyncArray_Hip()
121 //------------------------------------------------------------------------------
123 //------------------------------------------------------------------------------
128 impl->h_array = NULL; in CeedVectorSetAllInvalid_Hip()
129 impl->d_array = NULL; in CeedVectorSetAllInvalid_Hip()
133 //------------------------------------------------------------------------------
135 //------------------------------------------------------------------------------
140 *has_valid_array = impl->h_array || impl->d_array; in CeedVectorHasValidArray_Hip()
144 //------------------------------------------------------------------------------
146 //------------------------------------------------------------------------------
153 *has_array_of_type = impl->h_array_borrowed || impl->h_array_owned; in CeedVectorHasArrayOfType_Hip()
156 *has_array_of_type = impl->d_array_borrowed || impl->d_array_owned; in CeedVectorHasArrayOfType_Hip()
162 //------------------------------------------------------------------------------
164 //------------------------------------------------------------------------------
171 mem_type = impl->has_unified_addressing && !impl->h_array_borrowed ? CEED_MEM_DEVICE : mem_type; in CeedVectorHasBorrowedArrayOfType_Hip()
175 *has_borrowed_array_of_type = impl->h_array_borrowed; in CeedVectorHasBorrowedArrayOfType_Hip()
178 *has_borrowed_array_of_type = impl->d_array_borrowed; in CeedVectorHasBorrowedArrayOfType_Hip()
184 //------------------------------------------------------------------------------
186 //------------------------------------------------------------------------------
194 …CeedCallBackend(CeedSetHostCeedScalarArray(array, copy_mode, length, (const CeedScalar **)&impl->h… in CeedVectorSetArrayHost_Hip()
195 … (const CeedScalar **)&impl->h_array_borrowed, (const CeedScalar **)&impl->h_array)); in CeedVectorSetArrayHost_Hip()
199 //------------------------------------------------------------------------------
201 //------------------------------------------------------------------------------
211 …eviceCeedScalarArray_Hip(ceed, array, copy_mode, length, (const CeedScalar **)&impl->d_array_owned, in CeedVectorSetArrayDevice_Hip()
212 … (const CeedScalar **)&impl->d_array_borrowed, (const CeedScalar **)&impl->d_array)); in CeedVectorSetArrayDevice_Hip()
217 //------------------------------------------------------------------------------
219 //------------------------------------------------------------------------------
232 if (!impl->d_array) { in CeedVectorSetArrayUnifiedHostToDevice_Hip()
233 if (impl->d_array_borrowed) { in CeedVectorSetArrayUnifiedHostToDevice_Hip()
234 impl->d_array = impl->d_array_borrowed; in CeedVectorSetArrayUnifiedHostToDevice_Hip()
236 …if (!impl->d_array_owned) CeedCallHip(ceed, hipMalloc((void **)&impl->d_array_owned, sizeof(CeedSc… in CeedVectorSetArrayUnifiedHostToDevice_Hip()
237 impl->d_array = impl->d_array_owned; in CeedVectorSetArrayUnifiedHostToDevice_Hip()
240 …if (array) CeedCallHip(ceed, hipMemcpy(impl->d_array, array, sizeof(CeedScalar) * length, hipMemcp… in CeedVectorSetArrayUnifiedHostToDevice_Hip()
244 CeedCallHip(ceed, hipFree(impl->d_array_owned)); in CeedVectorSetArrayUnifiedHostToDevice_Hip()
245 CeedCallBackend(CeedFree(&impl->h_array_owned)); in CeedVectorSetArrayUnifiedHostToDevice_Hip()
246 impl->h_array_owned = NULL; in CeedVectorSetArrayUnifiedHostToDevice_Hip()
247 impl->h_array_borrowed = array; in CeedVectorSetArrayUnifiedHostToDevice_Hip()
248 impl->d_array = impl->h_array_borrowed; in CeedVectorSetArrayUnifiedHostToDevice_Hip()
254 //------------------------------------------------------------------------------
257 //------------------------------------------------------------------------------
265 if (impl->has_unified_addressing) { in CeedVectorSetArray_Hip()
276 //------------------------------------------------------------------------------
278 //------------------------------------------------------------------------------
284 //------------------------------------------------------------------------------
286 //------------------------------------------------------------------------------
289 //------------------------------------------------------------------------------
291 //------------------------------------------------------------------------------
304 if (stop == -1) stop = length; in CeedVectorCopyStrided_Hip()
306 if (impl->d_array) { in CeedVectorCopyStrided_Hip()
318 …CeedCallHipblas(ceed, hipblasScopy_64(handle, (int64_t)(stop - start), impl->d_array + start, (int… in CeedVectorCopyStrided_Hip()
320 …CeedCallHipblas(ceed, hipblasDcopy_64(handle, (int64_t)(stop - start), impl->d_array + start, (int… in CeedVectorCopyStrided_Hip()
324 CeedCallBackend(CeedDeviceCopyStrided_Hip(impl->d_array, start, stop, step, copy_array)); in CeedVectorCopyStrided_Hip()
327 impl->h_array = NULL; in CeedVectorCopyStrided_Hip()
329 } else if (impl->h_array) { in CeedVectorCopyStrided_Hip()
333 CeedCallBackend(CeedHostCopyStrided_Hip(impl->h_array, start, stop, step, copy_array)); in CeedVectorCopyStrided_Hip()
335 impl->d_array = NULL; in CeedVectorCopyStrided_Hip()
342 //------------------------------------------------------------------------------
344 //------------------------------------------------------------------------------
350 //------------------------------------------------------------------------------
352 //------------------------------------------------------------------------------
355 //------------------------------------------------------------------------------
357 //------------------------------------------------------------------------------
367 if (!impl->d_array && !impl->h_array) { in CeedVectorSetValue_Hip()
368 if (impl->d_array_borrowed) { in CeedVectorSetValue_Hip()
369 impl->d_array = impl->d_array_borrowed; in CeedVectorSetValue_Hip()
370 } else if (impl->h_array_borrowed) { in CeedVectorSetValue_Hip()
371 impl->h_array = impl->h_array_borrowed; in CeedVectorSetValue_Hip()
372 } else if (impl->d_array_owned) { in CeedVectorSetValue_Hip()
373 impl->d_array = impl->d_array_owned; in CeedVectorSetValue_Hip()
374 } else if (impl->h_array_owned) { in CeedVectorSetValue_Hip()
375 impl->h_array = impl->h_array_owned; in CeedVectorSetValue_Hip()
380 if (impl->d_array) { in CeedVectorSetValue_Hip()
381 if (val == 0 && !impl->h_array_borrowed) { in CeedVectorSetValue_Hip()
382 … CeedCallHip(CeedVectorReturnCeed(vec), hipMemset(impl->d_array, 0, length * sizeof(CeedScalar))); in CeedVectorSetValue_Hip()
384 CeedCallBackend(CeedDeviceSetValue_Hip(impl->d_array, length, val)); in CeedVectorSetValue_Hip()
386 impl->h_array = NULL; in CeedVectorSetValue_Hip()
387 } else if (impl->h_array) { in CeedVectorSetValue_Hip()
388 CeedCallBackend(CeedHostSetValue_Hip(impl->h_array, length, val)); in CeedVectorSetValue_Hip()
389 impl->d_array = NULL; in CeedVectorSetValue_Hip()
394 //------------------------------------------------------------------------------
396 //------------------------------------------------------------------------------
402 //------------------------------------------------------------------------------
404 //------------------------------------------------------------------------------
407 //------------------------------------------------------------------------------
409 //------------------------------------------------------------------------------
417 if (stop == -1) stop = length; in CeedVectorSetValueStrided_Hip()
418 if (impl->d_array) { in CeedVectorSetValueStrided_Hip()
419 CeedCallBackend(CeedDeviceSetValueStrided_Hip(impl->d_array, start, stop, step, val)); in CeedVectorSetValueStrided_Hip()
420 impl->h_array = NULL; in CeedVectorSetValueStrided_Hip()
421 } else if (impl->h_array) { in CeedVectorSetValueStrided_Hip()
422 CeedCallBackend(CeedHostSetValueStrided_Hip(impl->h_array, start, stop, step, val)); in CeedVectorSetValueStrided_Hip()
423 impl->d_array = NULL; in CeedVectorSetValueStrided_Hip()
430 //------------------------------------------------------------------------------
432 //------------------------------------------------------------------------------
444 (*array) = impl->h_array_borrowed; in CeedVectorTakeArray_Hip()
445 impl->h_array_borrowed = NULL; in CeedVectorTakeArray_Hip()
446 impl->h_array = NULL; in CeedVectorTakeArray_Hip()
449 (*array) = impl->d_array_borrowed; in CeedVectorTakeArray_Hip()
450 impl->d_array_borrowed = NULL; in CeedVectorTakeArray_Hip()
451 impl->d_array = NULL; in CeedVectorTakeArray_Hip()
457 //------------------------------------------------------------------------------
460 //------------------------------------------------------------------------------
467 mem_type = impl->has_unified_addressing && !impl->h_array_borrowed ? CEED_MEM_DEVICE : mem_type; in CeedVectorGetArrayCore_Hip()
475 *array = impl->h_array; in CeedVectorGetArrayCore_Hip()
478 *array = impl->d_array; in CeedVectorGetArrayCore_Hip()
484 //------------------------------------------------------------------------------
485 // Get read-only access to a vector via the specified mem_type
486 //------------------------------------------------------------------------------
491 //------------------------------------------------------------------------------
493 //------------------------------------------------------------------------------
500 mem_type = impl->has_unified_addressing && !impl->h_array_borrowed ? CEED_MEM_DEVICE : mem_type; in CeedVectorGetArray_Hip()
507 impl->h_array = *array; in CeedVectorGetArray_Hip()
508 if (impl->has_unified_addressing) impl->d_array = *array; in CeedVectorGetArray_Hip()
511 impl->d_array = *array; in CeedVectorGetArray_Hip()
517 //------------------------------------------------------------------------------
519 //------------------------------------------------------------------------------
529 mem_type = impl->has_unified_addressing && !impl->h_array_borrowed ? CEED_MEM_DEVICE : mem_type; in CeedVectorGetArrayWrite_Hip()
539 if (impl->h_array_borrowed) impl->h_array = impl->h_array_borrowed; in CeedVectorGetArrayWrite_Hip()
540 else impl->h_array = impl->h_array_owned; in CeedVectorGetArrayWrite_Hip()
543 if (impl->d_array_borrowed) impl->d_array = impl->d_array_borrowed; in CeedVectorGetArrayWrite_Hip()
544 else impl->d_array = impl->d_array_owned; in CeedVectorGetArrayWrite_Hip()
550 //------------------------------------------------------------------------------
552 //------------------------------------------------------------------------------
572 // With ROCm 6, we can use the 64-bit integer interface. Prior to that, in CeedVectorNorm_Hip()
585 #if (HIP_VERSION >= 60000000) // We have ROCm 6, and can use 64-bit integers in CeedVectorNorm_Hip()
594 CeedSize remaining_length = length - (CeedSize)(i)*INT_MAX; in CeedVectorNorm_Hip()
595 CeedInt sub_length = (i == num_calls - 1) ? (CeedInt)(remaining_length) : INT_MAX; in CeedVectorNorm_Hip()
612 CeedSize remaining_length = length - (CeedSize)(i)*INT_MAX; in CeedVectorNorm_Hip()
613 CeedInt sub_length = (i == num_calls - 1) ? (CeedInt)(remaining_length) : INT_MAX; in CeedVectorNorm_Hip()
634 CeedSize remaining_length = length - (CeedSize)(i)*INT_MAX; in CeedVectorNorm_Hip()
635 CeedInt sub_length = (i == num_calls - 1) ? (CeedInt)(remaining_length) : INT_MAX; in CeedVectorNorm_Hip()
653 CeedSize remaining_length = length - (CeedSize)(i)*INT_MAX; in CeedVectorNorm_Hip()
654 CeedInt sub_length = (i == num_calls - 1) ? (CeedInt)(remaining_length) : INT_MAX; in CeedVectorNorm_Hip()
672 …CeedCallHip(ceed, hipMemcpyAsync(&norm_no_abs, impl->d_array + index - 1, sizeof(CeedScalar), hipM… in CeedVectorNorm_Hip()
682 CeedSize remaining_length = length - (CeedSize)(i)*INT_MAX; in CeedVectorNorm_Hip()
683 CeedInt sub_length = (i == num_calls - 1) ? (CeedInt)(remaining_length) : INT_MAX; in CeedVectorNorm_Hip()
686 if (hip_data->has_unified_addressing) { in CeedVectorNorm_Hip()
688 sub_max = fabs(d_array[index - 1]); in CeedVectorNorm_Hip()
690 …CeedCallHip(ceed, hipMemcpyAsync(&sub_max, d_array_start + index - 1, sizeof(CeedScalar), hipMemcp… in CeedVectorNorm_Hip()
703 if (hip_data->has_unified_addressing) { in CeedVectorNorm_Hip()
705 norm_no_abs = fabs(d_array[index - 1]); in CeedVectorNorm_Hip()
707 …CeedCallHip(ceed, hipMemcpyAsync(&norm_no_abs, impl->d_array + index - 1, sizeof(CeedScalar), hipM… in CeedVectorNorm_Hip()
718 CeedSize remaining_length = length - (CeedSize)(i)*INT_MAX; in CeedVectorNorm_Hip()
719 CeedInt sub_length = (i == num_calls - 1) ? (CeedInt)(remaining_length) : INT_MAX; in CeedVectorNorm_Hip()
722 if (hip_data->has_unified_addressing) { in CeedVectorNorm_Hip()
724 sub_max = fabs(d_array[index - 1]); in CeedVectorNorm_Hip()
726 …CeedCallHip(ceed, hipMemcpyAsync(&sub_max, d_array_start + index - 1, sizeof(CeedScalar), hipMemcp… in CeedVectorNorm_Hip()
742 //------------------------------------------------------------------------------
744 //------------------------------------------------------------------------------
752 //------------------------------------------------------------------------------
754 //------------------------------------------------------------------------------
757 //------------------------------------------------------------------------------
759 //------------------------------------------------------------------------------
767 if (impl->d_array) CeedCallBackend(CeedDeviceReciprocal_Hip(impl->d_array, length)); in CeedVectorReciprocal_Hip()
768 if (impl->h_array) CeedCallBackend(CeedHostReciprocal_Hip(impl->h_array, length)); in CeedVectorReciprocal_Hip()
772 //------------------------------------------------------------------------------
774 //------------------------------------------------------------------------------
780 //------------------------------------------------------------------------------
782 //------------------------------------------------------------------------------
785 //------------------------------------------------------------------------------
787 //------------------------------------------------------------------------------
795 if (impl->d_array) { in CeedVectorScale_Hip()
803 …CeedCallHipblas(CeedVectorReturnCeed(x), hipblasSscal_64(handle, (int64_t)length, &alpha, impl->d_… in CeedVectorScale_Hip()
805 …CeedCallHipblas(CeedVectorReturnCeed(x), hipblasDscal_64(handle, (int64_t)length, &alpha, impl->d_… in CeedVectorScale_Hip()
809 CeedCallBackend(CeedDeviceScale_Hip(impl->d_array, alpha, length)); in CeedVectorScale_Hip()
811 impl->h_array = NULL; in CeedVectorScale_Hip()
813 if (impl->h_array) { in CeedVectorScale_Hip()
814 CeedCallBackend(CeedHostScale_Hip(impl->h_array, alpha, length)); in CeedVectorScale_Hip()
815 impl->d_array = NULL; in CeedVectorScale_Hip()
820 //------------------------------------------------------------------------------
822 //------------------------------------------------------------------------------
828 //------------------------------------------------------------------------------
830 //------------------------------------------------------------------------------
833 //------------------------------------------------------------------------------
835 //------------------------------------------------------------------------------
844 if (y_impl->d_array) { in CeedVectorAXPY_Hip()
853 …rnCeed(y), hipblasSaxpy_64(handle, (int64_t)length, &alpha, x_impl->d_array, 1, y_impl->d_array, 1… in CeedVectorAXPY_Hip()
855 …rnCeed(y), hipblasDaxpy_64(handle, (int64_t)length, &alpha, x_impl->d_array, 1, y_impl->d_array, 1… in CeedVectorAXPY_Hip()
859 CeedCallBackend(CeedDeviceAXPY_Hip(y_impl->d_array, alpha, x_impl->d_array, length)); in CeedVectorAXPY_Hip()
861 y_impl->h_array = NULL; in CeedVectorAXPY_Hip()
862 } else if (y_impl->h_array) { in CeedVectorAXPY_Hip()
864 CeedCallBackend(CeedHostAXPY_Hip(y_impl->h_array, alpha, x_impl->h_array, length)); in CeedVectorAXPY_Hip()
865 y_impl->d_array = NULL; in CeedVectorAXPY_Hip()
870 //------------------------------------------------------------------------------
872 //------------------------------------------------------------------------------
878 //------------------------------------------------------------------------------
880 //------------------------------------------------------------------------------
883 //------------------------------------------------------------------------------
885 //------------------------------------------------------------------------------
894 if (y_impl->d_array) { in CeedVectorAXPBY_Hip()
896 CeedCallBackend(CeedDeviceAXPBY_Hip(y_impl->d_array, alpha, beta, x_impl->d_array, length)); in CeedVectorAXPBY_Hip()
898 if (y_impl->h_array) { in CeedVectorAXPBY_Hip()
900 CeedCallBackend(CeedHostAXPBY_Hip(y_impl->h_array, alpha, beta, x_impl->h_array, length)); in CeedVectorAXPBY_Hip()
905 //------------------------------------------------------------------------------
907 //------------------------------------------------------------------------------
913 //------------------------------------------------------------------------------
915 //------------------------------------------------------------------------------
918 //------------------------------------------------------------------------------
920 //------------------------------------------------------------------------------
931 if (!w_impl->d_array && !w_impl->h_array) { in CeedVectorPointwiseMult_Hip()
934 if (w_impl->d_array) { in CeedVectorPointwiseMult_Hip()
937 …CeedCallBackend(CeedDevicePointwiseMult_Hip(w_impl->d_array, x_impl->d_array, y_impl->d_array, len… in CeedVectorPointwiseMult_Hip()
939 if (w_impl->h_array) { in CeedVectorPointwiseMult_Hip()
942 …CeedCallBackend(CeedHostPointwiseMult_Hip(w_impl->h_array, x_impl->h_array, y_impl->h_array, lengt… in CeedVectorPointwiseMult_Hip()
947 //------------------------------------------------------------------------------
949 //------------------------------------------------------------------------------
954 CeedCallHip(CeedVectorReturnCeed(vec), hipFree(impl->d_array_owned)); in CeedVectorDestroy_Hip()
955 CeedCallBackend(CeedFree(&impl->h_array_owned)); in CeedVectorDestroy_Hip()
960 //------------------------------------------------------------------------------
962 //------------------------------------------------------------------------------
990 impl->has_unified_addressing = hip_impl->has_unified_addressing; in CeedVectorCreate_Hip()
995 //------------------------------------------------------------------------------