32       *need_sync = has_valid_array && !impl->d_array;  in CeedVectorNeedSync_Hip()
53     impl->d_array = impl->d_array_borrowed;  in CeedVectorSyncH2D_Hip()
55     impl->d_array = impl->d_array_owned;  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()
74 …CeedCheck(impl->d_array, CeedVectorReturnCeed(vec), CEED_ERROR_BACKEND, "No valid device data to s…  in CeedVectorSyncD2H_Hip()
90 …CeedCallHip(CeedVectorReturnCeed(vec), hipMemcpy(impl->h_array, impl->d_array, bytes, hipMemcpyDev…  in CeedVectorSyncD2H_Hip()
129   impl->d_array = NULL;  in CeedVectorSetAllInvalid_Hip()
140   *has_valid_array = impl->h_array || impl->d_array;  in CeedVectorHasValidArray_Hip()
212 …               (const CeedScalar **)&impl->d_array_borrowed, (const CeedScalar **)&impl->d_array));  in CeedVectorSetArrayDevice_Hip()
232       if (!impl->d_array) {  in CeedVectorSetArrayUnifiedHostToDevice_Hip()
234           impl->d_array = impl->d_array_borrowed;  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()
248       impl->d_array          = impl->h_array_borrowed;  in CeedVectorSetArrayUnifiedHostToDevice_Hip()
287 int CeedDeviceCopyStrided_Hip(CeedScalar *d_array, CeedSize start, CeedSize stop, CeedSize step, Ce…
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()
335     impl->d_array = NULL;  in CeedVectorCopyStrided_Hip()
353 int CeedDeviceSetValue_Hip(CeedScalar *d_array, CeedSize length, CeedScalar val);
367   if (!impl->d_array && !impl->h_array) {  in CeedVectorSetValue_Hip()
369       impl->d_array = impl->d_array_borrowed;  in CeedVectorSetValue_Hip()
373       impl->d_array = impl->d_array_owned;  in CeedVectorSetValue_Hip()
380   if (impl->d_array) {  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()
389     impl->d_array = NULL;  in CeedVectorSetValue_Hip()
405 int CeedDeviceSetValueStrided_Hip(CeedScalar *d_array, CeedSize start, CeedSize stop, CeedSize step…
418   if (impl->d_array) {  in CeedVectorSetValueStrided_Hip()
419     CeedCallBackend(CeedDeviceSetValueStrided_Hip(impl->d_array, start, stop, step, val));  in CeedVectorSetValueStrided_Hip()
423     impl->d_array = NULL;  in CeedVectorSetValueStrided_Hip()
451       impl->d_array          = NULL;  in CeedVectorTakeArray_Hip()
478       *array = impl->d_array;  in CeedVectorGetArrayCore_Hip()
508       if (impl->has_unified_addressing) impl->d_array = *array;  in CeedVectorGetArray_Hip()
511       impl->d_array = *array;  in CeedVectorGetArray_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()
559   const CeedScalar *d_array;  in CeedVectorNorm_Hip()  local
580   CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &d_array));  in CeedVectorNorm_Hip()
586 …CeedCallHipblas(ceed, hipblasSasum_64(handle, (int64_t)length, (float *)d_array, 1, (float *)norm)…  in CeedVectorNorm_Hip()
593         d_array_start             = (float *)d_array + (CeedSize)(i)*INT_MAX;  in CeedVectorNorm_Hip()
604 …CeedCallHipblas(ceed, hipblasDasum_64(handle, (int64_t)length, (double *)d_array, 1, (double *)nor…  in CeedVectorNorm_Hip()
611         d_array_start             = (double *)d_array + (CeedSize)(i)*INT_MAX;  in CeedVectorNorm_Hip()
626 …CeedCallHipblas(ceed, hipblasSnrm2_64(handle, (int64_t)length, (float *)d_array, 1, (float *)norm)…  in CeedVectorNorm_Hip()
633         d_array_start             = (float *)d_array + (CeedSize)(i)*INT_MAX;  in CeedVectorNorm_Hip()
645 …CeedCallHipblas(ceed, hipblasDnrm2_64(handle, (int64_t)length, (double *)d_array, 1, (double *)nor…  in CeedVectorNorm_Hip()
652         d_array_start             = (double *)d_array + (CeedSize)(i)*INT_MAX;  in CeedVectorNorm_Hip()
671       CeedCallHipblas(ceed, hipblasIsamax_64(handle, (int64_t)length, (float *)d_array, 1, &index));  in CeedVectorNorm_Hip()
672 …CeedCallHip(ceed, hipMemcpyAsync(&norm_no_abs, impl->d_array + index - 1, sizeof(CeedScalar), hipM…  in CeedVectorNorm_Hip()
681         d_array_start             = (float *)d_array + (CeedSize)(i)*INT_MAX;  in CeedVectorNorm_Hip()
688           sub_max = fabs(d_array[index - 1]);  in CeedVectorNorm_Hip()
702 …    CeedCallHipblas(ceed, hipblasIdamax_64(handle, (int64_t)length, (double *)d_array, 1, &index));  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()
717         d_array_start             = (double *)d_array + (CeedSize)(i)*INT_MAX;  in CeedVectorNorm_Hip()
724           sub_max = fabs(d_array[index - 1]);  in CeedVectorNorm_Hip()
737   CeedCallBackend(CeedVectorRestoreArrayRead(vec, &d_array));  in CeedVectorNorm_Hip()
755 int CeedDeviceReciprocal_Hip(CeedScalar *d_array, CeedSize length);
767   if (impl->d_array) CeedCallBackend(CeedDeviceReciprocal_Hip(impl->d_array, length));  in CeedVectorReciprocal_Hip()
795   if (impl->d_array) {  in CeedVectorScale_Hip()
803 …pblas(CeedVectorReturnCeed(x), hipblasSscal_64(handle, (int64_t)length, &alpha, impl->d_array, 1));  in CeedVectorScale_Hip()
805 …pblas(CeedVectorReturnCeed(x), hipblasDscal_64(handle, (int64_t)length, &alpha, impl->d_array, 1));  in CeedVectorScale_Hip()
809     CeedCallBackend(CeedDeviceScale_Hip(impl->d_array, alpha, length));  in CeedVectorScale_Hip()
815     impl->d_array = NULL;  in CeedVectorScale_Hip()
844   if (y_impl->d_array) {  in CeedVectorAXPY_Hip()
853 …Ceed(y), hipblasSaxpy_64(handle, (int64_t)length, &alpha, x_impl->d_array, 1, y_impl->d_array, 1));  in CeedVectorAXPY_Hip()
855 …Ceed(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()
865     y_impl->d_array = NULL;  in CeedVectorAXPY_Hip()
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()
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()