1 // Copyright (c) 2017-2018, Lawrence Livermore National Security, LLC. 2 // Produced at the Lawrence Livermore National Laboratory. LLNL-CODE-734707. 3 // All Rights reserved. See files LICENSE and NOTICE for details. 4 // 5 // This file is part of CEED, a collection of benchmarks, miniapps, software 6 // libraries and APIs for efficient high-order finite element and spectral 7 // element discretizations for exascale applications. For more information and 8 // source code availability see http://github.com/ceed. 9 // 10 // The CEED research is supported by the Exascale Computing Project 17-SC-20-SC, 11 // a collaborative effort of two U.S. Department of Energy organizations (Office 12 // of Science and the National Nuclear Security Administration) responsible for 13 // the planning and preparation of a capable exascale ecosystem, including 14 // software, applications, hardware, advanced system engineering and early 15 // testbed platforms, in support of the nation's exascale computing imperative. 16 17 #include <ceed/ceed.h> 18 #include <ceed/backend.h> 19 #include <hip/hip_runtime.h> 20 #include <string.h> 21 #include "ceed-hip-ref.h" 22 23 //------------------------------------------------------------------------------ 24 // * Bytes used 25 //------------------------------------------------------------------------------ 26 static inline size_t bytes(const CeedQFunctionContext ctx) { 27 int ierr; 28 size_t ctxsize; 29 ierr = CeedQFunctionContextGetContextSize(ctx, &ctxsize); CeedChkBackend(ierr); 30 return ctxsize; 31 } 32 33 //------------------------------------------------------------------------------ 34 // Sync host to device 35 //------------------------------------------------------------------------------ 36 static inline int CeedQFunctionContextSyncH2D_Hip( 37 const CeedQFunctionContext ctx) { 38 int ierr; 39 Ceed ceed; 40 ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); 41 CeedQFunctionContext_Hip *impl; 42 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 43 44 if (!impl->h_data) 45 // LCOV_EXCL_START 46 return CeedError(ceed, CEED_ERROR_BACKEND, 47 "No valid host data to sync to device"); 48 // LCOV_EXCL_STOP 49 50 if (impl->d_data_borrowed) { 51 impl->d_data = impl->d_data_borrowed; 52 } else if (impl->d_data_owned) { 53 impl->d_data = impl->d_data_owned; 54 } else { 55 ierr = hipMalloc((void **)&impl->d_data_owned, bytes(ctx)); 56 CeedChk_Hip(ceed, ierr); 57 impl->d_data = impl->d_data_owned; 58 } 59 60 ierr = hipMemcpy(impl->d_data, impl->h_data, bytes(ctx), 61 hipMemcpyHostToDevice); CeedChk_Hip(ceed, ierr); 62 63 return CEED_ERROR_SUCCESS; 64 } 65 66 //------------------------------------------------------------------------------ 67 // Sync device to host 68 //------------------------------------------------------------------------------ 69 static inline int CeedQFunctionContextSyncD2H_Hip( 70 const CeedQFunctionContext ctx) { 71 int ierr; 72 Ceed ceed; 73 ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); 74 CeedQFunctionContext_Hip *impl; 75 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 76 77 if (!impl->d_data) 78 // LCOV_EXCL_START 79 return CeedError(ceed, CEED_ERROR_BACKEND, 80 "No valid device data to sync to host"); 81 // LCOV_EXCL_STOP 82 83 if (impl->h_data_borrowed) { 84 impl->h_data = impl->h_data_borrowed; 85 } else if (impl->h_data_owned) { 86 impl->h_data = impl->h_data_owned; 87 } else { 88 ierr = CeedMalloc(bytes(ctx), &impl->h_data_owned); 89 CeedChkBackend(ierr); 90 impl->h_data = impl->h_data_owned; 91 } 92 93 ierr = hipMemcpy(impl->h_data, impl->d_data, bytes(ctx), 94 hipMemcpyDeviceToHost); CeedChk_Hip(ceed, ierr); 95 96 return CEED_ERROR_SUCCESS; 97 } 98 99 //------------------------------------------------------------------------------ 100 // Sync data of type 101 //------------------------------------------------------------------------------ 102 static inline int CeedQFunctionContextSync_Hip(const CeedQFunctionContext ctx, 103 CeedMemType mem_type) { 104 switch (mem_type) { 105 case CEED_MEM_HOST: return CeedQFunctionContextSyncD2H_Hip(ctx); 106 case CEED_MEM_DEVICE: return CeedQFunctionContextSyncH2D_Hip(ctx); 107 } 108 return CEED_ERROR_UNSUPPORTED; 109 } 110 111 //------------------------------------------------------------------------------ 112 // Set all pointers as invalid 113 //------------------------------------------------------------------------------ 114 static inline int CeedQFunctionContextSetAllInvalid_Hip( 115 const CeedQFunctionContext ctx) { 116 int ierr; 117 CeedQFunctionContext_Hip *impl; 118 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 119 120 impl->h_data = NULL; 121 impl->d_data = NULL; 122 123 return CEED_ERROR_SUCCESS; 124 } 125 126 //------------------------------------------------------------------------------ 127 // Check for valid data 128 //------------------------------------------------------------------------------ 129 static inline int CeedQFunctionContextHasValidData_Hip( 130 const CeedQFunctionContext ctx, bool *has_valid_data) { 131 int ierr; 132 CeedQFunctionContext_Hip *impl; 133 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 134 135 *has_valid_data = !!impl->h_data || !!impl->d_data; 136 137 return CEED_ERROR_SUCCESS; 138 } 139 140 //------------------------------------------------------------------------------ 141 // Check if ctx has borrowed data 142 //------------------------------------------------------------------------------ 143 static inline int CeedQFunctionContextHasBorrowedDataOfType_Hip( 144 const CeedQFunctionContext ctx, CeedMemType mem_type, 145 bool *has_borrowed_data_of_type) { 146 int ierr; 147 CeedQFunctionContext_Hip *impl; 148 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 149 150 switch (mem_type) { 151 case CEED_MEM_HOST: 152 *has_borrowed_data_of_type = !!impl->h_data_borrowed; 153 break; 154 case CEED_MEM_DEVICE: 155 *has_borrowed_data_of_type = !!impl->d_data_borrowed; 156 break; 157 } 158 159 return CEED_ERROR_SUCCESS; 160 } 161 162 //------------------------------------------------------------------------------ 163 // Check if data of given type needs sync 164 //------------------------------------------------------------------------------ 165 static inline int CeedQFunctionContextNeedSync_Hip( 166 const CeedQFunctionContext ctx, CeedMemType mem_type, bool *need_sync) { 167 int ierr; 168 CeedQFunctionContext_Hip *impl; 169 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 170 171 bool has_valid_data = true; 172 ierr = CeedQFunctionContextHasValidData_Hip(ctx, &has_valid_data); 173 CeedChkBackend(ierr); 174 switch (mem_type) { 175 case CEED_MEM_HOST: 176 *need_sync = has_valid_data && !impl->h_data; 177 break; 178 case CEED_MEM_DEVICE: 179 *need_sync = has_valid_data && !impl->d_data; 180 break; 181 } 182 183 return CEED_ERROR_SUCCESS; 184 } 185 186 //------------------------------------------------------------------------------ 187 // Set data from host 188 //------------------------------------------------------------------------------ 189 static int CeedQFunctionContextSetDataHost_Hip(const CeedQFunctionContext ctx, 190 const CeedCopyMode copy_mode, void *data) { 191 int ierr; 192 CeedQFunctionContext_Hip *impl; 193 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 194 195 ierr = CeedFree(&impl->h_data_owned); CeedChkBackend(ierr); 196 switch (copy_mode) { 197 case CEED_COPY_VALUES: { 198 ierr = CeedMalloc(bytes(ctx), &impl->h_data_owned); CeedChkBackend(ierr); 199 impl->h_data_borrowed = NULL; 200 impl->h_data = impl->h_data_owned; 201 memcpy(impl->h_data, data, bytes(ctx)); 202 } break; 203 case CEED_OWN_POINTER: 204 impl->h_data_owned = data; 205 impl->h_data_borrowed = NULL; 206 impl->h_data = data; 207 break; 208 case CEED_USE_POINTER: 209 impl->h_data_borrowed = data; 210 impl->h_data = data; 211 break; 212 } 213 214 return CEED_ERROR_SUCCESS; 215 } 216 217 //------------------------------------------------------------------------------ 218 // Set data from device 219 //------------------------------------------------------------------------------ 220 static int CeedQFunctionContextSetDataDevice_Hip(const CeedQFunctionContext ctx, 221 const CeedCopyMode copy_mode, void *data) { 222 int ierr; 223 Ceed ceed; 224 ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); 225 CeedQFunctionContext_Hip *impl; 226 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 227 228 ierr = hipFree(impl->d_data_owned); CeedChk_Hip(ceed, ierr); 229 impl->d_data_owned = NULL; 230 switch (copy_mode) { 231 case CEED_COPY_VALUES: 232 ierr = hipMalloc((void **)&impl->d_data_owned, bytes(ctx)); 233 CeedChk_Hip(ceed, ierr); 234 impl->d_data_borrowed = NULL; 235 impl->d_data = impl->d_data_owned; 236 ierr = hipMemcpy(impl->d_data, data, bytes(ctx), 237 hipMemcpyDeviceToDevice); CeedChk_Hip(ceed, ierr); 238 break; 239 case CEED_OWN_POINTER: 240 impl->d_data_owned = data; 241 impl->d_data_borrowed = NULL; 242 impl->d_data = data; 243 break; 244 case CEED_USE_POINTER: 245 impl->d_data_owned = NULL; 246 impl->d_data_borrowed = data; 247 impl->d_data = data; 248 break; 249 } 250 251 return CEED_ERROR_SUCCESS; 252 } 253 254 //------------------------------------------------------------------------------ 255 // Set the data used by a user context, 256 // freeing any previously allocated data if applicable 257 //------------------------------------------------------------------------------ 258 static int CeedQFunctionContextSetData_Hip(const CeedQFunctionContext ctx, 259 const CeedMemType mem_type, const CeedCopyMode copy_mode, void *data) { 260 int ierr; 261 Ceed ceed; 262 ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); 263 264 ierr = CeedQFunctionContextSetAllInvalid_Hip(ctx); CeedChkBackend(ierr); 265 switch (mem_type) { 266 case CEED_MEM_HOST: 267 return CeedQFunctionContextSetDataHost_Hip(ctx, copy_mode, data); 268 case CEED_MEM_DEVICE: 269 return CeedQFunctionContextSetDataDevice_Hip(ctx, copy_mode, data); 270 } 271 272 return CEED_ERROR_UNSUPPORTED; 273 } 274 275 //------------------------------------------------------------------------------ 276 // Take data 277 //------------------------------------------------------------------------------ 278 static int CeedQFunctionContextTakeData_Hip(const CeedQFunctionContext ctx, 279 const CeedMemType mem_type, void *data) { 280 int ierr; 281 Ceed ceed; 282 ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); 283 CeedQFunctionContext_Hip *impl; 284 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 285 286 // Sync data to requested mem_type 287 bool need_sync = false; 288 ierr = CeedQFunctionContextNeedSync_Hip(ctx, mem_type, &need_sync); 289 CeedChkBackend(ierr); 290 if (need_sync) { 291 ierr = CeedQFunctionContextSync_Hip(ctx, mem_type); CeedChkBackend(ierr); 292 } 293 294 // Update pointer 295 switch (mem_type) { 296 case CEED_MEM_HOST: 297 *(void **)data = impl->h_data_borrowed; 298 impl->h_data_borrowed = NULL; 299 impl->h_data = NULL; 300 break; 301 case CEED_MEM_DEVICE: 302 *(void **)data = impl->d_data_borrowed; 303 impl->d_data_borrowed = NULL; 304 impl->d_data = NULL; 305 break; 306 } 307 308 return CEED_ERROR_SUCCESS; 309 } 310 311 //------------------------------------------------------------------------------ 312 // Core logic for GetData. 313 // If a different memory type is most up to date, this will perform a copy 314 //------------------------------------------------------------------------------ 315 static int CeedQFunctionContextGetDataCore_Hip(const CeedQFunctionContext ctx, 316 const CeedMemType mem_type, void *data) { 317 int ierr; 318 Ceed ceed; 319 ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); 320 CeedQFunctionContext_Hip *impl; 321 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 322 323 // Sync data to requested mem_type 324 bool need_sync = false; 325 ierr = CeedQFunctionContextNeedSync_Hip(ctx, mem_type, &need_sync); 326 CeedChkBackend(ierr); 327 if (need_sync) { 328 ierr = CeedQFunctionContextSync_Hip(ctx, mem_type); CeedChkBackend(ierr); 329 } 330 331 // Sync data to requested mem_type and update pointer 332 switch (mem_type) { 333 case CEED_MEM_HOST: 334 *(void **)data = impl->h_data; 335 break; 336 case CEED_MEM_DEVICE: 337 *(void **)data = impl->d_data; 338 break; 339 } 340 341 return CEED_ERROR_SUCCESS; 342 } 343 344 //------------------------------------------------------------------------------ 345 // Get read-only access to the data 346 //------------------------------------------------------------------------------ 347 static int CeedQFunctionContextGetDataRead_Hip(const CeedQFunctionContext ctx, 348 const CeedMemType mem_type, void *data) { 349 return CeedQFunctionContextGetDataCore_Hip(ctx, mem_type, data); 350 } 351 352 //------------------------------------------------------------------------------ 353 // Get read/write access to the data 354 //------------------------------------------------------------------------------ 355 static int CeedQFunctionContextGetData_Hip(const CeedQFunctionContext ctx, 356 const CeedMemType mem_type, void *data) { 357 int ierr; 358 CeedQFunctionContext_Hip *impl; 359 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 360 361 ierr = CeedQFunctionContextGetDataCore_Hip(ctx, mem_type, data); 362 CeedChkBackend(ierr); 363 364 // Mark only pointer for requested memory as valid 365 ierr = CeedQFunctionContextSetAllInvalid_Hip(ctx); CeedChkBackend(ierr); 366 switch (mem_type) { 367 case CEED_MEM_HOST: 368 impl->h_data = *(void **)data; 369 break; 370 case CEED_MEM_DEVICE: 371 impl->d_data = *(void **)data; 372 break; 373 } 374 375 return CEED_ERROR_SUCCESS; 376 } 377 378 //------------------------------------------------------------------------------ 379 // Destroy the user context 380 //------------------------------------------------------------------------------ 381 static int CeedQFunctionContextDestroy_Hip(const CeedQFunctionContext ctx) { 382 int ierr; 383 Ceed ceed; 384 ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); 385 CeedQFunctionContext_Hip *impl; 386 ierr = CeedQFunctionContextGetBackendData(ctx, &impl); CeedChkBackend(ierr); 387 388 ierr = hipFree(impl->d_data_owned); CeedChk_Hip(ceed, ierr); 389 ierr = CeedFree(&impl->h_data_owned); CeedChkBackend(ierr); 390 ierr = CeedFree(&impl); CeedChkBackend(ierr); 391 392 return CEED_ERROR_SUCCESS; 393 } 394 395 //------------------------------------------------------------------------------ 396 // QFunctionContext Create 397 //------------------------------------------------------------------------------ 398 int CeedQFunctionContextCreate_Hip(CeedQFunctionContext ctx) { 399 int ierr; 400 CeedQFunctionContext_Hip *impl; 401 Ceed ceed; 402 ierr = CeedQFunctionContextGetCeed(ctx, &ceed); CeedChkBackend(ierr); 403 404 ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "HasValidData", 405 CeedQFunctionContextHasValidData_Hip); 406 CeedChkBackend(ierr); 407 ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, 408 "HasBorrowedDataOfType", 409 CeedQFunctionContextHasBorrowedDataOfType_Hip); 410 CeedChkBackend(ierr); 411 ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "SetData", 412 CeedQFunctionContextSetData_Hip); CeedChkBackend(ierr); 413 ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "TakeData", 414 CeedQFunctionContextTakeData_Hip); CeedChkBackend(ierr); 415 ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "GetData", 416 CeedQFunctionContextGetData_Hip); CeedChkBackend(ierr); 417 ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "GetDataRead", 418 CeedQFunctionContextGetDataRead_Hip); CeedChkBackend(ierr); 419 ierr = CeedSetBackendFunction(ceed, "QFunctionContext", ctx, "Destroy", 420 CeedQFunctionContextDestroy_Hip); CeedChkBackend(ierr); 421 422 ierr = CeedCalloc(1, &impl); CeedChkBackend(ierr); 423 ierr = CeedQFunctionContextSetBackendData(ctx, impl); CeedChkBackend(ierr); 424 425 return CEED_ERROR_SUCCESS; 426 } 427 //------------------------------------------------------------------------------ 428