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