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 return OFFLOAD_FAIL; 290 } 291 DP("There are %" PRId64 " bytes allocated at target address " DPxMOD 292 " - is%s new\n", data_size, DPxPTR(TgtPtrBegin), 293 (IsNew ? "" : " not")); 294 295 if (arg_types[i] & OMP_TGT_MAPTYPE_RETURN_PARAM) { 296 uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase; 297 void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta); 298 DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase)); 299 args_base[i] = TgtPtrBase; 300 } 301 302 if (arg_types[i] & OMP_TGT_MAPTYPE_TO) { 303 bool copy = false; 304 if (!(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) || 305 HasCloseModifier) { 306 if (IsNew || (arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS)) { 307 copy = true; 308 } else if (arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) { 309 // Copy data only if the "parent" struct has RefCount==1. 310 int32_t parent_idx = member_of(arg_types[i]); 311 uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]); 312 assert(parent_rc > 0 && "parent struct not found"); 313 if (parent_rc == 1) { 314 copy = true; 315 } 316 } 317 } 318 319 if (copy && !IsHostPtr) { 320 DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n", 321 data_size, DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin)); 322 int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, data_size, 323 async_info_ptr); 324 if (rt != OFFLOAD_SUCCESS) { 325 DP("Copying data to device failed.\n"); 326 return OFFLOAD_FAIL; 327 } 328 } 329 } 330 331 if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) { 332 DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n", 333 DPxPTR(Pointer_TgtPtrBegin), DPxPTR(TgtPtrBegin)); 334 uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase; 335 void *TgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta); 336 int rt = Device.data_submit(Pointer_TgtPtrBegin, &TgtPtrBase, 337 sizeof(void *), async_info_ptr); 338 if (rt != OFFLOAD_SUCCESS) { 339 DP("Copying data to device failed.\n"); 340 return OFFLOAD_FAIL; 341 } 342 // create shadow pointers for this entry 343 Device.ShadowMtx.lock(); 344 Device.ShadowPtrMap[Pointer_HstPtrBegin] = {HstPtrBase, 345 Pointer_TgtPtrBegin, TgtPtrBase}; 346 Device.ShadowMtx.unlock(); 347 } 348 } 349 350 return OFFLOAD_SUCCESS; 351 } 352 353 /// Internal function to undo the mapping and retrieve the data from the device. 354 int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base, 355 void **args, int64_t *arg_sizes, int64_t *arg_types, 356 __tgt_async_info *async_info_ptr) { 357 // process each input. 358 for (int32_t i = arg_num - 1; i >= 0; --i) { 359 // Ignore private variables and arrays - there is no mapping for them. 360 // Also, ignore the use_device_ptr directive, it has no effect here. 361 if ((arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) || 362 (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE)) 363 continue; 364 365 void *HstPtrBegin = args[i]; 366 int64_t data_size = arg_sizes[i]; 367 // Adjust for proper alignment if this is a combined entry (for structs). 368 // Look at the next argument - if that is MEMBER_OF this one, then this one 369 // is a combined entry. 370 int64_t padding = 0; 371 const int next_i = i+1; 372 if (member_of(arg_types[i]) < 0 && next_i < arg_num && 373 member_of(arg_types[next_i]) == i) { 374 padding = (int64_t)HstPtrBegin % alignment; 375 if (padding) { 376 DP("Using a padding of %" PRId64 " bytes for begin address " DPxMOD 377 "\n", padding, DPxPTR(HstPtrBegin)); 378 HstPtrBegin = (char *) HstPtrBegin - padding; 379 data_size += padding; 380 } 381 } 382 383 bool IsLast, IsHostPtr; 384 bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) || 385 (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ); 386 bool ForceDelete = arg_types[i] & OMP_TGT_MAPTYPE_DELETE; 387 bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE; 388 389 // If PTR_AND_OBJ, HstPtrBegin is address of pointee 390 void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, data_size, IsLast, 391 UpdateRef, IsHostPtr); 392 DP("There are %" PRId64 " bytes allocated at target address " DPxMOD 393 " - is%s last\n", data_size, DPxPTR(TgtPtrBegin), 394 (IsLast ? "" : " not")); 395 396 bool DelEntry = IsLast || ForceDelete; 397 398 if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && 399 !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) { 400 DelEntry = false; // protect parent struct from being deallocated 401 } 402 403 if ((arg_types[i] & OMP_TGT_MAPTYPE_FROM) || DelEntry) { 404 // Move data back to the host 405 if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) { 406 bool Always = arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS; 407 bool CopyMember = false; 408 if (!(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) || 409 HasCloseModifier) { 410 if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && 411 !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) { 412 // Copy data only if the "parent" struct has RefCount==1. 413 int32_t parent_idx = member_of(arg_types[i]); 414 uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]); 415 assert(parent_rc > 0 && "parent struct not found"); 416 if (parent_rc == 1) { 417 CopyMember = true; 418 } 419 } 420 } 421 422 if ((DelEntry || Always || CopyMember) && 423 !(RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && 424 TgtPtrBegin == HstPtrBegin)) { 425 DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n", 426 data_size, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin)); 427 int rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, data_size, 428 async_info_ptr); 429 if (rt != OFFLOAD_SUCCESS) { 430 DP("Copying data from device failed.\n"); 431 return OFFLOAD_FAIL; 432 } 433 } 434 } 435 436 // If we copied back to the host a struct/array containing pointers, we 437 // need to restore the original host pointer values from their shadow 438 // copies. If the struct is going to be deallocated, remove any remaining 439 // shadow pointer entries for this struct. 440 uintptr_t lb = (uintptr_t) HstPtrBegin; 441 uintptr_t ub = (uintptr_t) HstPtrBegin + data_size; 442 Device.ShadowMtx.lock(); 443 for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin(); 444 it != Device.ShadowPtrMap.end();) { 445 void **ShadowHstPtrAddr = (void**) it->first; 446 447 // An STL map is sorted on its keys; use this property 448 // to quickly determine when to break out of the loop. 449 if ((uintptr_t) ShadowHstPtrAddr < lb) { 450 ++it; 451 continue; 452 } 453 if ((uintptr_t) ShadowHstPtrAddr >= ub) 454 break; 455 456 // If we copied the struct to the host, we need to restore the pointer. 457 if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) { 458 DP("Restoring original host pointer value " DPxMOD " for host " 459 "pointer " DPxMOD "\n", DPxPTR(it->second.HstPtrVal), 460 DPxPTR(ShadowHstPtrAddr)); 461 *ShadowHstPtrAddr = it->second.HstPtrVal; 462 } 463 // If the struct is to be deallocated, remove the shadow entry. 464 if (DelEntry) { 465 DP("Removing shadow pointer " DPxMOD "\n", DPxPTR(ShadowHstPtrAddr)); 466 it = Device.ShadowPtrMap.erase(it); 467 } else { 468 ++it; 469 } 470 } 471 Device.ShadowMtx.unlock(); 472 473 // Deallocate map 474 if (DelEntry) { 475 int rt = Device.deallocTgtPtr(HstPtrBegin, data_size, ForceDelete, 476 HasCloseModifier); 477 if (rt != OFFLOAD_SUCCESS) { 478 DP("Deallocating data from device failed.\n"); 479 return OFFLOAD_FAIL; 480 } 481 } 482 } 483 } 484 485 return OFFLOAD_SUCCESS; 486 } 487 488 /// Internal function to pass data to/from the target. 489 int target_data_update(DeviceTy &Device, int32_t arg_num, 490 void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types) { 491 // process each input. 492 for (int32_t i = 0; i < arg_num; ++i) { 493 if ((arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) || 494 (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE)) 495 continue; 496 497 void *HstPtrBegin = args[i]; 498 int64_t MapSize = arg_sizes[i]; 499 bool IsLast, IsHostPtr; 500 void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, MapSize, IsLast, 501 false, IsHostPtr); 502 if (!TgtPtrBegin) { 503 DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin)); 504 continue; 505 } 506 507 if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && 508 TgtPtrBegin == HstPtrBegin) { 509 DP("hst data:" DPxMOD " unified and shared, becomes a noop\n", 510 DPxPTR(HstPtrBegin)); 511 continue; 512 } 513 514 if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) { 515 DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n", 516 arg_sizes[i], DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin)); 517 int rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, MapSize, nullptr); 518 if (rt != OFFLOAD_SUCCESS) { 519 DP("Copying data from device failed.\n"); 520 return OFFLOAD_FAIL; 521 } 522 523 uintptr_t lb = (uintptr_t) HstPtrBegin; 524 uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize; 525 Device.ShadowMtx.lock(); 526 for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin(); 527 it != Device.ShadowPtrMap.end(); ++it) { 528 void **ShadowHstPtrAddr = (void**) it->first; 529 if ((uintptr_t) ShadowHstPtrAddr < lb) 530 continue; 531 if ((uintptr_t) ShadowHstPtrAddr >= ub) 532 break; 533 DP("Restoring original host pointer value " DPxMOD " for host pointer " 534 DPxMOD "\n", DPxPTR(it->second.HstPtrVal), 535 DPxPTR(ShadowHstPtrAddr)); 536 *ShadowHstPtrAddr = it->second.HstPtrVal; 537 } 538 Device.ShadowMtx.unlock(); 539 } 540 541 if (arg_types[i] & OMP_TGT_MAPTYPE_TO) { 542 DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n", 543 arg_sizes[i], DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin)); 544 int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, MapSize, nullptr); 545 if (rt != OFFLOAD_SUCCESS) { 546 DP("Copying data to device failed.\n"); 547 return OFFLOAD_FAIL; 548 } 549 550 uintptr_t lb = (uintptr_t) HstPtrBegin; 551 uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize; 552 Device.ShadowMtx.lock(); 553 for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin(); 554 it != Device.ShadowPtrMap.end(); ++it) { 555 void **ShadowHstPtrAddr = (void**) it->first; 556 if ((uintptr_t) ShadowHstPtrAddr < lb) 557 continue; 558 if ((uintptr_t) ShadowHstPtrAddr >= ub) 559 break; 560 DP("Restoring original target pointer value " DPxMOD " for target " 561 "pointer " DPxMOD "\n", DPxPTR(it->second.TgtPtrVal), 562 DPxPTR(it->second.TgtPtrAddr)); 563 rt = Device.data_submit(it->second.TgtPtrAddr, 564 &it->second.TgtPtrVal, sizeof(void *), nullptr); 565 if (rt != OFFLOAD_SUCCESS) { 566 DP("Copying data to device failed.\n"); 567 Device.ShadowMtx.unlock(); 568 return OFFLOAD_FAIL; 569 } 570 } 571 Device.ShadowMtx.unlock(); 572 } 573 } 574 return OFFLOAD_SUCCESS; 575 } 576 577 static const unsigned LambdaMapping = OMP_TGT_MAPTYPE_PTR_AND_OBJ | 578 OMP_TGT_MAPTYPE_LITERAL | 579 OMP_TGT_MAPTYPE_IMPLICIT; 580 static bool isLambdaMapping(int64_t Mapping) { 581 return (Mapping & LambdaMapping) == LambdaMapping; 582 } 583 584 /// performs the same actions as data_begin in case arg_num is 585 /// non-zero and initiates run of the offloaded region on the target platform; 586 /// if arg_num is non-zero after the region execution is done it also 587 /// performs the same action as data_update and data_end above. This function 588 /// returns 0 if it was able to transfer the execution to a target and an 589 /// integer different from zero otherwise. 590 int target(int64_t device_id, void *host_ptr, int32_t arg_num, 591 void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types, 592 int32_t team_num, int32_t thread_limit, int IsTeamConstruct) { 593 DeviceTy &Device = Devices[device_id]; 594 595 // Find the table information in the map or look it up in the translation 596 // tables. 597 TableMap *TM = 0; 598 TblMapMtx->lock(); 599 HostPtrToTableMapTy::iterator TableMapIt = HostPtrToTableMap->find(host_ptr); 600 if (TableMapIt == HostPtrToTableMap->end()) { 601 // We don't have a map. So search all the registered libraries. 602 TrlTblMtx->lock(); 603 for (HostEntriesBeginToTransTableTy::iterator 604 ii = HostEntriesBeginToTransTable->begin(), 605 ie = HostEntriesBeginToTransTable->end(); 606 !TM && ii != ie; ++ii) { 607 // get the translation table (which contains all the good info). 608 TranslationTable *TransTable = &ii->second; 609 // iterate over all the host table entries to see if we can locate the 610 // host_ptr. 611 __tgt_offload_entry *begin = TransTable->HostTable.EntriesBegin; 612 __tgt_offload_entry *end = TransTable->HostTable.EntriesEnd; 613 __tgt_offload_entry *cur = begin; 614 for (uint32_t i = 0; cur < end; ++cur, ++i) { 615 if (cur->addr != host_ptr) 616 continue; 617 // we got a match, now fill the HostPtrToTableMap so that we 618 // may avoid this search next time. 619 TM = &(*HostPtrToTableMap)[host_ptr]; 620 TM->Table = TransTable; 621 TM->Index = i; 622 break; 623 } 624 } 625 TrlTblMtx->unlock(); 626 } else { 627 TM = &TableMapIt->second; 628 } 629 TblMapMtx->unlock(); 630 631 // No map for this host pointer found! 632 if (!TM) { 633 DP("Host ptr " DPxMOD " does not have a matching target pointer.\n", 634 DPxPTR(host_ptr)); 635 return OFFLOAD_FAIL; 636 } 637 638 // get target table. 639 TrlTblMtx->lock(); 640 assert(TM->Table->TargetsTable.size() > (size_t)device_id && 641 "Not expecting a device ID outside the table's bounds!"); 642 __tgt_target_table *TargetTable = TM->Table->TargetsTable[device_id]; 643 TrlTblMtx->unlock(); 644 assert(TargetTable && "Global data has not been mapped\n"); 645 646 __tgt_async_info AsyncInfo; 647 648 // Move data to device. 649 int rc = target_data_begin(Device, arg_num, args_base, args, arg_sizes, 650 arg_types, &AsyncInfo); 651 if (rc != OFFLOAD_SUCCESS) { 652 DP("Call to target_data_begin failed, abort target.\n"); 653 return OFFLOAD_FAIL; 654 } 655 656 std::vector<void *> tgt_args; 657 std::vector<ptrdiff_t> tgt_offsets; 658 659 // List of (first-)private arrays allocated for this target region 660 std::vector<void *> fpArrays; 661 std::vector<int> tgtArgsPositions(arg_num, -1); 662 663 for (int32_t i = 0; i < arg_num; ++i) { 664 if (!(arg_types[i] & OMP_TGT_MAPTYPE_TARGET_PARAM)) { 665 // This is not a target parameter, do not push it into tgt_args. 666 // Check for lambda mapping. 667 if (isLambdaMapping(arg_types[i])) { 668 assert((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && 669 "PTR_AND_OBJ must be also MEMBER_OF."); 670 unsigned idx = member_of(arg_types[i]); 671 int tgtIdx = tgtArgsPositions[idx]; 672 assert(tgtIdx != -1 && "Base address must be translated already."); 673 // The parent lambda must be processed already and it must be the last 674 // in tgt_args and tgt_offsets arrays. 675 void *HstPtrVal = args[i]; 676 void *HstPtrBegin = args_base[i]; 677 void *HstPtrBase = args[idx]; 678 bool IsLast, IsHostPtr; // unused. 679 void *TgtPtrBase = 680 (void *)((intptr_t)tgt_args[tgtIdx] + tgt_offsets[tgtIdx]); 681 DP("Parent lambda base " DPxMOD "\n", DPxPTR(TgtPtrBase)); 682 uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase; 683 void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta); 684 void *Pointer_TgtPtrBegin = 685 Device.getTgtPtrBegin(HstPtrVal, arg_sizes[i], IsLast, false, 686 IsHostPtr); 687 if (!Pointer_TgtPtrBegin) { 688 DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n", 689 DPxPTR(HstPtrVal)); 690 continue; 691 } 692 if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && 693 TgtPtrBegin == HstPtrBegin) { 694 DP("Unified memory is active, no need to map lambda captured" 695 "variable (" DPxMOD ")\n", DPxPTR(HstPtrVal)); 696 continue; 697 } 698 DP("Update lambda reference (" DPxMOD ") -> [" DPxMOD "]\n", 699 DPxPTR(Pointer_TgtPtrBegin), DPxPTR(TgtPtrBegin)); 700 int rt = Device.data_submit(TgtPtrBegin, &Pointer_TgtPtrBegin, 701 sizeof(void *), &AsyncInfo); 702 if (rt != OFFLOAD_SUCCESS) { 703 DP("Copying data to device failed.\n"); 704 return OFFLOAD_FAIL; 705 } 706 } 707 continue; 708 } 709 void *HstPtrBegin = args[i]; 710 void *HstPtrBase = args_base[i]; 711 void *TgtPtrBegin; 712 ptrdiff_t TgtBaseOffset; 713 bool IsLast, IsHostPtr; // unused. 714 if (arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) { 715 DP("Forwarding first-private value " DPxMOD " to the target construct\n", 716 DPxPTR(HstPtrBase)); 717 TgtPtrBegin = HstPtrBase; 718 TgtBaseOffset = 0; 719 } else if (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE) { 720 // Allocate memory for (first-)private array 721 TgtPtrBegin = Device.RTL->data_alloc(Device.RTLDeviceID, 722 arg_sizes[i], HstPtrBegin); 723 if (!TgtPtrBegin) { 724 DP ("Data allocation for %sprivate array " DPxMOD " failed, " 725 "abort target.\n", 726 (arg_types[i] & OMP_TGT_MAPTYPE_TO ? "first-" : ""), 727 DPxPTR(HstPtrBegin)); 728 return OFFLOAD_FAIL; 729 } 730 fpArrays.push_back(TgtPtrBegin); 731 TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin; 732 #ifdef OMPTARGET_DEBUG 733 void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset); 734 DP("Allocated %" PRId64 " bytes of target memory at " DPxMOD " for " 735 "%sprivate array " DPxMOD " - pushing target argument " DPxMOD "\n", 736 arg_sizes[i], DPxPTR(TgtPtrBegin), 737 (arg_types[i] & OMP_TGT_MAPTYPE_TO ? "first-" : ""), 738 DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBase)); 739 #endif 740 // If first-private, copy data from host 741 if (arg_types[i] & OMP_TGT_MAPTYPE_TO) { 742 int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, arg_sizes[i], 743 &AsyncInfo); 744 if (rt != OFFLOAD_SUCCESS) { 745 DP("Copying data to device failed, failed.\n"); 746 return OFFLOAD_FAIL; 747 } 748 } 749 } else if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { 750 TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBase, sizeof(void *), IsLast, 751 false, IsHostPtr); 752 TgtBaseOffset = 0; // no offset for ptrs. 753 DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD " to " 754 "object " DPxMOD "\n", DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBase), 755 DPxPTR(HstPtrBase)); 756 } else { 757 TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast, 758 false, IsHostPtr); 759 TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin; 760 #ifdef OMPTARGET_DEBUG 761 void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset); 762 DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD "\n", 763 DPxPTR(TgtPtrBase), DPxPTR(HstPtrBegin)); 764 #endif 765 } 766 tgtArgsPositions[i] = tgt_args.size(); 767 tgt_args.push_back(TgtPtrBegin); 768 tgt_offsets.push_back(TgtBaseOffset); 769 } 770 771 assert(tgt_args.size() == tgt_offsets.size() && 772 "Size mismatch in arguments and offsets"); 773 774 // Pop loop trip count 775 uint64_t ltc = 0; 776 TblMapMtx->lock(); 777 auto I = Device.LoopTripCnt.find(__kmpc_global_thread_num(NULL)); 778 if (I != Device.LoopTripCnt.end()) { 779 ltc = I->second; 780 Device.LoopTripCnt.erase(I); 781 DP("loop trip count is %lu.\n", ltc); 782 } 783 TblMapMtx->unlock(); 784 785 // Launch device execution. 786 DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n", 787 TargetTable->EntriesBegin[TM->Index].name, 788 DPxPTR(TargetTable->EntriesBegin[TM->Index].addr), TM->Index); 789 if (IsTeamConstruct) { 790 rc = Device.run_team_region(TargetTable->EntriesBegin[TM->Index].addr, 791 &tgt_args[0], &tgt_offsets[0], tgt_args.size(), 792 team_num, thread_limit, ltc, &AsyncInfo); 793 } else { 794 rc = Device.run_region(TargetTable->EntriesBegin[TM->Index].addr, 795 &tgt_args[0], &tgt_offsets[0], tgt_args.size(), 796 &AsyncInfo); 797 } 798 if (rc != OFFLOAD_SUCCESS) { 799 DP ("Executing target region abort target.\n"); 800 return OFFLOAD_FAIL; 801 } 802 803 // Deallocate (first-)private arrays 804 for (auto it : fpArrays) { 805 int rt = Device.RTL->data_delete(Device.RTLDeviceID, it); 806 if (rt != OFFLOAD_SUCCESS) { 807 DP("Deallocation of (first-)private arrays failed.\n"); 808 return OFFLOAD_FAIL; 809 } 810 } 811 812 // Move data from device. 813 int rt = target_data_end(Device, arg_num, args_base, args, arg_sizes, 814 arg_types, &AsyncInfo); 815 if (rt != OFFLOAD_SUCCESS) { 816 DP("Call to target_data_end failed, abort targe.\n"); 817 return OFFLOAD_FAIL; 818 } 819 820 if (Device.RTL->synchronize) 821 return Device.RTL->synchronize(device_id, &AsyncInfo); 822 823 return OFFLOAD_SUCCESS; 824 } 825