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