1 // Copyright (c) 2017-2025, 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.h> 9 #include <ceed/backend.h> 10 #include <ceed/jit-source/hip/hip-types.h> 11 #include <stddef.h> 12 #include <hip/hiprtc.h> 13 14 #include "../hip/ceed-hip-common.h" 15 #include "../hip/ceed-hip-compile.h" 16 #include "ceed-hip-gen-operator-build.h" 17 #include "ceed-hip-gen.h" 18 19 //------------------------------------------------------------------------------ 20 // Destroy operator 21 //------------------------------------------------------------------------------ 22 static int CeedOperatorDestroy_Hip_gen(CeedOperator op) { 23 Ceed ceed; 24 CeedOperator_Hip_gen *impl; 25 26 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 27 CeedCallBackend(CeedOperatorGetData(op, &impl)); 28 if (impl->points.num_per_elem) CeedCallHip(ceed, hipFree((void **)impl->points.num_per_elem)); 29 CeedCallBackend(CeedFree(&impl)); 30 CeedCallBackend(CeedDestroy(&ceed)); 31 return CEED_ERROR_SUCCESS; 32 } 33 34 //------------------------------------------------------------------------------ 35 // Apply and add to output 36 //------------------------------------------------------------------------------ 37 static int CeedOperatorApplyAddCore_Hip_gen(CeedOperator op, hipStream_t stream, const CeedScalar *input_arr, CeedScalar *output_arr, 38 bool *is_run_good, CeedRequest *request) { 39 bool is_at_points, is_tensor; 40 Ceed ceed; 41 CeedInt num_elem, num_input_fields, num_output_fields; 42 CeedEvalMode eval_mode; 43 CeedQFunctionField *qf_input_fields, *qf_output_fields; 44 CeedQFunction_Hip_gen *qf_data; 45 CeedQFunction qf; 46 CeedOperatorField *op_input_fields, *op_output_fields; 47 CeedOperator_Hip_gen *data; 48 49 // Creation of the operator 50 CeedCallBackend(CeedOperatorBuildKernel_Hip_gen(op, is_run_good)); 51 if (!(*is_run_good)) return CEED_ERROR_SUCCESS; 52 53 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 54 CeedCallBackend(CeedOperatorGetData(op, &data)); 55 CeedCallBackend(CeedOperatorGetQFunction(op, &qf)); 56 CeedCallBackend(CeedQFunctionGetData(qf, &qf_data)); 57 CeedCallBackend(CeedOperatorGetNumElements(op, &num_elem)); 58 CeedCallBackend(CeedOperatorGetFields(op, &num_input_fields, &op_input_fields, &num_output_fields, &op_output_fields)); 59 CeedCallBackend(CeedQFunctionGetFields(qf, NULL, &qf_input_fields, NULL, &qf_output_fields)); 60 61 // Input vectors 62 for (CeedInt i = 0; i < num_input_fields; i++) { 63 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_input_fields[i], &eval_mode)); 64 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip 65 data->fields.inputs[i] = NULL; 66 } else { 67 bool is_active; 68 CeedVector vec; 69 70 // Get input vector 71 CeedCallBackend(CeedOperatorFieldGetVector(op_input_fields[i], &vec)); 72 is_active = vec == CEED_VECTOR_ACTIVE; 73 if (is_active) data->fields.inputs[i] = input_arr; 74 else CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->fields.inputs[i])); 75 CeedCallBackend(CeedVectorDestroy(&vec)); 76 } 77 } 78 79 // Output vectors 80 for (CeedInt i = 0; i < num_output_fields; i++) { 81 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_output_fields[i], &eval_mode)); 82 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip 83 data->fields.outputs[i] = NULL; 84 } else { 85 bool is_active; 86 CeedVector vec; 87 88 // Get output vector 89 CeedCallBackend(CeedOperatorFieldGetVector(op_output_fields[i], &vec)); 90 is_active = vec == CEED_VECTOR_ACTIVE; 91 if (is_active) data->fields.outputs[i] = output_arr; 92 else CeedCallBackend(CeedVectorGetArray(vec, CEED_MEM_DEVICE, &data->fields.outputs[i])); 93 CeedCallBackend(CeedVectorDestroy(&vec)); 94 } 95 } 96 97 // Point coordinates, if needed 98 CeedCallBackend(CeedOperatorIsAtPoints(op, &is_at_points)); 99 if (is_at_points) { 100 // Coords 101 CeedVector vec; 102 103 CeedCallBackend(CeedOperatorAtPointsGetPoints(op, NULL, &vec)); 104 CeedCallBackend(CeedVectorGetArrayRead(vec, CEED_MEM_DEVICE, &data->points.coords)); 105 CeedCallBackend(CeedVectorDestroy(&vec)); 106 107 // Points per elem 108 if (num_elem != data->points.num_elem) { 109 CeedInt *points_per_elem; 110 const CeedInt num_bytes = num_elem * sizeof(CeedInt); 111 CeedElemRestriction rstr_points = NULL; 112 113 data->points.num_elem = num_elem; 114 CeedCallBackend(CeedOperatorAtPointsGetPoints(op, &rstr_points, NULL)); 115 CeedCallBackend(CeedCalloc(num_elem, &points_per_elem)); 116 for (CeedInt e = 0; e < num_elem; e++) { 117 CeedInt num_points_elem; 118 119 CeedCallBackend(CeedElemRestrictionGetNumPointsInElement(rstr_points, e, &num_points_elem)); 120 points_per_elem[e] = num_points_elem; 121 } 122 if (data->points.num_per_elem) CeedCallHip(ceed, hipFree((void **)data->points.num_per_elem)); 123 CeedCallHip(ceed, hipMalloc((void **)&data->points.num_per_elem, num_bytes)); 124 CeedCallHip(ceed, hipMemcpy((void *)data->points.num_per_elem, points_per_elem, num_bytes, hipMemcpyHostToDevice)); 125 CeedCallBackend(CeedElemRestrictionDestroy(&rstr_points)); 126 CeedCallBackend(CeedFree(&points_per_elem)); 127 } 128 } 129 130 // Get context data 131 CeedCallBackend(CeedQFunctionGetInnerContextData(qf, CEED_MEM_DEVICE, &qf_data->d_c)); 132 133 // Apply operator 134 void *opargs[] = {(void *)&num_elem, &qf_data->d_c, &data->indices, &data->fields, &data->B, &data->G, &data->W, &data->points}; 135 const CeedInt dim = data->dim; 136 const CeedInt Q_1d = data->Q_1d; 137 const CeedInt P_1d = data->max_P_1d; 138 const CeedInt thread_1d = CeedIntMax(Q_1d, P_1d); 139 140 CeedCallBackend(CeedOperatorHasTensorBases(op, &is_tensor)); 141 CeedInt block_sizes[3] = {thread_1d, ((!is_tensor || dim == 1) ? 1 : thread_1d), -1}; 142 143 if (is_tensor) { 144 CeedCallBackend(BlockGridCalculate_Hip_gen(is_tensor ? dim : 1, num_elem, P_1d, Q_1d, block_sizes)); 145 if (is_at_points) block_sizes[2] = 1; 146 } else { 147 CeedInt elems_per_block = 64 * thread_1d > 256 ? 256 / thread_1d : 64; 148 149 elems_per_block = elems_per_block > 0 ? elems_per_block : 1; 150 block_sizes[2] = elems_per_block; 151 } 152 if (dim == 1 || !is_tensor) { 153 CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); 154 CeedInt sharedMem = block_sizes[2] * thread_1d * sizeof(CeedScalar); 155 156 CeedCallBackend( 157 CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_sizes[1], block_sizes[2], sharedMem, is_run_good, opargs)); 158 } else if (dim == 2) { 159 CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); 160 CeedInt sharedMem = block_sizes[2] * thread_1d * thread_1d * sizeof(CeedScalar); 161 162 CeedCallBackend( 163 CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_sizes[1], block_sizes[2], sharedMem, is_run_good, opargs)); 164 } else if (dim == 3) { 165 CeedInt grid = num_elem / block_sizes[2] + ((num_elem / block_sizes[2] * block_sizes[2] < num_elem) ? 1 : 0); 166 CeedInt sharedMem = block_sizes[2] * thread_1d * thread_1d * sizeof(CeedScalar); 167 168 CeedCallBackend( 169 CeedTryRunKernelDimShared_Hip(ceed, data->op, stream, grid, block_sizes[0], block_sizes[1], block_sizes[2], sharedMem, is_run_good, opargs)); 170 } 171 172 // Restore input arrays 173 for (CeedInt i = 0; i < num_input_fields; i++) { 174 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_input_fields[i], &eval_mode)); 175 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip 176 } else { 177 bool is_active; 178 CeedVector vec; 179 180 CeedCallBackend(CeedOperatorFieldGetVector(op_input_fields[i], &vec)); 181 is_active = vec == CEED_VECTOR_ACTIVE; 182 if (!is_active) CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->fields.inputs[i])); 183 CeedCallBackend(CeedVectorDestroy(&vec)); 184 } 185 } 186 187 // Restore output arrays 188 for (CeedInt i = 0; i < num_output_fields; i++) { 189 CeedCallBackend(CeedQFunctionFieldGetEvalMode(qf_output_fields[i], &eval_mode)); 190 if (eval_mode == CEED_EVAL_WEIGHT) { // Skip 191 } else { 192 bool is_active; 193 CeedVector vec; 194 195 CeedCallBackend(CeedOperatorFieldGetVector(op_output_fields[i], &vec)); 196 is_active = vec == CEED_VECTOR_ACTIVE; 197 if (!is_active) CeedCallBackend(CeedVectorRestoreArray(vec, &data->fields.outputs[i])); 198 CeedCallBackend(CeedVectorDestroy(&vec)); 199 } 200 } 201 202 // Restore point coordinates, if needed 203 if (is_at_points) { 204 CeedVector vec; 205 206 CeedCallBackend(CeedOperatorAtPointsGetPoints(op, NULL, &vec)); 207 CeedCallBackend(CeedVectorRestoreArrayRead(vec, &data->points.coords)); 208 CeedCallBackend(CeedVectorDestroy(&vec)); 209 } 210 211 // Restore context data 212 CeedCallBackend(CeedQFunctionRestoreInnerContextData(qf, &qf_data->d_c)); 213 214 // Cleanup 215 CeedCallBackend(CeedDestroy(&ceed)); 216 CeedCallBackend(CeedQFunctionDestroy(&qf)); 217 if (!(*is_run_good)) data->use_fallback = true; 218 return CEED_ERROR_SUCCESS; 219 } 220 221 static int CeedOperatorApplyAdd_Hip_gen(CeedOperator op, CeedVector input_vec, CeedVector output_vec, CeedRequest *request) { 222 bool is_run_good = false; 223 const CeedScalar *input_arr = NULL; 224 CeedScalar *output_arr = NULL; 225 226 // Try to run kernel 227 if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(input_vec, CEED_MEM_DEVICE, &input_arr)); 228 if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArray(output_vec, CEED_MEM_DEVICE, &output_arr)); 229 CeedCallBackend(CeedOperatorApplyAddCore_Hip_gen(op, NULL, input_arr, output_arr, &is_run_good, request)); 230 if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArrayRead(input_vec, &input_arr)); 231 if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArray(output_vec, &output_arr)); 232 233 // Fallback on unsuccessful run 234 if (!is_run_good) { 235 CeedOperator op_fallback; 236 237 CeedDebug256(CeedOperatorReturnCeed(op), CEED_DEBUG_COLOR_SUCCESS, "Falling back to /gpu/hip/ref CeedOperator"); 238 CeedCallBackend(CeedOperatorGetFallback(op, &op_fallback)); 239 CeedCallBackend(CeedOperatorApplyAdd(op_fallback, input_vec, output_vec, request)); 240 } 241 return CEED_ERROR_SUCCESS; 242 } 243 244 static int CeedOperatorApplyAddComposite_Hip_gen(CeedOperator op, CeedVector input_vec, CeedVector output_vec, CeedRequest *request) { 245 bool is_run_good[CEED_COMPOSITE_MAX] = {false}; 246 CeedInt num_suboperators; 247 const CeedScalar *input_arr = NULL; 248 CeedScalar *output_arr = NULL; 249 Ceed ceed; 250 CeedOperator *sub_operators; 251 252 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 253 CeedCall(CeedCompositeOperatorGetNumSub(op, &num_suboperators)); 254 CeedCall(CeedCompositeOperatorGetSubList(op, &sub_operators)); 255 if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArrayRead(input_vec, CEED_MEM_DEVICE, &input_arr)); 256 if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorGetArray(output_vec, CEED_MEM_DEVICE, &output_arr)); 257 for (CeedInt i = 0; i < num_suboperators; i++) { 258 CeedInt num_elem = 0; 259 260 CeedCall(CeedOperatorGetNumElements(sub_operators[i], &num_elem)); 261 if (num_elem > 0) { 262 hipStream_t stream = NULL; 263 264 CeedCallHip(ceed, hipStreamCreate(&stream)); 265 CeedCallBackend(CeedOperatorApplyAddCore_Hip_gen(sub_operators[i], stream, input_arr, output_arr, &is_run_good[i], request)); 266 CeedCallHip(ceed, hipStreamDestroy(stream)); 267 } 268 } 269 if (input_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArrayRead(input_vec, &input_arr)); 270 if (output_vec != CEED_VECTOR_NONE) CeedCallBackend(CeedVectorRestoreArray(output_vec, &output_arr)); 271 CeedCallHip(ceed, hipDeviceSynchronize()); 272 273 // Fallback on unsuccessful run 274 for (CeedInt i = 0; i < num_suboperators; i++) { 275 if (!is_run_good[i]) { 276 CeedOperator op_fallback; 277 278 CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Falling back to /gpu/hip/ref CeedOperator"); 279 CeedCallBackend(CeedOperatorGetFallback(sub_operators[i], &op_fallback)); 280 CeedCallBackend(CeedOperatorApplyAdd(op_fallback, input_vec, output_vec, request)); 281 } 282 } 283 CeedCallBackend(CeedDestroy(&ceed)); 284 return CEED_ERROR_SUCCESS; 285 } 286 287 //------------------------------------------------------------------------------ 288 // Create operator 289 //------------------------------------------------------------------------------ 290 int CeedOperatorCreate_Hip_gen(CeedOperator op) { 291 bool is_composite; 292 Ceed ceed; 293 CeedOperator_Hip_gen *impl; 294 295 CeedCallBackend(CeedOperatorGetCeed(op, &ceed)); 296 CeedCallBackend(CeedCalloc(1, &impl)); 297 CeedCallBackend(CeedOperatorSetData(op, impl)); 298 CeedCall(CeedOperatorIsComposite(op, &is_composite)); 299 if (is_composite) { 300 CeedCallBackend(CeedSetBackendFunction(ceed, "Operator", op, "ApplyAddComposite", CeedOperatorApplyAddComposite_Hip_gen)); 301 } else { 302 CeedCallBackend(CeedSetBackendFunction(ceed, "Operator", op, "ApplyAdd", CeedOperatorApplyAdd_Hip_gen)); 303 } 304 CeedCallBackend(CeedSetBackendFunction(ceed, "Operator", op, "Destroy", CeedOperatorDestroy_Hip_gen)); 305 CeedCallBackend(CeedDestroy(&ceed)); 306 return CEED_ERROR_SUCCESS; 307 } 308 309 //------------------------------------------------------------------------------ 310