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.
3 //
4 // SPDX-License-Identifier: BSD-2-Clause
5 //
6 // This file is part of CEED: http://github.com/ceed
7
8 #include "ceed-hip-common.h"
9
10 #include <ceed.h>
11 #include <ceed/backend.h>
12 #include <stdlib.h>
13 #include <string.h>
14
15 //------------------------------------------------------------------------------
16 // Device information backend init
17 //------------------------------------------------------------------------------
CeedInit_Hip(Ceed ceed,const char * resource)18 int CeedInit_Hip(Ceed ceed, const char *resource) {
19 Ceed_Hip *data;
20 const char *device_spec = strstr(resource, ":device_id=");
21 const int device_id = (device_spec) ? atoi(device_spec + 11) : -1;
22 int current_device_id, xnack_value;
23 const char *xnack;
24
25 CeedCallHip(ceed, hipGetDevice(¤t_device_id));
26 if (device_id >= 0 && current_device_id != device_id) {
27 CeedCallHip(ceed, hipSetDevice(device_id));
28 current_device_id = device_id;
29 }
30
31 CeedCallBackend(CeedGetData(ceed, &data));
32 data->device_id = current_device_id;
33 CeedCallHip(ceed, hipGetDeviceProperties(&data->device_prop, current_device_id));
34 xnack = getenv("HSA_XNACK");
35 xnack_value = !!xnack ? atol(xnack) : 0;
36 data->has_unified_addressing = xnack_value > 0 ? data->device_prop.unifiedAddressing : 0;
37 if (data->has_unified_addressing) {
38 CeedDebug(ceed, "Using unified memory addressing");
39 }
40 data->opt_block_size = 256;
41 return CEED_ERROR_SUCCESS;
42 }
43
44 //------------------------------------------------------------------------------
45 // Backend Destroy
46 //------------------------------------------------------------------------------
CeedDestroy_Hip(Ceed ceed)47 int CeedDestroy_Hip(Ceed ceed) {
48 Ceed_Hip *data;
49
50 CeedCallBackend(CeedGetData(ceed, &data));
51 if (data->hipblas_handle) CeedCallHipblas(ceed, hipblasDestroy(data->hipblas_handle));
52 CeedCallBackend(CeedFree(&data));
53 return CEED_ERROR_SUCCESS;
54 }
55
56 //------------------------------------------------------------------------------
57 // Memory transfer utilities
58 //------------------------------------------------------------------------------
CeedSetDeviceGenericArray_Hip(Ceed ceed,const void * source_array,CeedCopyMode copy_mode,size_t size_unit,CeedSize num_values,void * target_array_owned,void * target_array_borrowed,void * target_array)59 static inline int CeedSetDeviceGenericArray_Hip(Ceed ceed, const void *source_array, CeedCopyMode copy_mode, size_t size_unit, CeedSize num_values,
60 void *target_array_owned, void *target_array_borrowed, void *target_array) {
61 switch (copy_mode) {
62 case CEED_COPY_VALUES:
63 if (!*(void **)target_array) {
64 if (*(void **)target_array_borrowed) {
65 *(void **)target_array = *(void **)target_array_borrowed;
66 } else {
67 if (!*(void **)target_array_owned) CeedCallHip(ceed, hipMalloc(target_array_owned, size_unit * num_values));
68 *(void **)target_array = *(void **)target_array_owned;
69 }
70 }
71 if (source_array) CeedCallHip(ceed, hipMemcpy(*(void **)target_array, source_array, size_unit * num_values, hipMemcpyDeviceToDevice));
72 break;
73 case CEED_OWN_POINTER:
74 CeedCallHip(ceed, hipFree(*(void **)target_array_owned));
75 *(void **)target_array_owned = (void *)source_array;
76 *(void **)target_array_borrowed = NULL;
77 *(void **)target_array = *(void **)target_array_owned;
78 break;
79 case CEED_USE_POINTER:
80 CeedCallHip(ceed, hipFree(*(void **)target_array_owned));
81 *(void **)target_array_owned = NULL;
82 *(void **)target_array_borrowed = (void *)source_array;
83 *(void **)target_array = *(void **)target_array_borrowed;
84 }
85 return CEED_ERROR_SUCCESS;
86 }
87
CeedSetDeviceBoolArray_Hip(Ceed ceed,const bool * source_array,CeedCopyMode copy_mode,CeedSize num_values,const bool ** target_array_owned,const bool ** target_array_borrowed,const bool ** target_array)88 int CeedSetDeviceBoolArray_Hip(Ceed ceed, const bool *source_array, CeedCopyMode copy_mode, CeedSize num_values, const bool **target_array_owned,
89 const bool **target_array_borrowed, const bool **target_array) {
90 CeedCallBackend(CeedSetDeviceGenericArray_Hip(ceed, source_array, copy_mode, sizeof(bool), num_values, target_array_owned, target_array_borrowed,
91 target_array));
92 return CEED_ERROR_SUCCESS;
93 }
94
CeedSetDeviceCeedInt8Array_Hip(Ceed ceed,const CeedInt8 * source_array,CeedCopyMode copy_mode,CeedSize num_values,const CeedInt8 ** target_array_owned,const CeedInt8 ** target_array_borrowed,const CeedInt8 ** target_array)95 int CeedSetDeviceCeedInt8Array_Hip(Ceed ceed, const CeedInt8 *source_array, CeedCopyMode copy_mode, CeedSize num_values,
96 const CeedInt8 **target_array_owned, const CeedInt8 **target_array_borrowed, const CeedInt8 **target_array) {
97 CeedCallBackend(CeedSetDeviceGenericArray_Hip(ceed, source_array, copy_mode, sizeof(CeedInt8), num_values, target_array_owned,
98 target_array_borrowed, target_array));
99 return CEED_ERROR_SUCCESS;
100 }
101
CeedSetDeviceCeedIntArray_Hip(Ceed ceed,const CeedInt * source_array,CeedCopyMode copy_mode,CeedSize num_values,const CeedInt ** target_array_owned,const CeedInt ** target_array_borrowed,const CeedInt ** target_array)102 int CeedSetDeviceCeedIntArray_Hip(Ceed ceed, const CeedInt *source_array, CeedCopyMode copy_mode, CeedSize num_values,
103 const CeedInt **target_array_owned, const CeedInt **target_array_borrowed, const CeedInt **target_array) {
104 CeedCallBackend(CeedSetDeviceGenericArray_Hip(ceed, source_array, copy_mode, sizeof(CeedInt), num_values, target_array_owned, target_array_borrowed,
105 target_array));
106 return CEED_ERROR_SUCCESS;
107 }
108
CeedSetDeviceCeedScalarArray_Hip(Ceed ceed,const CeedScalar * source_array,CeedCopyMode copy_mode,CeedSize num_values,const CeedScalar ** target_array_owned,const CeedScalar ** target_array_borrowed,const CeedScalar ** target_array)109 int CeedSetDeviceCeedScalarArray_Hip(Ceed ceed, const CeedScalar *source_array, CeedCopyMode copy_mode, CeedSize num_values,
110 const CeedScalar **target_array_owned, const CeedScalar **target_array_borrowed,
111 const CeedScalar **target_array) {
112 CeedCallBackend(CeedSetDeviceGenericArray_Hip(ceed, source_array, copy_mode, sizeof(CeedScalar), num_values, target_array_owned,
113 target_array_borrowed, target_array));
114 return CEED_ERROR_SUCCESS;
115 }
116
117 //------------------------------------------------------------------------------
118