1d275d636SJeremy L Thompson // Copyright (c) 2017-2025, Lawrence Livermore National Security, LLC and other CEED contributors. 23d8e8822SJeremy L Thompson // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 37fcac036SJeremy L Thompson // 43d8e8822SJeremy L Thompson // SPDX-License-Identifier: BSD-2-Clause 57fcac036SJeremy L Thompson // 63d8e8822SJeremy L Thompson // This file is part of CEED: http://github.com/ceed 7509d4af6SJeremy L Thompson #pragma once 87fcac036SJeremy L Thompson 949aac155SJeremy L Thompson #include <ceed.h> 107fcac036SJeremy L Thompson #include <ceed/backend.h> 11ce5711eeSJeremy L Thompson #include <hip/hip_runtime.h> 120df8cb37SJeremy L Thompson #if (HIP_VERSION >= 50200000) 1349aac155SJeremy L Thompson #include <hipblas/hipblas.h> // IWYU pragma: export 140df8cb37SJeremy L Thompson #else 1549aac155SJeremy L Thompson #include <hipblas.h> // IWYU pragma: export 160df8cb37SJeremy L Thompson #endif 177fcac036SJeremy L Thompson 187fcac036SJeremy L Thompson #define QUOTE(...) #__VA_ARGS__ 197fcac036SJeremy L Thompson 207fcac036SJeremy L Thompson #define CeedChk_Hip(ceed, x) \ 217fcac036SJeremy L Thompson do { \ 227fcac036SJeremy L Thompson hipError_t hip_result = x; \ 237fcac036SJeremy L Thompson if (hip_result != hipSuccess) { \ 247fcac036SJeremy L Thompson const char *msg = hipGetErrorName(hip_result); \ 257fcac036SJeremy L Thompson return CeedError((ceed), CEED_ERROR_BACKEND, msg); \ 267fcac036SJeremy L Thompson } \ 277fcac036SJeremy L Thompson } while (0) 287fcac036SJeremy L Thompson 297fcac036SJeremy L Thompson #define CeedChk_Hipblas(ceed, x) \ 307fcac036SJeremy L Thompson do { \ 317fcac036SJeremy L Thompson hipblasStatus_t hipblas_result = x; \ 327fcac036SJeremy L Thompson if (hipblas_result != HIPBLAS_STATUS_SUCCESS) { \ 337fcac036SJeremy L Thompson const char *msg = hipblasGetErrorName(hipblas_result); \ 347fcac036SJeremy L Thompson return CeedError((ceed), CEED_ERROR_BACKEND, msg); \ 357fcac036SJeremy L Thompson } \ 367fcac036SJeremy L Thompson } while (0) 377fcac036SJeremy L Thompson 382b730f8bSJeremy L Thompson #define CeedCallHip(ceed, ...) \ 392b730f8bSJeremy L Thompson do { \ 402b730f8bSJeremy L Thompson hipError_t ierr_q_ = __VA_ARGS__; \ 412b730f8bSJeremy L Thompson CeedChk_Hip(ceed, ierr_q_); \ 426574a04fSJeremy L Thompson } while (0) 432b730f8bSJeremy L Thompson 442b730f8bSJeremy L Thompson #define CeedCallHipblas(ceed, ...) \ 452b730f8bSJeremy L Thompson do { \ 462b730f8bSJeremy L Thompson hipblasStatus_t ierr_q_ = __VA_ARGS__; \ 472b730f8bSJeremy L Thompson CeedChk_Hipblas(ceed, ierr_q_); \ 486574a04fSJeremy L Thompson } while (0) 492b730f8bSJeremy L Thompson 502b730f8bSJeremy L Thompson #define CASE(name) \ 512b730f8bSJeremy L Thompson case name: \ 522b730f8bSJeremy L Thompson return #name 537fcac036SJeremy L Thompson // LCOV_EXCL_START 547fcac036SJeremy L Thompson CEED_UNUSED static const char *hipblasGetErrorName(hipblasStatus_t error) { 557fcac036SJeremy L Thompson switch (error) { 567fcac036SJeremy L Thompson CASE(HIPBLAS_STATUS_SUCCESS); 577fcac036SJeremy L Thompson CASE(HIPBLAS_STATUS_NOT_INITIALIZED); 587fcac036SJeremy L Thompson CASE(HIPBLAS_STATUS_ALLOC_FAILED); 597fcac036SJeremy L Thompson CASE(HIPBLAS_STATUS_INVALID_VALUE); 607fcac036SJeremy L Thompson CASE(HIPBLAS_STATUS_ARCH_MISMATCH); 617fcac036SJeremy L Thompson CASE(HIPBLAS_STATUS_MAPPING_ERROR); 627fcac036SJeremy L Thompson CASE(HIPBLAS_STATUS_EXECUTION_FAILED); 637fcac036SJeremy L Thompson CASE(HIPBLAS_STATUS_INTERNAL_ERROR); 642b730f8bSJeremy L Thompson default: 652b730f8bSJeremy L Thompson return "HIPBLAS_STATUS_UNKNOWN_ERROR"; 667fcac036SJeremy L Thompson } 677fcac036SJeremy L Thompson } 687fcac036SJeremy L Thompson // LCOV_EXCL_STOP 697fcac036SJeremy L Thompson 707fcac036SJeremy L Thompson typedef struct { 710d0321e0SJeremy L Thompson int device_id; 720d0321e0SJeremy L Thompson hipblasHandle_t hipblas_handle; 73b2165e7aSSebastian Grimberg struct hipDeviceProp_t device_prop; 74b2165e7aSSebastian Grimberg int opt_block_size; 75*a3b195efSJeremy L Thompson int has_unified_addressing; 767fcac036SJeremy L Thompson } Ceed_Hip; 777fcac036SJeremy L Thompson 78eb7e6cafSJeremy L Thompson CEED_INTERN int CeedInit_Hip(Ceed ceed, const char *resource); 797fcac036SJeremy L Thompson 807fcac036SJeremy L Thompson CEED_INTERN int CeedDestroy_Hip(Ceed ceed); 817fcac036SJeremy L Thompson 82f5d1e504SJeremy L Thompson CEED_INTERN int CeedSetDeviceBoolArray_Hip(Ceed ceed, const bool *source_array, CeedCopyMode copy_mode, CeedSize num_values, 83f5d1e504SJeremy L Thompson const bool **target_array_owned, const bool **target_array_borrowed, const bool **target_array); 84f5d1e504SJeremy L Thompson CEED_INTERN int CeedSetDeviceCeedInt8Array_Hip(Ceed ceed, const CeedInt8 *source_array, CeedCopyMode copy_mode, CeedSize num_values, 85f5d1e504SJeremy L Thompson const CeedInt8 **target_array_owned, const CeedInt8 **target_array_borrowed, 86f5d1e504SJeremy L Thompson const CeedInt8 **target_array); 87f5d1e504SJeremy L Thompson CEED_INTERN int CeedSetDeviceCeedIntArray_Hip(Ceed ceed, const CeedInt *source_array, CeedCopyMode copy_mode, CeedSize num_values, 88f5d1e504SJeremy L Thompson const CeedInt **target_array_owned, const CeedInt **target_array_borrowed, 89f5d1e504SJeremy L Thompson const CeedInt **target_array); 90f5d1e504SJeremy L Thompson CEED_INTERN int CeedSetDeviceCeedScalarArray_Hip(Ceed ceed, const CeedScalar *source_array, CeedCopyMode copy_mode, CeedSize num_values, 91f5d1e504SJeremy L Thompson const CeedScalar **target_array_owned, const CeedScalar **target_array_borrowed, 92f5d1e504SJeremy L Thompson const CeedScalar **target_array); 93