13d8e8822SJeremy L Thompson // Copyright (c) 2017-2022, 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. 30d0321e0SJeremy L Thompson // 43d8e8822SJeremy L Thompson // SPDX-License-Identifier: BSD-2-Clause 50d0321e0SJeremy L Thompson // 63d8e8822SJeremy L Thompson // This file is part of CEED: http://github.com/ceed 70d0321e0SJeremy L Thompson 80d0321e0SJeremy L Thompson #include <ceed/ceed.h> 90d0321e0SJeremy L Thompson #include <ceed/backend.h> 100d0321e0SJeremy L Thompson #include <string.h> 110d0321e0SJeremy L Thompson #include <stdlib.h> 120d0321e0SJeremy L Thompson #include "ceed-hip-ref.h" 130d0321e0SJeremy L Thompson 140d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 150d0321e0SJeremy L Thompson // HIP preferred MemType 160d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 170d0321e0SJeremy L Thompson static int CeedGetPreferredMemType_Hip(CeedMemType *type) { 180d0321e0SJeremy L Thompson *type = CEED_MEM_DEVICE; 190d0321e0SJeremy L Thompson return CEED_ERROR_SUCCESS; 200d0321e0SJeremy L Thompson } 210d0321e0SJeremy L Thompson 220d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 230d0321e0SJeremy L Thompson // Get hipBLAS handle 240d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 250d0321e0SJeremy L Thompson int CeedHipGetHipblasHandle(Ceed ceed, hipblasHandle_t *handle) { 260d0321e0SJeremy L Thompson int ierr; 270d0321e0SJeremy L Thompson Ceed_Hip *data; 280d0321e0SJeremy L Thompson ierr = CeedGetData(ceed, &data); CeedChkBackend(ierr); 290d0321e0SJeremy L Thompson 300d0321e0SJeremy L Thompson if (!data->hipblas_handle) { 310d0321e0SJeremy L Thompson ierr = hipblasCreate(&data->hipblas_handle); CeedChk_Hipblas(ceed, ierr); 320d0321e0SJeremy L Thompson } 330d0321e0SJeremy L Thompson *handle = data->hipblas_handle; 340d0321e0SJeremy L Thompson return CEED_ERROR_SUCCESS; 350d0321e0SJeremy L Thompson } 360d0321e0SJeremy L Thompson 370d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 380d0321e0SJeremy L Thompson // Backend Init 390d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 400d0321e0SJeremy L Thompson static int CeedInit_Hip(const char *resource, Ceed ceed) { 410d0321e0SJeremy L Thompson int ierr; 420d0321e0SJeremy L Thompson 43*b11824b3SJeremy L Thompson char *resource_root; 44*b11824b3SJeremy L Thompson ierr = CeedHipGetResourceRoot(ceed, resource, &resource_root); 45*b11824b3SJeremy L Thompson CeedChkBackend(ierr); 46*b11824b3SJeremy L Thompson if (strcmp(resource_root, "/gpu/hip/ref")) 470d0321e0SJeremy L Thompson // LCOV_EXCL_START 480d0321e0SJeremy L Thompson return CeedError(ceed, CEED_ERROR_BACKEND, 490d0321e0SJeremy L Thompson "Hip backend cannot use resource: %s", resource); 500d0321e0SJeremy L Thompson // LCOV_EXCL_STOP 51*b11824b3SJeremy L Thompson ierr = CeedFree(&resource_root); CeedChkBackend(ierr); 52*b11824b3SJeremy L Thompson ierr = CeedSetDeterministic(ceed, true); CeedChkBackend(ierr); 530d0321e0SJeremy L Thompson 540d0321e0SJeremy L Thompson Ceed_Hip *data; 550d0321e0SJeremy L Thompson ierr = CeedCalloc(1, &data); CeedChkBackend(ierr); 560d0321e0SJeremy L Thompson ierr = CeedSetData(ceed, data); CeedChkBackend(ierr); 570d0321e0SJeremy L Thompson ierr = CeedHipInit(ceed, resource); CeedChkBackend(ierr); 580d0321e0SJeremy L Thompson 590d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "GetPreferredMemType", 600d0321e0SJeremy L Thompson CeedGetPreferredMemType_Hip); CeedChkBackend(ierr); 610d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "VectorCreate", 620d0321e0SJeremy L Thompson CeedVectorCreate_Hip); CeedChkBackend(ierr); 630d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "BasisCreateTensorH1", 640d0321e0SJeremy L Thompson CeedBasisCreateTensorH1_Hip); CeedChkBackend(ierr); 650d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "BasisCreateH1", 660d0321e0SJeremy L Thompson CeedBasisCreateH1_Hip); CeedChkBackend(ierr); 670d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "ElemRestrictionCreate", 680d0321e0SJeremy L Thompson CeedElemRestrictionCreate_Hip); CeedChkBackend(ierr); 690d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, 700d0321e0SJeremy L Thompson "ElemRestrictionCreateBlocked", 710d0321e0SJeremy L Thompson CeedElemRestrictionCreateBlocked_Hip); 720d0321e0SJeremy L Thompson CeedChkBackend(ierr); 730d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "QFunctionCreate", 740d0321e0SJeremy L Thompson CeedQFunctionCreate_Hip); CeedChkBackend(ierr); 750d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "QFunctionContextCreate", 760d0321e0SJeremy L Thompson CeedQFunctionContextCreate_Hip); CeedChkBackend(ierr); 770d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "OperatorCreate", 780d0321e0SJeremy L Thompson CeedOperatorCreate_Hip); CeedChkBackend(ierr); 790d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "CompositeOperatorCreate", 800d0321e0SJeremy L Thompson CeedCompositeOperatorCreate_Hip); CeedChkBackend(ierr); 810d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "Destroy", 820d0321e0SJeremy L Thompson CeedDestroy_Hip); CeedChkBackend(ierr); 830d0321e0SJeremy L Thompson return CEED_ERROR_SUCCESS; 840d0321e0SJeremy L Thompson } 850d0321e0SJeremy L Thompson 860d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 870d0321e0SJeremy L Thompson // Backend Register 880d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 890d0321e0SJeremy L Thompson CEED_INTERN int CeedRegister_Hip(void) { 900d0321e0SJeremy L Thompson return CeedRegister("/gpu/hip/ref", CeedInit_Hip, 40); 910d0321e0SJeremy L Thompson } 920d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 93