1 //===------ omptarget.cpp - Target independent OpenMP target RTL -- C++ -*-===// 2 // 3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 // See https://llvm.org/LICENSE.txt for license information. 5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 // 7 //===----------------------------------------------------------------------===// 8 // 9 // Implementation of the interface to be used by Clang during the codegen of a 10 // target region. 11 // 12 //===----------------------------------------------------------------------===// 13 14 #include <omptarget.h> 15 16 #include "device.h" 17 #include "private.h" 18 #include "rtl.h" 19 20 #include <cassert> 21 #include <vector> 22 23 #ifdef OMPTARGET_DEBUG 24 int DebugLevel = 0; 25 #endif // OMPTARGET_DEBUG 26 27 28 29 /* All begin addresses for partially mapped structs must be 8-aligned in order 30 * to ensure proper alignment of members. E.g. 31 * 32 * struct S { 33 * int a; // 4-aligned 34 * int b; // 4-aligned 35 * int *p; // 8-aligned 36 * } s1; 37 * ... 38 * #pragma omp target map(tofrom: s1.b, s1.p[0:N]) 39 * { 40 * s1.b = 5; 41 * for (int i...) s1.p[i] = ...; 42 * } 43 * 44 * Here we are mapping s1 starting from member b, so BaseAddress=&s1=&s1.a and 45 * BeginAddress=&s1.b. Let's assume that the struct begins at address 0x100, 46 * then &s1.a=0x100, &s1.b=0x104, &s1.p=0x108. Each member obeys the alignment 47 * requirements for its type. Now, when we allocate memory on the device, in 48 * CUDA's case cuMemAlloc() returns an address which is at least 256-aligned. 49 * This means that the chunk of the struct on the device will start at a 50 * 256-aligned address, let's say 0x200. Then the address of b will be 0x200 and 51 * address of p will be a misaligned 0x204 (on the host there was no need to add 52 * padding between b and p, so p comes exactly 4 bytes after b). If the device 53 * kernel tries to access s1.p, a misaligned address error occurs (as reported 54 * by the CUDA plugin). By padding the begin address down to a multiple of 8 and 55 * extending the size of the allocated chuck accordingly, the chuck on the 56 * device will start at 0x200 with the padding (4 bytes), then &s1.b=0x204 and 57 * &s1.p=0x208, as they should be to satisfy the alignment requirements. 58 */ 59 static const int64_t alignment = 8; 60 61 /// Map global data and execute pending ctors 62 static int InitLibrary(DeviceTy& Device) { 63 /* 64 * Map global data 65 */ 66 int32_t device_id = Device.DeviceID; 67 int rc = OFFLOAD_SUCCESS; 68 69 Device.PendingGlobalsMtx.lock(); 70 TrlTblMtx->lock(); 71 for (HostEntriesBeginToTransTableTy::iterator 72 ii = HostEntriesBeginToTransTable->begin(); 73 ii != HostEntriesBeginToTransTable->end(); ++ii) { 74 TranslationTable *TransTable = &ii->second; 75 if (TransTable->HostTable.EntriesBegin == 76 TransTable->HostTable.EntriesEnd) { 77 // No host entry so no need to proceed 78 continue; 79 } 80 if (TransTable->TargetsTable[device_id] != 0) { 81 // Library entries have already been processed 82 continue; 83 } 84 85 // 1) get image. 86 assert(TransTable->TargetsImages.size() > (size_t)device_id && 87 "Not expecting a device ID outside the table's bounds!"); 88 __tgt_device_image *img = TransTable->TargetsImages[device_id]; 89 if (!img) { 90 DP("No image loaded for device id %d.\n", device_id); 91 rc = OFFLOAD_FAIL; 92 break; 93 } 94 // 2) load image into the target table. 95 __tgt_target_table *TargetTable = 96 TransTable->TargetsTable[device_id] = Device.load_binary(img); 97 // Unable to get table for this image: invalidate image and fail. 98 if (!TargetTable) { 99 DP("Unable to generate entries table for device id %d.\n", device_id); 100 TransTable->TargetsImages[device_id] = 0; 101 rc = OFFLOAD_FAIL; 102 break; 103 } 104 105 // Verify whether the two table sizes match. 106 size_t hsize = 107 TransTable->HostTable.EntriesEnd - TransTable->HostTable.EntriesBegin; 108 size_t tsize = TargetTable->EntriesEnd - TargetTable->EntriesBegin; 109 110 // Invalid image for these host entries! 111 if (hsize != tsize) { 112 DP("Host and Target tables mismatch for device id %d [%zx != %zx].\n", 113 device_id, hsize, tsize); 114 TransTable->TargetsImages[device_id] = 0; 115 TransTable->TargetsTable[device_id] = 0; 116 rc = OFFLOAD_FAIL; 117 break; 118 } 119 120 // process global data that needs to be mapped. 121 Device.DataMapMtx.lock(); 122 __tgt_target_table *HostTable = &TransTable->HostTable; 123 for (__tgt_offload_entry *CurrDeviceEntry = TargetTable->EntriesBegin, 124 *CurrHostEntry = HostTable->EntriesBegin, 125 *EntryDeviceEnd = TargetTable->EntriesEnd; 126 CurrDeviceEntry != EntryDeviceEnd; 127 CurrDeviceEntry++, CurrHostEntry++) { 128 if (CurrDeviceEntry->size != 0) { 129 // has data. 130 assert(CurrDeviceEntry->size == CurrHostEntry->size && 131 "data size mismatch"); 132 133 // Fortran may use multiple weak declarations for the same symbol, 134 // therefore we must allow for multiple weak symbols to be loaded from 135 // the fat binary. Treat these mappings as any other "regular" mapping. 136 // Add entry to map. 137 if (Device.getTgtPtrBegin(CurrHostEntry->addr, CurrHostEntry->size)) 138 continue; 139 DP("Add mapping from host " DPxMOD " to device " DPxMOD " with size %zu" 140 "\n", DPxPTR(CurrHostEntry->addr), DPxPTR(CurrDeviceEntry->addr), 141 CurrDeviceEntry->size); 142 Device.HostDataToTargetMap.push_front(HostDataToTargetTy( 143 (uintptr_t)CurrHostEntry->addr /*HstPtrBase*/, 144 (uintptr_t)CurrHostEntry->addr /*HstPtrBegin*/, 145 (uintptr_t)CurrHostEntry->addr + CurrHostEntry->size /*HstPtrEnd*/, 146 (uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/, 147 true /*IsRefCountINF*/)); 148 } 149 } 150 Device.DataMapMtx.unlock(); 151 } 152 TrlTblMtx->unlock(); 153 154 if (rc != OFFLOAD_SUCCESS) { 155 Device.PendingGlobalsMtx.unlock(); 156 return rc; 157 } 158 159 /* 160 * Run ctors for static objects 161 */ 162 if (!Device.PendingCtorsDtors.empty()) { 163 // Call all ctors for all libraries registered so far 164 for (auto &lib : Device.PendingCtorsDtors) { 165 if (!lib.second.PendingCtors.empty()) { 166 DP("Has pending ctors... call now\n"); 167 for (auto &entry : lib.second.PendingCtors) { 168 void *ctor = entry; 169 int rc = target(device_id, ctor, 0, NULL, NULL, NULL, 170 NULL, 1, 1, true /*team*/); 171 if (rc != OFFLOAD_SUCCESS) { 172 DP("Running ctor " DPxMOD " failed.\n", DPxPTR(ctor)); 173 Device.PendingGlobalsMtx.unlock(); 174 return OFFLOAD_FAIL; 175 } 176 } 177 // Clear the list to indicate that this device has been used 178 lib.second.PendingCtors.clear(); 179 DP("Done with pending ctors for lib " DPxMOD "\n", DPxPTR(lib.first)); 180 } 181 } 182 } 183 Device.HasPendingGlobals = false; 184 Device.PendingGlobalsMtx.unlock(); 185 186 return OFFLOAD_SUCCESS; 187 } 188 189 // Check whether a device has been initialized, global ctors have been 190 // executed and global data has been mapped; do so if not already done. 191 int CheckDeviceAndCtors(int64_t device_id) { 192 // Is device ready? 193 if (!device_is_ready(device_id)) { 194 DP("Device %" PRId64 " is not ready.\n", device_id); 195 return OFFLOAD_FAIL; 196 } 197 198 // Get device info. 199 DeviceTy &Device = Devices[device_id]; 200 201 // Check whether global data has been mapped for this device 202 Device.PendingGlobalsMtx.lock(); 203 bool hasPendingGlobals = Device.HasPendingGlobals; 204 Device.PendingGlobalsMtx.unlock(); 205 if (hasPendingGlobals && InitLibrary(Device) != OFFLOAD_SUCCESS) { 206 DP("Failed to init globals on device %" PRId64 "\n", device_id); 207 return OFFLOAD_FAIL; 208 } 209 210 return OFFLOAD_SUCCESS; 211 } 212 213 static int32_t member_of(int64_t type) { 214 return ((type & OMP_TGT_MAPTYPE_MEMBER_OF) >> 48) - 1; 215 } 216 217 /// Internal function to do the mapping and transfer the data to the device 218 int target_data_begin(DeviceTy &Device, int32_t arg_num, void **args_base, 219 void **args, int64_t *arg_sizes, int64_t *arg_types, 220 __tgt_async_info *async_info_ptr) { 221 // process each input. 222 for (int32_t i = 0; i < arg_num; ++i) { 223 // Ignore private variables and arrays - there is no mapping for them. 224 if ((arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) || 225 (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE)) 226 continue; 227 228 void *HstPtrBegin = args[i]; 229 void *HstPtrBase = args_base[i]; 230 int64_t data_size = arg_sizes[i]; 231 232 // Adjust for proper alignment if this is a combined entry (for structs). 233 // Look at the next argument - if that is MEMBER_OF this one, then this one 234 // is a combined entry. 235 int64_t padding = 0; 236 const int next_i = i+1; 237 if (member_of(arg_types[i]) < 0 && next_i < arg_num && 238 member_of(arg_types[next_i]) == i) { 239 padding = (int64_t)HstPtrBegin % alignment; 240 if (padding) { 241 DP("Using a padding of %" PRId64 " bytes for begin address " DPxMOD 242 "\n", padding, DPxPTR(HstPtrBegin)); 243 HstPtrBegin = (char *) HstPtrBegin - padding; 244 data_size += padding; 245 } 246 } 247 248 // Address of pointer on the host and device, respectively. 249 void *Pointer_HstPtrBegin, *Pointer_TgtPtrBegin; 250 bool IsNew, Pointer_IsNew; 251 bool IsHostPtr = false; 252 bool IsImplicit = arg_types[i] & OMP_TGT_MAPTYPE_IMPLICIT; 253 // Force the creation of a device side copy of the data when: 254 // a close map modifier was associated with a map that contained a to. 255 bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE; 256 // UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we 257 // have reached this point via __tgt_target_data_begin and not __tgt_target 258 // then no argument is marked as TARGET_PARAM ("omp target data map" is not 259 // associated with a target region, so there are no target parameters). This 260 // may be considered a hack, we could revise the scheme in the future. 261 bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF); 262 if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { 263 DP("Has a pointer entry: \n"); 264 // base is address of pointer. 265 Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBase, HstPtrBase, 266 sizeof(void *), Pointer_IsNew, IsHostPtr, IsImplicit, UpdateRef, 267 HasCloseModifier); 268 if (!Pointer_TgtPtrBegin) { 269 DP("Call to getOrAllocTgtPtr returned null pointer (device failure or " 270 "illegal mapping).\n"); 271 return OFFLOAD_FAIL; 272 } 273 DP("There are %zu bytes allocated at target address " DPxMOD " - is%s new" 274 "\n", sizeof(void *), DPxPTR(Pointer_TgtPtrBegin), 275 (Pointer_IsNew ? "" : " not")); 276 Pointer_HstPtrBegin = HstPtrBase; 277 // modify current entry. 278 HstPtrBase = *(void **)HstPtrBase; 279 UpdateRef = true; // subsequently update ref count of pointee 280 } 281 282 void *TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBegin, HstPtrBase, 283 data_size, IsNew, IsHostPtr, IsImplicit, UpdateRef, HasCloseModifier); 284 if (!TgtPtrBegin && data_size) { 285 // If data_size==0, then the argument could be a zero-length pointer to 286 // NULL, so getOrAlloc() returning NULL is not an error. 287 DP("Call to getOrAllocTgtPtr returned null pointer (device failure or " 288 "illegal mapping).\n"); 289 } 290 DP("There are %" PRId64 " bytes allocated at target address " DPxMOD 291 " - is%s new\n", data_size, DPxPTR(TgtPtrBegin), 292 (IsNew ? "" : " not")); 293 294 if (arg_types[i] & OMP_TGT_MAPTYPE_RETURN_PARAM) { 295 uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase; 296 void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta); 297 DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase)); 298 args_base[i] = TgtPtrBase; 299 } 300 301 if (arg_types[i] & OMP_TGT_MAPTYPE_TO) { 302 bool copy = false; 303 if (!(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) || 304 HasCloseModifier) { 305 if (IsNew || (arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS)) { 306 copy = true; 307 } else if (arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) { 308 // Copy data only if the "parent" struct has RefCount==1. 309 int32_t parent_idx = member_of(arg_types[i]); 310 uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]); 311 assert(parent_rc > 0 && "parent struct not found"); 312 if (parent_rc == 1) { 313 copy = true; 314 } 315 } 316 } 317 318 if (copy && !IsHostPtr) { 319 DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n", 320 data_size, DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin)); 321 int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, data_size, 322 async_info_ptr); 323 if (rt != OFFLOAD_SUCCESS) { 324 DP("Copying data to device failed.\n"); 325 return OFFLOAD_FAIL; 326 } 327 } 328 } 329 330 if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) { 331 DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n", 332 DPxPTR(Pointer_TgtPtrBegin), DPxPTR(TgtPtrBegin)); 333 uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase; 334 void *TgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta); 335 int rt = Device.data_submit(Pointer_TgtPtrBegin, &TgtPtrBase, 336 sizeof(void *), async_info_ptr); 337 if (rt != OFFLOAD_SUCCESS) { 338 DP("Copying data to device failed.\n"); 339 return OFFLOAD_FAIL; 340 } 341 // create shadow pointers for this entry 342 Device.ShadowMtx.lock(); 343 Device.ShadowPtrMap[Pointer_HstPtrBegin] = {HstPtrBase, 344 Pointer_TgtPtrBegin, TgtPtrBase}; 345 Device.ShadowMtx.unlock(); 346 } 347 } 348 349 return OFFLOAD_SUCCESS; 350 } 351 352 /// Internal function to undo the mapping and retrieve the data from the device. 353 int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base, 354 void **args, int64_t *arg_sizes, int64_t *arg_types, 355 __tgt_async_info *async_info_ptr) { 356 // process each input. 357 for (int32_t i = arg_num - 1; i >= 0; --i) { 358 // Ignore private variables and arrays - there is no mapping for them. 359 // Also, ignore the use_device_ptr directive, it has no effect here. 360 if ((arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) || 361 (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE)) 362 continue; 363 364 void *HstPtrBegin = args[i]; 365 int64_t data_size = arg_sizes[i]; 366 // Adjust for proper alignment if this is a combined entry (for structs). 367 // Look at the next argument - if that is MEMBER_OF this one, then this one 368 // is a combined entry. 369 int64_t padding = 0; 370 const int next_i = i+1; 371 if (member_of(arg_types[i]) < 0 && next_i < arg_num && 372 member_of(arg_types[next_i]) == i) { 373 padding = (int64_t)HstPtrBegin % alignment; 374 if (padding) { 375 DP("Using a padding of %" PRId64 " bytes for begin address " DPxMOD 376 "\n", padding, DPxPTR(HstPtrBegin)); 377 HstPtrBegin = (char *) HstPtrBegin - padding; 378 data_size += padding; 379 } 380 } 381 382 bool IsLast, IsHostPtr; 383 bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) || 384 (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ); 385 bool ForceDelete = arg_types[i] & OMP_TGT_MAPTYPE_DELETE; 386 bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE; 387 388 // If PTR_AND_OBJ, HstPtrBegin is address of pointee 389 void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, data_size, IsLast, 390 UpdateRef, IsHostPtr); 391 DP("There are %" PRId64 " bytes allocated at target address " DPxMOD 392 " - is%s last\n", data_size, DPxPTR(TgtPtrBegin), 393 (IsLast ? "" : " not")); 394 395 bool DelEntry = IsLast || ForceDelete; 396 397 if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && 398 !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) { 399 DelEntry = false; // protect parent struct from being deallocated 400 } 401 402 if ((arg_types[i] & OMP_TGT_MAPTYPE_FROM) || DelEntry) { 403 // Move data back to the host 404 if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) { 405 bool Always = arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS; 406 bool CopyMember = false; 407 if (!(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) || 408 HasCloseModifier) { 409 if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && 410 !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) { 411 // Copy data only if the "parent" struct has RefCount==1. 412 int32_t parent_idx = member_of(arg_types[i]); 413 uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]); 414 assert(parent_rc > 0 && "parent struct not found"); 415 if (parent_rc == 1) { 416 CopyMember = true; 417 } 418 } 419 } 420 421 if ((DelEntry || Always || CopyMember) && 422 !(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && 423 TgtPtrBegin == HstPtrBegin)) { 424 DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n", 425 data_size, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin)); 426 int rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, data_size, 427 async_info_ptr); 428 if (rt != OFFLOAD_SUCCESS) { 429 DP("Copying data from device failed.\n"); 430 return OFFLOAD_FAIL; 431 } 432 } 433 } 434 435 // If we copied back to the host a struct/array containing pointers, we 436 // need to restore the original host pointer values from their shadow 437 // copies. If the struct is going to be deallocated, remove any remaining 438 // shadow pointer entries for this struct. 439 uintptr_t lb = (uintptr_t) HstPtrBegin; 440 uintptr_t ub = (uintptr_t) HstPtrBegin + data_size; 441 Device.ShadowMtx.lock(); 442 for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin(); 443 it != Device.ShadowPtrMap.end();) { 444 void **ShadowHstPtrAddr = (void**) it->first; 445 446 // An STL map is sorted on its keys; use this property 447 // to quickly determine when to break out of the loop. 448 if ((uintptr_t) ShadowHstPtrAddr < lb) { 449 ++it; 450 continue; 451 } 452 if ((uintptr_t) ShadowHstPtrAddr >= ub) 453 break; 454 455 // If we copied the struct to the host, we need to restore the pointer. 456 if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) { 457 DP("Restoring original host pointer value " DPxMOD " for host " 458 "pointer " DPxMOD "\n", DPxPTR(it->second.HstPtrVal), 459 DPxPTR(ShadowHstPtrAddr)); 460 *ShadowHstPtrAddr = it->second.HstPtrVal; 461 } 462 // If the struct is to be deallocated, remove the shadow entry. 463 if (DelEntry) { 464 DP("Removing shadow pointer " DPxMOD "\n", DPxPTR(ShadowHstPtrAddr)); 465 it = Device.ShadowPtrMap.erase(it); 466 } else { 467 ++it; 468 } 469 } 470 Device.ShadowMtx.unlock(); 471 472 // Deallocate map 473 if (DelEntry) { 474 int rt = Device.deallocTgtPtr(HstPtrBegin, data_size, ForceDelete, 475 HasCloseModifier); 476 if (rt != OFFLOAD_SUCCESS) { 477 DP("Deallocating data from device failed.\n"); 478 return OFFLOAD_FAIL; 479 } 480 } 481 } 482 } 483 484 return OFFLOAD_SUCCESS; 485 } 486 487 /// Internal function to pass data to/from the target. 488 int target_data_update(DeviceTy &Device, int32_t arg_num, 489 void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types) { 490 // process each input. 491 for (int32_t i = 0; i < arg_num; ++i) { 492 if ((arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) || 493 (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE)) 494 continue; 495 496 void *HstPtrBegin = args[i]; 497 int64_t MapSize = arg_sizes[i]; 498 bool IsLast, IsHostPtr; 499 void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, MapSize, IsLast, 500 false, IsHostPtr); 501 if (!TgtPtrBegin) { 502 DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin)); 503 continue; 504 } 505 506 if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && 507 TgtPtrBegin == HstPtrBegin) { 508 DP("hst data:" DPxMOD " unified and shared, becomes a noop\n", 509 DPxPTR(HstPtrBegin)); 510 continue; 511 } 512 513 if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) { 514 DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n", 515 arg_sizes[i], DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin)); 516 int rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, MapSize, nullptr); 517 if (rt != OFFLOAD_SUCCESS) { 518 DP("Copying data from device failed.\n"); 519 return OFFLOAD_FAIL; 520 } 521 522 uintptr_t lb = (uintptr_t) HstPtrBegin; 523 uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize; 524 Device.ShadowMtx.lock(); 525 for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin(); 526 it != Device.ShadowPtrMap.end(); ++it) { 527 void **ShadowHstPtrAddr = (void**) it->first; 528 if ((uintptr_t) ShadowHstPtrAddr < lb) 529 continue; 530 if ((uintptr_t) ShadowHstPtrAddr >= ub) 531 break; 532 DP("Restoring original host pointer value " DPxMOD " for host pointer " 533 DPxMOD "\n", DPxPTR(it->second.HstPtrVal), 534 DPxPTR(ShadowHstPtrAddr)); 535 *ShadowHstPtrAddr = it->second.HstPtrVal; 536 } 537 Device.ShadowMtx.unlock(); 538 } 539 540 if (arg_types[i] & OMP_TGT_MAPTYPE_TO) { 541 DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n", 542 arg_sizes[i], DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin)); 543 int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, MapSize, nullptr); 544 if (rt != OFFLOAD_SUCCESS) { 545 DP("Copying data to device failed.\n"); 546 return OFFLOAD_FAIL; 547 } 548 549 uintptr_t lb = (uintptr_t) HstPtrBegin; 550 uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize; 551 Device.ShadowMtx.lock(); 552 for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin(); 553 it != Device.ShadowPtrMap.end(); ++it) { 554 void **ShadowHstPtrAddr = (void**) it->first; 555 if ((uintptr_t) ShadowHstPtrAddr < lb) 556 continue; 557 if ((uintptr_t) ShadowHstPtrAddr >= ub) 558 break; 559 DP("Restoring original target pointer value " DPxMOD " for target " 560 "pointer " DPxMOD "\n", DPxPTR(it->second.TgtPtrVal), 561 DPxPTR(it->second.TgtPtrAddr)); 562 rt = Device.data_submit(it->second.TgtPtrAddr, 563 &it->second.TgtPtrVal, sizeof(void *), nullptr); 564 if (rt != OFFLOAD_SUCCESS) { 565 DP("Copying data to device failed.\n"); 566 Device.ShadowMtx.unlock(); 567 return OFFLOAD_FAIL; 568 } 569 } 570 Device.ShadowMtx.unlock(); 571 } 572 } 573 return OFFLOAD_SUCCESS; 574 } 575 576 static const unsigned LambdaMapping = OMP_TGT_MAPTYPE_PTR_AND_OBJ | 577 OMP_TGT_MAPTYPE_LITERAL | 578 OMP_TGT_MAPTYPE_IMPLICIT; 579 static bool isLambdaMapping(int64_t Mapping) { 580 return (Mapping & LambdaMapping) == LambdaMapping; 581 } 582 583 /// performs the same actions as data_begin in case arg_num is 584 /// non-zero and initiates run of the offloaded region on the target platform; 585 /// if arg_num is non-zero after the region execution is done it also 586 /// performs the same action as data_update and data_end above. This function 587 /// returns 0 if it was able to transfer the execution to a target and an 588 /// integer different from zero otherwise. 589 int target(int64_t device_id, void *host_ptr, int32_t arg_num, 590 void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types, 591 int32_t team_num, int32_t thread_limit, int IsTeamConstruct) { 592 DeviceTy &Device = Devices[device_id]; 593 594 // Find the table information in the map or look it up in the translation 595 // tables. 596 TableMap *TM = 0; 597 TblMapMtx->lock(); 598 HostPtrToTableMapTy::iterator TableMapIt = HostPtrToTableMap->find(host_ptr); 599 if (TableMapIt == HostPtrToTableMap->end()) { 600 // We don't have a map. So search all the registered libraries. 601 TrlTblMtx->lock(); 602 for (HostEntriesBeginToTransTableTy::iterator 603 ii = HostEntriesBeginToTransTable->begin(), 604 ie = HostEntriesBeginToTransTable->end(); 605 !TM && ii != ie; ++ii) { 606 // get the translation table (which contains all the good info). 607 TranslationTable *TransTable = &ii->second; 608 // iterate over all the host table entries to see if we can locate the 609 // host_ptr. 610 __tgt_offload_entry *begin = TransTable->HostTable.EntriesBegin; 611 __tgt_offload_entry *end = TransTable->HostTable.EntriesEnd; 612 __tgt_offload_entry *cur = begin; 613 for (uint32_t i = 0; cur < end; ++cur, ++i) { 614 if (cur->addr != host_ptr) 615 continue; 616 // we got a match, now fill the HostPtrToTableMap so that we 617 // may avoid this search next time. 618 TM = &(*HostPtrToTableMap)[host_ptr]; 619 TM->Table = TransTable; 620 TM->Index = i; 621 break; 622 } 623 } 624 TrlTblMtx->unlock(); 625 } else { 626 TM = &TableMapIt->second; 627 } 628 TblMapMtx->unlock(); 629 630 // No map for this host pointer found! 631 if (!TM) { 632 DP("Host ptr " DPxMOD " does not have a matching target pointer.\n", 633 DPxPTR(host_ptr)); 634 return OFFLOAD_FAIL; 635 } 636 637 // get target table. 638 TrlTblMtx->lock(); 639 assert(TM->Table->TargetsTable.size() > (size_t)device_id && 640 "Not expecting a device ID outside the table's bounds!"); 641 __tgt_target_table *TargetTable = TM->Table->TargetsTable[device_id]; 642 TrlTblMtx->unlock(); 643 assert(TargetTable && "Global data has not been mapped\n"); 644 645 __tgt_async_info AsyncInfo; 646 647 // Move data to device. 648 int rc = target_data_begin(Device, arg_num, args_base, args, arg_sizes, 649 arg_types, &AsyncInfo); 650 if (rc != OFFLOAD_SUCCESS) { 651 DP("Call to target_data_begin failed, abort target.\n"); 652 return OFFLOAD_FAIL; 653 } 654 655 std::vector<void *> tgt_args; 656 std::vector<ptrdiff_t> tgt_offsets; 657 658 // List of (first-)private arrays allocated for this target region 659 std::vector<void *> fpArrays; 660 std::vector<int> tgtArgsPositions(arg_num, -1); 661 662 for (int32_t i = 0; i < arg_num; ++i) { 663 if (!(arg_types[i] & OMP_TGT_MAPTYPE_TARGET_PARAM)) { 664 // This is not a target parameter, do not push it into tgt_args. 665 // Check for lambda mapping. 666 if (isLambdaMapping(arg_types[i])) { 667 assert((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && 668 "PTR_AND_OBJ must be also MEMBER_OF."); 669 unsigned idx = member_of(arg_types[i]); 670 int tgtIdx = tgtArgsPositions[idx]; 671 assert(tgtIdx != -1 && "Base address must be translated already."); 672 // The parent lambda must be processed already and it must be the last 673 // in tgt_args and tgt_offsets arrays. 674 void *HstPtrVal = args[i]; 675 void *HstPtrBegin = args_base[i]; 676 void *HstPtrBase = args[idx]; 677 bool IsLast, IsHostPtr; // unused. 678 void *TgtPtrBase = 679 (void *)((intptr_t)tgt_args[tgtIdx] + tgt_offsets[tgtIdx]); 680 DP("Parent lambda base " DPxMOD "\n", DPxPTR(TgtPtrBase)); 681 uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase; 682 void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta); 683 void *Pointer_TgtPtrBegin = 684 Device.getTgtPtrBegin(HstPtrVal, arg_sizes[i], IsLast, false, 685 IsHostPtr); 686 if (!Pointer_TgtPtrBegin) { 687 DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n", 688 DPxPTR(HstPtrVal)); 689 continue; 690 } 691 if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && 692 TgtPtrBegin == HstPtrBegin) { 693 DP("Unified memory is active, no need to map lambda captured" 694 "variable (" DPxMOD ")\n", DPxPTR(HstPtrVal)); 695 continue; 696 } 697 DP("Update lambda reference (" DPxMOD ") -> [" DPxMOD "]\n", 698 DPxPTR(Pointer_TgtPtrBegin), DPxPTR(TgtPtrBegin)); 699 int rt = Device.data_submit(TgtPtrBegin, &Pointer_TgtPtrBegin, 700 sizeof(void *), &AsyncInfo); 701 if (rt != OFFLOAD_SUCCESS) { 702 DP("Copying data to device failed.\n"); 703 return OFFLOAD_FAIL; 704 } 705 } 706 continue; 707 } 708 void *HstPtrBegin = args[i]; 709 void *HstPtrBase = args_base[i]; 710 void *TgtPtrBegin; 711 ptrdiff_t TgtBaseOffset; 712 bool IsLast, IsHostPtr; // unused. 713 if (arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) { 714 DP("Forwarding first-private value " DPxMOD " to the target construct\n", 715 DPxPTR(HstPtrBase)); 716 TgtPtrBegin = HstPtrBase; 717 TgtBaseOffset = 0; 718 } else if (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE) { 719 // Allocate memory for (first-)private array 720 TgtPtrBegin = Device.RTL->data_alloc(Device.RTLDeviceID, 721 arg_sizes[i], HstPtrBegin); 722 if (!TgtPtrBegin) { 723 DP ("Data allocation for %sprivate array " DPxMOD " failed, " 724 "abort target.\n", 725 (arg_types[i] & OMP_TGT_MAPTYPE_TO ? "first-" : ""), 726 DPxPTR(HstPtrBegin)); 727 return OFFLOAD_FAIL; 728 } 729 fpArrays.push_back(TgtPtrBegin); 730 TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin; 731 #ifdef OMPTARGET_DEBUG 732 void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset); 733 DP("Allocated %" PRId64 " bytes of target memory at " DPxMOD " for " 734 "%sprivate array " DPxMOD " - pushing target argument " DPxMOD "\n", 735 arg_sizes[i], DPxPTR(TgtPtrBegin), 736 (arg_types[i] & OMP_TGT_MAPTYPE_TO ? "first-" : ""), 737 DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBase)); 738 #endif 739 // If first-private, copy data from host 740 if (arg_types[i] & OMP_TGT_MAPTYPE_TO) { 741 int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, arg_sizes[i], 742 &AsyncInfo); 743 if (rt != OFFLOAD_SUCCESS) { 744 DP("Copying data to device failed, failed.\n"); 745 return OFFLOAD_FAIL; 746 } 747 } 748 } else if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { 749 TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBase, sizeof(void *), IsLast, 750 false, IsHostPtr); 751 TgtBaseOffset = 0; // no offset for ptrs. 752 DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD " to " 753 "object " DPxMOD "\n", DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBase), 754 DPxPTR(HstPtrBase)); 755 } else { 756 TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast, 757 false, IsHostPtr); 758 TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin; 759 #ifdef OMPTARGET_DEBUG 760 void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset); 761 DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD "\n", 762 DPxPTR(TgtPtrBase), DPxPTR(HstPtrBegin)); 763 #endif 764 } 765 tgtArgsPositions[i] = tgt_args.size(); 766 tgt_args.push_back(TgtPtrBegin); 767 tgt_offsets.push_back(TgtBaseOffset); 768 } 769 770 assert(tgt_args.size() == tgt_offsets.size() && 771 "Size mismatch in arguments and offsets"); 772 773 // Pop loop trip count 774 uint64_t ltc = 0; 775 TblMapMtx->lock(); 776 auto I = Device.LoopTripCnt.find(__kmpc_global_thread_num(NULL)); 777 if (I != Device.LoopTripCnt.end()) { 778 ltc = I->second; 779 Device.LoopTripCnt.erase(I); 780 DP("loop trip count is %lu.\n", ltc); 781 } 782 TblMapMtx->unlock(); 783 784 // Launch device execution. 785 DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n", 786 TargetTable->EntriesBegin[TM->Index].name, 787 DPxPTR(TargetTable->EntriesBegin[TM->Index].addr), TM->Index); 788 if (IsTeamConstruct) { 789 rc = Device.run_team_region(TargetTable->EntriesBegin[TM->Index].addr, 790 &tgt_args[0], &tgt_offsets[0], tgt_args.size(), 791 team_num, thread_limit, ltc, &AsyncInfo); 792 } else { 793 rc = Device.run_region(TargetTable->EntriesBegin[TM->Index].addr, 794 &tgt_args[0], &tgt_offsets[0], tgt_args.size(), 795 &AsyncInfo); 796 } 797 if (rc != OFFLOAD_SUCCESS) { 798 DP ("Executing target region abort target.\n"); 799 return OFFLOAD_FAIL; 800 } 801 802 // Deallocate (first-)private arrays 803 for (auto it : fpArrays) { 804 int rt = Device.RTL->data_delete(Device.RTLDeviceID, it); 805 if (rt != OFFLOAD_SUCCESS) { 806 DP("Deallocation of (first-)private arrays failed.\n"); 807 return OFFLOAD_FAIL; 808 } 809 } 810 811 // Move data from device. 812 int rt = target_data_end(Device, arg_num, args_base, args, arg_sizes, 813 arg_types, &AsyncInfo); 814 if (rt != OFFLOAD_SUCCESS) { 815 DP("Call to target_data_end failed, abort targe.\n"); 816 return OFFLOAD_FAIL; 817 } 818 819 if (Device.RTL->synchronize) 820 return Device.RTL->synchronize(device_id, &AsyncInfo); 821 822 return OFFLOAD_SUCCESS; 823 } 824