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 "ceed-sycl-ref.hpp"
17 //------------------------------------------------------------------------------
19 //------------------------------------------------------------------------------
28 *need_sync = has_valid_array && !impl->h_array; in CeedVectorNeedSync_Sycl()
31 *need_sync = has_valid_array && !impl->d_array; in CeedVectorNeedSync_Sycl()
37 //------------------------------------------------------------------------------
39 //------------------------------------------------------------------------------
50 CeedCheck(impl->h_array, ceed, CEED_ERROR_BACKEND, "No valid host data to sync to device"); in CeedVectorSyncH2D_Sycl()
53 if (impl->d_array_borrowed) { in CeedVectorSyncH2D_Sycl()
54 impl->d_array = impl->d_array_borrowed; in CeedVectorSyncH2D_Sycl()
55 } else if (impl->d_array_owned) { in CeedVectorSyncH2D_Sycl()
56 impl->d_array = impl->d_array_owned; in CeedVectorSyncH2D_Sycl()
58 …CeedCallSycl(ceed, impl->d_array_owned = sycl::malloc_device<CeedScalar>(length, data->sycl_device… in CeedVectorSyncH2D_Sycl()
59 impl->d_array = impl->d_array_owned; in CeedVectorSyncH2D_Sycl()
65 if (!data->sycl_queue.is_in_order()) e = {data->sycl_queue.ext_oneapi_submit_barrier()}; in CeedVectorSyncH2D_Sycl()
66 …CeedCallSycl(ceed, data->sycl_queue.copy<CeedScalar>(impl->h_array, impl->d_array, length, e).wait… in CeedVectorSyncH2D_Sycl()
71 //------------------------------------------------------------------------------
73 //------------------------------------------------------------------------------
84 CeedCheck(impl->d_array, ceed, CEED_ERROR_BACKEND, "No valid device data to sync to host"); in CeedVectorSyncD2H_Sycl()
87 if (impl->h_array_borrowed) { in CeedVectorSyncD2H_Sycl()
88 impl->h_array = impl->h_array_borrowed; in CeedVectorSyncD2H_Sycl()
89 } else if (impl->h_array_owned) { in CeedVectorSyncD2H_Sycl()
90 impl->h_array = impl->h_array_owned; in CeedVectorSyncD2H_Sycl()
92 CeedCallBackend(CeedCalloc(length, &impl->h_array_owned)); in CeedVectorSyncD2H_Sycl()
93 impl->h_array = impl->h_array_owned; in CeedVectorSyncD2H_Sycl()
99 if (!data->sycl_queue.is_in_order()) e = {data->sycl_queue.ext_oneapi_submit_barrier()}; in CeedVectorSyncD2H_Sycl()
100 …CeedCallSycl(ceed, data->sycl_queue.copy<CeedScalar>(impl->d_array, impl->h_array, length, e).wait… in CeedVectorSyncD2H_Sycl()
105 //------------------------------------------------------------------------------
107 //------------------------------------------------------------------------------
124 //------------------------------------------------------------------------------
126 //------------------------------------------------------------------------------
131 impl->h_array = NULL; in CeedVectorSetAllInvalid_Sycl()
132 impl->d_array = NULL; in CeedVectorSetAllInvalid_Sycl()
136 //------------------------------------------------------------------------------
138 //------------------------------------------------------------------------------
143 *has_valid_array = impl->h_array || impl->d_array; in CeedVectorHasValidArray_Sycl()
147 //------------------------------------------------------------------------------
149 //------------------------------------------------------------------------------
156 *has_array_of_type = impl->h_array_borrowed || impl->h_array_owned; in CeedVectorHasArrayOfType_Sycl()
159 *has_array_of_type = impl->d_array_borrowed || impl->d_array_owned; in CeedVectorHasArrayOfType_Sycl()
165 //------------------------------------------------------------------------------
167 //------------------------------------------------------------------------------
174 *has_borrowed_array_of_type = impl->h_array_borrowed; in CeedVectorHasBorrowedArrayOfType_Sycl()
177 *has_borrowed_array_of_type = impl->d_array_borrowed; in CeedVectorHasBorrowedArrayOfType_Sycl()
183 //------------------------------------------------------------------------------
185 //------------------------------------------------------------------------------
193 …CeedCallBackend(CeedSetHostCeedScalarArray(array, copy_mode, length, (const CeedScalar **)&impl->h… in CeedVectorSetArrayHost_Sycl()
194 … (const CeedScalar **)&impl->h_array_borrowed, (const CeedScalar **)&impl->h_array)); in CeedVectorSetArrayHost_Sycl()
198 //------------------------------------------------------------------------------
200 //------------------------------------------------------------------------------
215 if (!data->sycl_queue.is_in_order()) e = {data->sycl_queue.ext_oneapi_submit_barrier()}; in CeedVectorSetArrayDevice_Sycl()
219 if (!impl->d_array_owned) { in CeedVectorSetArrayDevice_Sycl()
220 …CeedCallSycl(ceed, impl->d_array_owned = sycl::malloc_device<CeedScalar>(length, data->sycl_device… in CeedVectorSetArrayDevice_Sycl()
224 …CeedCallSycl(ceed, data->sycl_queue.copy<CeedScalar>(array, impl->d_array_owned, length, e).wait_a… in CeedVectorSetArrayDevice_Sycl()
226 impl->d_array_borrowed = NULL; in CeedVectorSetArrayDevice_Sycl()
227 impl->d_array = impl->d_array_owned; in CeedVectorSetArrayDevice_Sycl()
230 if (impl->d_array_owned) { in CeedVectorSetArrayDevice_Sycl()
232 CeedCallSycl(ceed, data->sycl_queue.wait_and_throw()); in CeedVectorSetArrayDevice_Sycl()
233 CeedCallSycl(ceed, sycl::free(impl->d_array_owned, data->sycl_context)); in CeedVectorSetArrayDevice_Sycl()
235 impl->d_array_owned = array; in CeedVectorSetArrayDevice_Sycl()
236 impl->d_array_borrowed = NULL; in CeedVectorSetArrayDevice_Sycl()
237 impl->d_array = impl->d_array_owned; in CeedVectorSetArrayDevice_Sycl()
240 if (impl->d_array_owned) { in CeedVectorSetArrayDevice_Sycl()
242 CeedCallSycl(ceed, data->sycl_queue.wait_and_throw()); in CeedVectorSetArrayDevice_Sycl()
243 CeedCallSycl(ceed, sycl::free(impl->d_array_owned, data->sycl_context)); in CeedVectorSetArrayDevice_Sycl()
245 impl->d_array_owned = NULL; in CeedVectorSetArrayDevice_Sycl()
246 impl->d_array_borrowed = array; in CeedVectorSetArrayDevice_Sycl()
247 impl->d_array = impl->d_array_borrowed; in CeedVectorSetArrayDevice_Sycl()
254 //------------------------------------------------------------------------------
257 //------------------------------------------------------------------------------
273 //------------------------------------------------------------------------------
275 //------------------------------------------------------------------------------
281 //------------------------------------------------------------------------------
283 //------------------------------------------------------------------------------
292 //------------------------------------------------------------------------------
294 //------------------------------------------------------------------------------
308 if (!impl->d_array && !impl->h_array) { in CeedVectorSetValue_Sycl()
309 if (impl->d_array_borrowed) { in CeedVectorSetValue_Sycl()
310 impl->d_array = impl->d_array_borrowed; in CeedVectorSetValue_Sycl()
311 } else if (impl->h_array_borrowed) { in CeedVectorSetValue_Sycl()
312 impl->h_array = impl->h_array_borrowed; in CeedVectorSetValue_Sycl()
313 } else if (impl->d_array_owned) { in CeedVectorSetValue_Sycl()
314 impl->d_array = impl->d_array_owned; in CeedVectorSetValue_Sycl()
315 } else if (impl->h_array_owned) { in CeedVectorSetValue_Sycl()
316 impl->h_array = impl->h_array_owned; in CeedVectorSetValue_Sycl()
321 if (impl->d_array) { in CeedVectorSetValue_Sycl()
322 CeedCallBackend(CeedDeviceSetValue_Sycl(data->sycl_queue, impl->d_array, length, val)); in CeedVectorSetValue_Sycl()
323 impl->h_array = NULL; in CeedVectorSetValue_Sycl()
325 if (impl->h_array) { in CeedVectorSetValue_Sycl()
326 CeedCallBackend(CeedHostSetValue_Sycl(impl->h_array, length, val)); in CeedVectorSetValue_Sycl()
327 impl->d_array = NULL; in CeedVectorSetValue_Sycl()
332 //------------------------------------------------------------------------------
334 //------------------------------------------------------------------------------
347 if (!data->sycl_queue.is_in_order()) data->sycl_queue.ext_oneapi_submit_barrier(); in CeedVectorTakeArray_Sycl()
355 (*array) = impl->h_array_borrowed; in CeedVectorTakeArray_Sycl()
356 impl->h_array_borrowed = NULL; in CeedVectorTakeArray_Sycl()
357 impl->h_array = NULL; in CeedVectorTakeArray_Sycl()
360 (*array) = impl->d_array_borrowed; in CeedVectorTakeArray_Sycl()
361 impl->d_array_borrowed = NULL; in CeedVectorTakeArray_Sycl()
362 impl->d_array = NULL; in CeedVectorTakeArray_Sycl()
368 //------------------------------------------------------------------------------
371 //------------------------------------------------------------------------------
383 *array = impl->h_array; in CeedVectorGetArrayCore_Sycl()
386 *array = impl->d_array; in CeedVectorGetArrayCore_Sycl()
392 //------------------------------------------------------------------------------
393 // Get read-only access to a vector via the specified mem_type
394 //------------------------------------------------------------------------------
399 //------------------------------------------------------------------------------
401 //------------------------------------------------------------------------------
410 impl->h_array = *array; in CeedVectorGetArray_Sycl()
413 impl->d_array = *array; in CeedVectorGetArray_Sycl()
419 //------------------------------------------------------------------------------
421 //------------------------------------------------------------------------------
435 if (impl->h_array_borrowed) impl->h_array = impl->h_array_borrowed; in CeedVectorGetArrayWrite_Sycl()
436 else impl->h_array = impl->h_array_owned; in CeedVectorGetArrayWrite_Sycl()
439 if (impl->d_array_borrowed) impl->d_array = impl->d_array_borrowed; in CeedVectorGetArrayWrite_Sycl()
440 else impl->d_array = impl->d_array_owned; in CeedVectorGetArrayWrite_Sycl()
446 //------------------------------------------------------------------------------
448 //------------------------------------------------------------------------------
467 if (!data->sycl_queue.is_in_order()) e = {data->sycl_queue.ext_oneapi_submit_barrier()}; in CeedVectorNorm_Sycl()
472 …auto sumReduction = sycl::reduction(impl->reduction_norm, sycl::plus<>(), {sycl::property::reducti… in CeedVectorNorm_Sycl()
473 …data->sycl_queue.parallel_for(length, e, sumReduction, [=](sycl::id<1> i, auto &sum) { sum += abs(… in CeedVectorNorm_Sycl()
477 …auto sumReduction = sycl::reduction(impl->reduction_norm, sycl::plus<>(), {sycl::property::reducti… in CeedVectorNorm_Sycl()
478 …data->sycl_queue.parallel_for(length, e, sumReduction, [=](sycl::id<1> i, auto &sum) { sum += (d_a… in CeedVectorNorm_Sycl()
482 …auto maxReduction = sycl::reduction(impl->reduction_norm, sycl::maximum<>(), {sycl::property::redu… in CeedVectorNorm_Sycl()
483 …data->sycl_queue.parallel_for(length, e, maxReduction, [=](sycl::id<1> i, auto &max) { max.combine… in CeedVectorNorm_Sycl()
486 // L2 norm - square root over reduced value in CeedVectorNorm_Sycl()
487 if (type == CEED_NORM_2) *norm = sqrt(*impl->reduction_norm); in CeedVectorNorm_Sycl()
488 else *norm = *impl->reduction_norm; in CeedVectorNorm_Sycl()
493 //------------------------------------------------------------------------------
495 //------------------------------------------------------------------------------
503 //------------------------------------------------------------------------------
505 //------------------------------------------------------------------------------
516 //------------------------------------------------------------------------------
518 //------------------------------------------------------------------------------
532 …if (impl->d_array) CeedCallBackend(CeedDeviceReciprocal_Sycl(data->sycl_queue, impl->d_array, leng… in CeedVectorReciprocal_Sycl()
533 if (impl->h_array) CeedCallBackend(CeedHostReciprocal_Sycl(impl->h_array, length)); in CeedVectorReciprocal_Sycl()
537 //------------------------------------------------------------------------------
539 //------------------------------------------------------------------------------
545 //------------------------------------------------------------------------------
547 //------------------------------------------------------------------------------
556 //------------------------------------------------------------------------------
558 //------------------------------------------------------------------------------
572 …if (x_impl->d_array) CeedCallBackend(CeedDeviceScale_Sycl(data->sycl_queue, x_impl->d_array, alpha… in CeedVectorScale_Sycl()
573 if (x_impl->h_array) CeedCallBackend(CeedHostScale_Sycl(x_impl->h_array, alpha, length)); in CeedVectorScale_Sycl()
577 //------------------------------------------------------------------------------
579 //------------------------------------------------------------------------------
585 //------------------------------------------------------------------------------
587 //------------------------------------------------------------------------------
596 //------------------------------------------------------------------------------
598 //------------------------------------------------------------------------------
613 if (y_impl->d_array) { in CeedVectorAXPY_Sycl()
615 …CeedCallBackend(CeedDeviceAXPY_Sycl(data->sycl_queue, y_impl->d_array, alpha, x_impl->d_array, len… in CeedVectorAXPY_Sycl()
617 if (y_impl->h_array) { in CeedVectorAXPY_Sycl()
619 CeedCallBackend(CeedHostAXPY_Sycl(y_impl->h_array, alpha, x_impl->h_array, length)); in CeedVectorAXPY_Sycl()
624 //------------------------------------------------------------------------------
626 //------------------------------------------------------------------------------
632 //------------------------------------------------------------------------------
634 //------------------------------------------------------------------------------
643 //------------------------------------------------------------------------------
645 //------------------------------------------------------------------------------
661 if (!w_impl->d_array && !w_impl->h_array) { in CeedVectorPointwiseMult_Sycl()
664 if (w_impl->d_array) { in CeedVectorPointwiseMult_Sycl()
667 …CeedCallBackend(CeedDevicePointwiseMult_Sycl(data->sycl_queue, w_impl->d_array, x_impl->d_array, y… in CeedVectorPointwiseMult_Sycl()
669 if (w_impl->h_array) { in CeedVectorPointwiseMult_Sycl()
672 …CeedCallBackend(CeedHostPointwiseMult_Sycl(w_impl->h_array, x_impl->h_array, y_impl->h_array, leng… in CeedVectorPointwiseMult_Sycl()
677 //------------------------------------------------------------------------------
679 //------------------------------------------------------------------------------
690 CeedCallSycl(ceed, data->sycl_queue.wait_and_throw()); in CeedVectorDestroy_Sycl()
691 CeedCallSycl(ceed, sycl::free(impl->d_array_owned, data->sycl_context)); in CeedVectorDestroy_Sycl()
692 CeedCallSycl(ceed, sycl::free(impl->reduction_norm, data->sycl_context)); in CeedVectorDestroy_Sycl()
694 CeedCallBackend(CeedFree(&impl->h_array_owned)); in CeedVectorDestroy_Sycl()
700 //------------------------------------------------------------------------------
702 //------------------------------------------------------------------------------
711 CeedCallSycl(ceed, impl->reduction_norm = sycl::malloc_host<CeedScalar>(1, data->sycl_context)); in CeedVectorCreate_Sycl()