1*3d8e8822SJeremy L Thompson // Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors. 2*3d8e8822SJeremy L Thompson // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 30d0321e0SJeremy L Thompson // 4*3d8e8822SJeremy L Thompson // SPDX-License-Identifier: BSD-2-Clause 50d0321e0SJeremy L Thompson // 6*3d8e8822SJeremy 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 430d0321e0SJeremy L Thompson if (strcmp(resource, "/gpu/hip/ref")) 440d0321e0SJeremy L Thompson // LCOV_EXCL_START 450d0321e0SJeremy L Thompson return CeedError(ceed, CEED_ERROR_BACKEND, 460d0321e0SJeremy L Thompson "Hip backend cannot use resource: %s", resource); 470d0321e0SJeremy L Thompson // LCOV_EXCL_STOP 480d0321e0SJeremy L Thompson ierr = CeedSetDeterministic(ceed, true); CeedChk(ierr); 490d0321e0SJeremy L Thompson 500d0321e0SJeremy L Thompson Ceed_Hip *data; 510d0321e0SJeremy L Thompson ierr = CeedCalloc(1, &data); CeedChkBackend(ierr); 520d0321e0SJeremy L Thompson ierr = CeedSetData(ceed, data); CeedChkBackend(ierr); 530d0321e0SJeremy L Thompson ierr = CeedHipInit(ceed, resource); CeedChkBackend(ierr); 540d0321e0SJeremy L Thompson 550d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "GetPreferredMemType", 560d0321e0SJeremy L Thompson CeedGetPreferredMemType_Hip); CeedChkBackend(ierr); 570d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "VectorCreate", 580d0321e0SJeremy L Thompson CeedVectorCreate_Hip); CeedChkBackend(ierr); 590d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "BasisCreateTensorH1", 600d0321e0SJeremy L Thompson CeedBasisCreateTensorH1_Hip); CeedChkBackend(ierr); 610d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "BasisCreateH1", 620d0321e0SJeremy L Thompson CeedBasisCreateH1_Hip); CeedChkBackend(ierr); 630d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "ElemRestrictionCreate", 640d0321e0SJeremy L Thompson CeedElemRestrictionCreate_Hip); CeedChkBackend(ierr); 650d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, 660d0321e0SJeremy L Thompson "ElemRestrictionCreateBlocked", 670d0321e0SJeremy L Thompson CeedElemRestrictionCreateBlocked_Hip); 680d0321e0SJeremy L Thompson CeedChkBackend(ierr); 690d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "QFunctionCreate", 700d0321e0SJeremy L Thompson CeedQFunctionCreate_Hip); CeedChkBackend(ierr); 710d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "QFunctionContextCreate", 720d0321e0SJeremy L Thompson CeedQFunctionContextCreate_Hip); CeedChkBackend(ierr); 730d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "OperatorCreate", 740d0321e0SJeremy L Thompson CeedOperatorCreate_Hip); CeedChkBackend(ierr); 750d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "CompositeOperatorCreate", 760d0321e0SJeremy L Thompson CeedCompositeOperatorCreate_Hip); CeedChkBackend(ierr); 770d0321e0SJeremy L Thompson ierr = CeedSetBackendFunction(ceed, "Ceed", ceed, "Destroy", 780d0321e0SJeremy L Thompson CeedDestroy_Hip); CeedChkBackend(ierr); 790d0321e0SJeremy L Thompson return CEED_ERROR_SUCCESS; 800d0321e0SJeremy L Thompson } 810d0321e0SJeremy L Thompson 820d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 830d0321e0SJeremy L Thompson // Backend Register 840d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 850d0321e0SJeremy L Thompson CEED_INTERN int CeedRegister_Hip(void) { 860d0321e0SJeremy L Thompson return CeedRegister("/gpu/hip/ref", CeedInit_Hip, 40); 870d0321e0SJeremy L Thompson } 880d0321e0SJeremy L Thompson //------------------------------------------------------------------------------ 89