diff options
Diffstat (limited to 'offload')
53 files changed, 1707 insertions, 813 deletions
diff --git a/offload/DeviceRTL/CMakeLists.txt b/offload/DeviceRTL/CMakeLists.txt deleted file mode 100644 index e4916f4d4975..000000000000 --- a/offload/DeviceRTL/CMakeLists.txt +++ /dev/null @@ -1,188 +0,0 @@ -set(LIBOMPTARGET_BUILD_DEVICERTL_BCLIB TRUE CACHE BOOL - "Can be set to false to disable building this library.") - -if (NOT LIBOMPTARGET_BUILD_DEVICERTL_BCLIB) - message(STATUS "Not building DeviceRTL: Disabled by LIBOMPTARGET_BUILD_DEVICERTL_BCLIB") - return() -endif() - -# Check to ensure the host system is a supported host architecture. -if(NOT ${CMAKE_SIZEOF_VOID_P} EQUAL "8") - message(STATUS "Not building DeviceRTL: Runtime does not support 32-bit hosts") - return() -endif() - -if (LLVM_DIR) - # Builds that use pre-installed LLVM have LLVM_DIR set. - # A standalone or LLVM_ENABLE_RUNTIMES=openmp build takes this route - find_program(CLANG_TOOL clang PATHS ${LLVM_TOOLS_BINARY_DIR} NO_DEFAULT_PATH) -elseif (LLVM_TOOL_CLANG_BUILD AND NOT CMAKE_CROSSCOMPILING AND NOT OPENMP_STANDALONE_BUILD) - # LLVM in-tree builds may use CMake target names to discover the tools. - # A LLVM_ENABLE_PROJECTS=openmp build takes this route - set(CLANG_TOOL $<TARGET_FILE:clang>) -else() - message(STATUS "Not building DeviceRTL. No appropriate clang found") - return() -endif() - -set(devicertl_base_directory ${CMAKE_CURRENT_SOURCE_DIR}) -set(include_directory ${devicertl_base_directory}/include) -set(source_directory ${devicertl_base_directory}/src) - -set(include_files - ${include_directory}/Allocator.h - ${include_directory}/Configuration.h - ${include_directory}/Debug.h - ${include_directory}/Interface.h - ${include_directory}/LibC.h - ${include_directory}/Mapping.h - ${include_directory}/Profiling.h - ${include_directory}/State.h - ${include_directory}/Synchronization.h - ${include_directory}/DeviceTypes.h - ${include_directory}/DeviceUtils.h - ${include_directory}/Workshare.h -) - -set(src_files - ${source_directory}/Allocator.cpp - ${source_directory}/Configuration.cpp - ${source_directory}/Debug.cpp - ${source_directory}/Kernel.cpp - ${source_directory}/LibC.cpp - ${source_directory}/Mapping.cpp - ${source_directory}/Misc.cpp - ${source_directory}/Parallelism.cpp - ${source_directory}/Profiling.cpp - ${source_directory}/Reduction.cpp - ${source_directory}/State.cpp - ${source_directory}/Synchronization.cpp - ${source_directory}/Tasking.cpp - ${source_directory}/DeviceUtils.cpp - ${source_directory}/Workshare.cpp -) - -# We disable the slp vectorizer during the runtime optimization to avoid -# vectorized accesses to the shared state. Generally, those are "good" but -# the optimizer pipeline (esp. Attributor) does not fully support vectorized -# instructions yet and we end up missing out on way more important constant -# propagation. That said, we will run the vectorizer again after the runtime -# has been linked into the user program. -set(clang_opt_flags -O3 -mllvm -openmp-opt-disable -DSHARED_SCRATCHPAD_SIZE=512 -mllvm -vectorize-slp=false ) - -# If the user built with the GPU C library enabled we will use that instead. -if(${LIBOMPTARGET_GPU_LIBC_SUPPORT}) - list(APPEND clang_opt_flags -DOMPTARGET_HAS_LIBC) -endif() - -# Set flags for LLVM Bitcode compilation. -set(bc_flags -c -flto -std=c++17 -fvisibility=hidden - ${clang_opt_flags} -nogpulib -nostdlibinc - -fno-rtti -fno-exceptions -fconvergent-functions - -Wno-unknown-cuda-version - -DOMPTARGET_DEVICE_RUNTIME - -I${include_directory} - -I${devicertl_base_directory}/../include - -I${devicertl_base_directory}/../../libc -) - -# first create an object target -function(compileDeviceRTLLibrary target_name target_triple) - set(target_bc_flags ${ARGN}) - - foreach(src ${src_files}) - get_filename_component(infile ${src} ABSOLUTE) - get_filename_component(outfile ${src} NAME) - set(outfile "${outfile}-${target_name}.o") - set(depfile "${outfile}.d") - - # Passing an empty CPU to -march= suppressed target specific metadata. - add_custom_command(OUTPUT ${outfile} - COMMAND ${CLANG_TOOL} - ${bc_flags} - --target=${target_triple} - ${target_bc_flags} - -MD -MF ${depfile} - ${infile} -o ${outfile} - DEPENDS ${infile} - DEPFILE ${depfile} - COMMENT "Building LLVM bitcode ${outfile}" - VERBATIM - ) - if(TARGET clang) - # Add a file-level dependency to ensure that clang is up-to-date. - # By default, add_custom_command only builds clang if the - # executable is missing. - add_custom_command(OUTPUT ${outfile} - DEPENDS clang - APPEND - ) - endif() - set_property(DIRECTORY APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${outfile}) - - list(APPEND obj_files ${CMAKE_CURRENT_BINARY_DIR}/${outfile}) - endforeach() - # Trick to combine these into a bitcode file via the linker's LTO pass. This - # is used to provide the legacy `libomptarget-<name>.bc` files. Hack this - # through as an executable to get it to use the relocatable link. - add_executable(libomptarget-${target_name} ${obj_files}) - set_target_properties(libomptarget-${target_name} PROPERTIES - RUNTIME_OUTPUT_DIRECTORY ${LIBOMPTARGET_LLVM_LIBRARY_INTDIR} - LINKER_LANGUAGE CXX - BUILD_RPATH "" - INSTALL_RPATH "" - RUNTIME_OUTPUT_NAME libomptarget-${target_name}.bc) - target_compile_options(libomptarget-${target_name} PRIVATE - "--target=${target_triple}" "-fuse-ld=lld" "-march=" "-mcpu=" - "-Wno-unused-command-line-argument") - target_link_options(libomptarget-${target_name} PRIVATE "--target=${target_triple}" - "-r" "-nostdlib" "-flto" "-Wl,--lto-emit-llvm" - "-fuse-ld=lld" "-march=" "-mcpu=") - install(TARGETS libomptarget-${target_name} - PERMISSIONS OWNER_WRITE OWNER_READ GROUP_READ WORLD_READ - DESTINATION "lib${LLVM_LIBDIR_SUFFIX}/${target_triple}") - - add_library(omptarget.${target_name}.all_objs OBJECT IMPORTED) - set_property(TARGET omptarget.${target_name}.all_objs APPEND PROPERTY IMPORTED_OBJECTS - ${LIBOMPTARGET_LLVM_LIBRARY_INTDIR}/libomptarget-${target_name}.bc) - add_dependencies(omptarget.${target_name}.all_objs libomptarget-${target_name}) - - # Archive all the object files generated above into a static library - add_library(omptarget.${target_name} STATIC) - set_target_properties(omptarget.${target_name} PROPERTIES - ARCHIVE_OUTPUT_DIRECTORY "${LIBOMPTARGET_LLVM_LIBRARY_INTDIR}/${target_triple}" - ARCHIVE_OUTPUT_NAME ompdevice - LINKER_LANGUAGE CXX - ) - target_link_libraries(omptarget.${target_name} PRIVATE omptarget.${target_name}.all_objs) - target_link_options(omptarget.${target_name} PRIVATE "--target=${target_triple}" - "-Wno-unused-command-line-argument" "-r" "-nostdlib" "-flto" - "-Wl,--lto-emit-llvm" "-fuse-ld=lld" "-march=" "-mcpu=") - - install(TARGETS omptarget.${target_name} - ARCHIVE DESTINATION "lib${LLVM_LIBDIR_SUFFIX}/${target_triple}") - - if (CMAKE_EXPORT_COMPILE_COMMANDS) - set(ide_target_name omptarget-ide-${target_name}) - add_library(${ide_target_name} STATIC EXCLUDE_FROM_ALL ${src_files}) - target_compile_options(${ide_target_name} PRIVATE - -fvisibility=hidden --target=${target_triple} - -nogpulib -nostdlibinc -Wno-unknown-cuda-version - ) - target_compile_definitions(${ide_target_name} PRIVATE SHARED_SCRATCHPAD_SIZE=512) - target_include_directories(${ide_target_name} PRIVATE - ${include_directory} - ${devicertl_base_directory}/../../libc - ${devicertl_base_directory}/../include - ) - install(TARGETS ${ide_target_name} EXCLUDE_FROM_ALL) - endif() -endfunction() - -if(NOT LLVM_TARGETS_TO_BUILD OR "AMDGPU" IN_LIST LLVM_TARGETS_TO_BUILD) - compileDeviceRTLLibrary(amdgpu amdgcn-amd-amdhsa -Xclang -mcode-object-version=none) -endif() - -if(NOT LLVM_TARGETS_TO_BUILD OR "NVPTX" IN_LIST LLVM_TARGETS_TO_BUILD) - compileDeviceRTLLibrary(nvptx nvptx64-nvidia-cuda --cuda-feature=+ptx63) -endif() diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h index 93c1e56905ae..45bd9c6e7da8 100644 --- a/offload/include/OpenMP/Mapping.h +++ b/offload/include/OpenMP/Mapping.h @@ -49,9 +49,46 @@ public: /// Information about shadow pointers. struct ShadowPtrInfoTy { void **HstPtrAddr = nullptr; - void *HstPtrVal = nullptr; void **TgtPtrAddr = nullptr; - void *TgtPtrVal = nullptr; + int64_t PtrSize = sizeof(void *); // Size of the pointer/descriptor + + // Store the complete contents for both host and target pointers/descriptors. + // 96 bytes is chosen as the "Small" size to cover simple Fortran + // descriptors of up to 3 dimensions. + llvm::SmallVector<char, 96> HstPtrContent; + llvm::SmallVector<char, 96> TgtPtrContent; + + ShadowPtrInfoTy(void **HstPtrAddr, void **TgtPtrAddr, void *TgtPteeBase, + int64_t PtrSize) + : HstPtrAddr(HstPtrAddr), TgtPtrAddr(TgtPtrAddr), PtrSize(PtrSize), + HstPtrContent(PtrSize), TgtPtrContent(PtrSize) { + constexpr int64_t VoidPtrSize = sizeof(void *); + assert(HstPtrAddr != nullptr && "HstPtrAddr is nullptr"); + assert(TgtPtrAddr != nullptr && "TgtPtrAddr is nullptr"); + assert(PtrSize >= VoidPtrSize && "PtrSize is less than sizeof(void *)"); + + void *HstPteeBase = *HstPtrAddr; + // The first VoidPtrSize bytes for HstPtrContent/TgtPtrContent are from + // HstPteeBase/TgtPteeBase. + std::memcpy(HstPtrContent.data(), &HstPteeBase, VoidPtrSize); + std::memcpy(TgtPtrContent.data(), &TgtPteeBase, VoidPtrSize); + + // If we are not dealing with Fortran descriptors (pointers larger than + // VoidPtrSize), then that's that. + if (PtrSize <= VoidPtrSize) + return; + + // For larger pointers, i.e. Fortran descriptors, the remaining contents of + // the descriptor come from the host descriptor, i.e. HstPtrAddr. + std::memcpy(HstPtrContent.data() + VoidPtrSize, + reinterpret_cast<char *>(HstPtrAddr) + VoidPtrSize, + PtrSize - VoidPtrSize); + std::memcpy(TgtPtrContent.data() + VoidPtrSize, + reinterpret_cast<char *>(HstPtrAddr) + VoidPtrSize, + PtrSize - VoidPtrSize); + } + + ShadowPtrInfoTy() = delete; bool operator==(const ShadowPtrInfoTy &Other) const { return HstPtrAddr == Other.HstPtrAddr; @@ -243,9 +280,25 @@ public: auto Pair = States->ShadowPtrInfos.insert(ShadowPtrInfo); if (Pair.second) return true; + // Check for a stale entry, if found, replace the old one. - if ((*Pair.first).TgtPtrVal == ShadowPtrInfo.TgtPtrVal) + + // For Fortran descriptors, we need to compare their full contents, + // as the starting address may be the same while other fields have + // been updated. e.g. + // + // !$omp target enter data map(x(1:100)) ! (1) + // p => x(10: 19) + // !$omp target enter data map(p, p(:)) ! (2) + // p => x(5: 9) + // !$omp target enter data map(attach(always): p(:)) ! (3) + // + // While &desc_p and &p(1) (TgtPtrAddr and first "sizeof(void*)" bytes of + // TgtPtrContent) are same for (2) and (3), the pointer attachment for (3) + // needs to update the bounds information in the descriptor of p on device. + if ((*Pair.first).TgtPtrContent == ShadowPtrInfo.TgtPtrContent) return false; + States->ShadowPtrInfos.erase(ShadowPtrInfo); return addShadowPointer(ShadowPtrInfo); } diff --git a/offload/include/omptarget.h b/offload/include/omptarget.h index 8fd722bb1502..794b79e07674 100644 --- a/offload/include/omptarget.h +++ b/offload/include/omptarget.h @@ -101,8 +101,6 @@ enum TargetAllocTy : int32_t { TARGET_ALLOC_HOST, TARGET_ALLOC_SHARED, TARGET_ALLOC_DEFAULT, - /// The allocation will not block on other streams. - TARGET_ALLOC_DEVICE_NON_BLOCKING, }; inline KernelArgsTy CTorDTorKernelArgs = { diff --git a/offload/liboffload/API/Memory.td b/offload/liboffload/API/Memory.td index cc98b672a26a..79e803833004 100644 --- a/offload/liboffload/API/Memory.td +++ b/offload/liboffload/API/Memory.td @@ -21,6 +21,9 @@ def ol_alloc_type_t : Enum { def olMemAlloc : Function { let desc = "Creates a memory allocation on the specified device."; + let details = [ + "All allocations through olMemAlloc regardless of source share a single virtual address range. There is no risk of multiple devices returning equal pointers to different memory." + ]; let params = [ Param<"ol_device_handle_t", "Device", "handle of the device to allocate on", PARAM_IN>, Param<"ol_alloc_type_t", "Type", "type of the allocation", PARAM_IN>, @@ -42,6 +45,56 @@ def olMemFree : Function { let returns = []; } +def ol_mem_info_t : Enum { + let desc = "Supported memory info."; + let is_typed = 1; + let etors = [ + TaggedEtor<"DEVICE", "ol_device_handle_t", "The handle of the device associated with the allocation.">, + TaggedEtor<"BASE", "void *", "Base address of this allocation.">, + TaggedEtor<"SIZE", "size_t", "Size of this allocation in bytes.">, + TaggedEtor<"TYPE", "ol_alloc_type_t", "Type of this allocation.">, + ]; +} + +def olGetMemInfo : Function { + let desc = "Queries the given property of a memory allocation allocated with olMemAlloc."; + let details = [ + "`olGetMemInfoSize` can be used to query the storage size required for the given query.", + "The provided pointer can point to any location inside the allocation.", + ]; + let params = [ + Param<"const void *", "Ptr", "pointer to the allocated memory", PARAM_IN>, + Param<"ol_mem_info_t", "PropName", "type of the info to retrieve", PARAM_IN>, + Param<"size_t", "PropSize", "the number of bytes pointed to by PropValue.", PARAM_IN>, + TypeTaggedParam<"void*", "PropValue", "array of bytes holding the info. " + "If Size is not equal to or greater to the real number of bytes needed to return the info " + "then the OL_ERRC_INVALID_SIZE error is returned and pPlatformInfo is not used.", PARAM_OUT, + TypeInfo<"PropName" , "PropSize">> + ]; + let returns = [ + Return<"OL_ERRC_INVALID_SIZE", [ + "`PropSize == 0`", + "If `PropSize` is less than the real number of bytes needed to return the info." + ]>, + Return<"OL_ERRC_NOT_FOUND", ["memory was not allocated by liboffload"]> + ]; +} + +def olGetMemInfoSize : Function { + let desc = "Returns the storage size of the given queue query."; + let details = [ + "The provided pointer can point to any location inside the allocation.", + ]; + let params = [ + Param<"const void *", "Ptr", "pointer to the allocated memory", PARAM_IN>, + Param<"ol_mem_info_t", "PropName", "type of the info to query", PARAM_IN>, + Param<"size_t*", "PropSizeRet", "pointer to the number of bytes required to store the query", PARAM_OUT> + ]; + let returns = [ + Return<"OL_ERRC_NOT_FOUND", ["memory was not allocated by liboffload"]> + ]; +} + def olMemcpy : Function { let desc = "Enqueue a memcpy operation."; let details = [ diff --git a/offload/liboffload/API/Program.td b/offload/liboffload/API/Program.td index 1f48f650cab7..7e11b3d8e331 100644 --- a/offload/liboffload/API/Program.td +++ b/offload/liboffload/API/Program.td @@ -24,6 +24,18 @@ def olCreateProgram : Function { let returns = []; } +def olIsValidBinary : Function { + let desc = "Validate if the binary image pointed to by `ProgData` is compatible with the device."; + let details = ["The provided `ProgData` will not be loaded onto the device"]; + let params = [ + Param<"ol_device_handle_t", "Device", "handle of the device", PARAM_IN>, + Param<"const void*", "ProgData", "pointer to the program binary data", PARAM_IN>, + Param<"size_t", "ProgDataSize", "size of the program binary in bytes", PARAM_IN>, + Param<"bool*", "Valid", "output is true if the image is compatible", PARAM_OUT> + ]; + let returns = []; +} + def olDestroyProgram : Function { let desc = "Destroy the program and free all underlying resources."; let details = []; diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp index 7e8e297831f4..c549ae04361d 100644 --- a/offload/liboffload/src/OffloadImpl.cpp +++ b/offload/liboffload/src/OffloadImpl.cpp @@ -39,12 +39,28 @@ using namespace llvm::omp::target; using namespace llvm::omp::target::plugin; using namespace error; +struct ol_platform_impl_t { + ol_platform_impl_t(std::unique_ptr<GenericPluginTy> Plugin, + ol_platform_backend_t BackendType) + : Plugin(std::move(Plugin)), BackendType(BackendType) {} + std::unique_ptr<GenericPluginTy> Plugin; + llvm::SmallVector<std::unique_ptr<ol_device_impl_t>> Devices; + ol_platform_backend_t BackendType; + + /// Complete all pending work for this platform and perform any needed + /// cleanup. + /// + /// After calling this function, no liboffload functions should be called with + /// this platform handle. + llvm::Error destroy(); +}; + // Handle type definitions. Ideally these would be 1:1 with the plugins, but // we add some additional data here for now to avoid churn in the plugin // interface. struct ol_device_impl_t { ol_device_impl_t(int DeviceNum, GenericDeviceTy *Device, - ol_platform_handle_t Platform, InfoTreeNode &&DevInfo) + ol_platform_impl_t &Platform, InfoTreeNode &&DevInfo) : DeviceNum(DeviceNum), Device(Device), Platform(Platform), Info(std::forward<InfoTreeNode>(DevInfo)) {} @@ -55,7 +71,7 @@ struct ol_device_impl_t { int DeviceNum; GenericDeviceTy *Device; - ol_platform_handle_t Platform; + ol_platform_impl_t &Platform; InfoTreeNode Info; llvm::SmallVector<__tgt_async_info *> OutstandingQueues; @@ -102,31 +118,17 @@ struct ol_device_impl_t { } }; -struct ol_platform_impl_t { - ol_platform_impl_t(std::unique_ptr<GenericPluginTy> Plugin, - ol_platform_backend_t BackendType) - : Plugin(std::move(Plugin)), BackendType(BackendType) {} - std::unique_ptr<GenericPluginTy> Plugin; - llvm::SmallVector<std::unique_ptr<ol_device_impl_t>> Devices; - ol_platform_backend_t BackendType; - - /// Complete all pending work for this platform and perform any needed - /// cleanup. - /// - /// After calling this function, no liboffload functions should be called with - /// this platform handle. - llvm::Error destroy() { - llvm::Error Result = Plugin::success(); - for (auto &D : Devices) - if (auto Err = D->destroy()) - Result = llvm::joinErrors(std::move(Result), std::move(Err)); +llvm::Error ol_platform_impl_t::destroy() { + llvm::Error Result = Plugin::success(); + for (auto &D : Devices) + if (auto Err = D->destroy()) + Result = llvm::joinErrors(std::move(Result), std::move(Err)); - if (auto Res = Plugin->deinit()) - Result = llvm::joinErrors(std::move(Result), std::move(Res)); + if (auto Res = Plugin->deinit()) + Result = llvm::joinErrors(std::move(Result), std::move(Res)); - return Result; - } -}; + return Result; +} struct ol_queue_impl_t { ol_queue_impl_t(__tgt_async_info *AsyncInfo, ol_device_handle_t Device) @@ -157,14 +159,11 @@ struct ol_event_impl_t { struct ol_program_impl_t { ol_program_impl_t(plugin::DeviceImageTy *Image, - std::unique_ptr<llvm::MemoryBuffer> ImageData, - const __tgt_device_image &DeviceImage) - : Image(Image), ImageData(std::move(ImageData)), - DeviceImage(DeviceImage) {} + llvm::MemoryBufferRef DeviceImage) + : Image(Image), DeviceImage(DeviceImage) {} plugin::DeviceImageTy *Image; - std::unique_ptr<llvm::MemoryBuffer> ImageData; std::mutex SymbolListMutex; - __tgt_device_image DeviceImage; + llvm::MemoryBufferRef DeviceImage; llvm::StringMap<std::unique_ptr<ol_symbol_impl_t>> KernelSymbols; llvm::StringMap<std::unique_ptr<ol_symbol_impl_t>> GlobalSymbols; }; @@ -185,6 +184,9 @@ namespace offload { struct AllocInfo { ol_device_handle_t Device; ol_alloc_type_t Type; + void *Start; + // One byte past the end + void *End; }; // Global shared state for liboffload @@ -203,12 +205,15 @@ struct OffloadContext { bool ValidationEnabled = true; DenseMap<void *, AllocInfo> AllocInfoMap{}; std::mutex AllocInfoMapMutex{}; - SmallVector<ol_platform_impl_t, 4> Platforms{}; + // Partitioned list of memory base addresses. Each element in this list is a + // key in AllocInfoMap + llvm::SmallVector<void *> AllocBases{}; + SmallVector<std::unique_ptr<ol_platform_impl_t>, 4> Platforms{}; size_t RefCount; ol_device_handle_t HostDevice() { // The host platform is always inserted last - return Platforms.back().Devices[0].get(); + return Platforms.back()->Devices[0].get(); } static OffloadContext &get() { @@ -247,38 +252,35 @@ Error initPlugins(OffloadContext &Context) { // Attempt to create an instance of each supported plugin. #define PLUGIN_TARGET(Name) \ do { \ - Context.Platforms.emplace_back(ol_platform_impl_t{ \ - std::unique_ptr<GenericPluginTy>(createPlugin_##Name()), \ - pluginNameToBackend(#Name)}); \ + if (StringRef(#Name) != "host") \ + Context.Platforms.emplace_back(std::make_unique<ol_platform_impl_t>( \ + std::unique_ptr<GenericPluginTy>(createPlugin_##Name()), \ + pluginNameToBackend(#Name))); \ } while (false); #include "Shared/Targets.def" // Preemptively initialize all devices in the plugin for (auto &Platform : Context.Platforms) { - // Do not use the host plugin - it isn't supported. - if (Platform.BackendType == OL_PLATFORM_BACKEND_UNKNOWN) - continue; - auto Err = Platform.Plugin->init(); + auto Err = Platform->Plugin->init(); [[maybe_unused]] std::string InfoMsg = toString(std::move(Err)); - for (auto DevNum = 0; DevNum < Platform.Plugin->number_of_devices(); + for (auto DevNum = 0; DevNum < Platform->Plugin->number_of_devices(); DevNum++) { - if (Platform.Plugin->init_device(DevNum) == OFFLOAD_SUCCESS) { - auto Device = &Platform.Plugin->getDevice(DevNum); + if (Platform->Plugin->init_device(DevNum) == OFFLOAD_SUCCESS) { + auto Device = &Platform->Plugin->getDevice(DevNum); auto Info = Device->obtainInfoImpl(); if (auto Err = Info.takeError()) return Err; - Platform.Devices.emplace_back(std::make_unique<ol_device_impl_t>( - DevNum, Device, &Platform, std::move(*Info))); + Platform->Devices.emplace_back(std::make_unique<ol_device_impl_t>( + DevNum, Device, *Platform, std::move(*Info))); } } } // Add the special host device auto &HostPlatform = Context.Platforms.emplace_back( - ol_platform_impl_t{nullptr, OL_PLATFORM_BACKEND_HOST}); - HostPlatform.Devices.emplace_back( - std::make_unique<ol_device_impl_t>(-1, nullptr, nullptr, InfoTreeNode{})); - Context.HostDevice()->Platform = &HostPlatform; + std::make_unique<ol_platform_impl_t>(nullptr, OL_PLATFORM_BACKEND_HOST)); + HostPlatform->Devices.emplace_back(std::make_unique<ol_device_impl_t>( + -1, nullptr, *HostPlatform, InfoTreeNode{})); Context.TracingEnabled = std::getenv("OFFLOAD_TRACE"); Context.ValidationEnabled = !std::getenv("OFFLOAD_DISABLE_VALIDATION"); @@ -315,10 +317,10 @@ Error olShutDown_impl() { for (auto &P : OldContext->Platforms) { // Host plugin is nullptr and has no deinit - if (!P.Plugin || !P.Plugin->is_initialized()) + if (!P->Plugin || !P->Plugin->is_initialized()) continue; - if (auto Res = P.destroy()) + if (auto Res = P->destroy()) Result = llvm::joinErrors(std::move(Result), std::move(Res)); } @@ -383,7 +385,7 @@ Error olGetDeviceInfoImplDetail(ol_device_handle_t Device, // These are not implemented by the plugin interface switch (PropName) { case OL_DEVICE_INFO_PLATFORM: - return Info.write<void *>(Device->Platform); + return Info.write<void *>(&Device->Platform); case OL_DEVICE_INFO_TYPE: return Info.write<ol_device_type_t>(OL_DEVICE_TYPE_GPU); @@ -516,7 +518,7 @@ Error olGetDeviceInfoImplDetailHost(ol_device_handle_t Device, switch (PropName) { case OL_DEVICE_INFO_PLATFORM: - return Info.write<void *>(Device->Platform); + return Info.write<void *>(&Device->Platform); case OL_DEVICE_INFO_TYPE: return Info.write<ol_device_type_t>(OL_DEVICE_TYPE_HOST); case OL_DEVICE_INFO_NAME: @@ -594,7 +596,7 @@ Error olGetDeviceInfoSize_impl(ol_device_handle_t Device, Error olIterateDevices_impl(ol_device_iterate_cb_t Callback, void *UserData) { for (auto &Platform : OffloadContext::get().Platforms) { - for (auto &Device : Platform.Devices) { + for (auto &Device : Platform->Devices) { if (!Callback(Device.get(), UserData)) { break; } @@ -616,20 +618,61 @@ TargetAllocTy convertOlToPluginAllocTy(ol_alloc_type_t Type) { } } +constexpr size_t MAX_ALLOC_TRIES = 50; Error olMemAlloc_impl(ol_device_handle_t Device, ol_alloc_type_t Type, size_t Size, void **AllocationOut) { - auto Alloc = - Device->Device->dataAlloc(Size, nullptr, convertOlToPluginAllocTy(Type)); - if (!Alloc) - return Alloc.takeError(); + SmallVector<void *> Rejects; + + // Repeat the allocation up to a certain amount of times. If it happens to + // already be allocated (e.g. by a device from another vendor) throw it away + // and try again. + for (size_t Count = 0; Count < MAX_ALLOC_TRIES; Count++) { + auto NewAlloc = Device->Device->dataAlloc(Size, nullptr, + convertOlToPluginAllocTy(Type)); + if (!NewAlloc) + return NewAlloc.takeError(); + + void *NewEnd = &static_cast<char *>(*NewAlloc)[Size]; + auto &AllocBases = OffloadContext::get().AllocBases; + auto &AllocInfoMap = OffloadContext::get().AllocInfoMap; + { + std::lock_guard<std::mutex> Lock(OffloadContext::get().AllocInfoMapMutex); + + // Check that this memory region doesn't overlap another one + // That is, the start of this allocation needs to be after another + // allocation's end point, and the end of this allocation needs to be + // before the next one's start. + // `Gap` is the first alloc who ends after the new alloc's start point. + auto Gap = + std::lower_bound(AllocBases.begin(), AllocBases.end(), *NewAlloc, + [&](const void *Iter, const void *Val) { + return AllocInfoMap.at(Iter).End <= Val; + }); + if (Gap == AllocBases.end() || NewEnd <= AllocInfoMap.at(*Gap).Start) { + // Success, no conflict + AllocInfoMap.insert_or_assign( + *NewAlloc, AllocInfo{Device, Type, *NewAlloc, NewEnd}); + AllocBases.insert( + std::lower_bound(AllocBases.begin(), AllocBases.end(), *NewAlloc), + *NewAlloc); + *AllocationOut = *NewAlloc; + + for (void *R : Rejects) + if (auto Err = + Device->Device->dataDelete(R, convertOlToPluginAllocTy(Type))) + return Err; + return Error::success(); + } - *AllocationOut = *Alloc; - { - std::lock_guard<std::mutex> Lock(OffloadContext::get().AllocInfoMapMutex); - OffloadContext::get().AllocInfoMap.insert_or_assign( - *Alloc, AllocInfo{Device, Type}); + // To avoid the next attempt allocating the same memory we just freed, we + // hold onto it until we complete the allocation + Rejects.push_back(*NewAlloc); + } } - return Error::success(); + + // We've tried multiple times, and can't allocate a non-overlapping region. + return createOffloadError(ErrorCode::BACKEND_FAILURE, + "failed to allocate non-overlapping memory"); } Error olMemFree_impl(void *Address) { @@ -645,6 +688,9 @@ Error olMemFree_impl(void *Address) { Device = AllocInfo.Device; Type = AllocInfo.Type; OffloadContext::get().AllocInfoMap.erase(Address); + + auto &Bases = OffloadContext::get().AllocBases; + Bases.erase(std::lower_bound(Bases.begin(), Bases.end(), Address)); } if (auto Res = @@ -654,6 +700,60 @@ Error olMemFree_impl(void *Address) { return Error::success(); } +Error olGetMemInfoImplDetail(const void *Ptr, ol_mem_info_t PropName, + size_t PropSize, void *PropValue, + size_t *PropSizeRet) { + InfoWriter Info(PropSize, PropValue, PropSizeRet); + std::lock_guard<std::mutex> Lock(OffloadContext::get().AllocInfoMapMutex); + + auto &AllocBases = OffloadContext::get().AllocBases; + auto &AllocInfoMap = OffloadContext::get().AllocInfoMap; + const AllocInfo *Alloc = nullptr; + if (AllocInfoMap.contains(Ptr)) { + // Fast case, we have been given the base pointer directly + Alloc = &AllocInfoMap.at(Ptr); + } else { + // Slower case, we need to look up the base pointer first + // Find the first memory allocation whose end is after the target pointer, + // and then check to see if it is in range + auto Loc = std::lower_bound(AllocBases.begin(), AllocBases.end(), Ptr, + [&](const void *Iter, const void *Val) { + return AllocInfoMap.at(Iter).End <= Val; + }); + if (Loc == AllocBases.end() || Ptr < AllocInfoMap.at(*Loc).Start) + return Plugin::error(ErrorCode::NOT_FOUND, + "allocated memory information not found"); + Alloc = &AllocInfoMap.at(*Loc); + } + + switch (PropName) { + case OL_MEM_INFO_DEVICE: + return Info.write<ol_device_handle_t>(Alloc->Device); + case OL_MEM_INFO_BASE: + return Info.write<void *>(Alloc->Start); + case OL_MEM_INFO_SIZE: + return Info.write<size_t>(static_cast<char *>(Alloc->End) - + static_cast<char *>(Alloc->Start)); + case OL_MEM_INFO_TYPE: + return Info.write<ol_alloc_type_t>(Alloc->Type); + default: + return createOffloadError(ErrorCode::INVALID_ENUMERATION, + "olGetMemInfo enum '%i' is invalid", PropName); + } + + return Error::success(); +} + +Error olGetMemInfo_impl(const void *Ptr, ol_mem_info_t PropName, + size_t PropSize, void *PropValue) { + return olGetMemInfoImplDetail(Ptr, PropName, PropSize, PropValue, nullptr); +} + +Error olGetMemInfoSize_impl(const void *Ptr, ol_mem_info_t PropName, + size_t *PropSizeRet) { + return olGetMemInfoImplDetail(Ptr, PropName, 0, nullptr, PropSizeRet); +} + Error olCreateQueue_impl(ol_device_handle_t Device, ol_queue_handle_t *Queue) { auto CreatedQueue = std::make_unique<ol_queue_impl_t>(nullptr, Device); @@ -890,29 +990,23 @@ Error olMemFill_impl(ol_queue_handle_t Queue, void *Ptr, size_t PatternSize, Error olCreateProgram_impl(ol_device_handle_t Device, const void *ProgData, size_t ProgDataSize, ol_program_handle_t *Program) { - // Make a copy of the program binary in case it is released by the caller. - auto ImageData = MemoryBuffer::getMemBufferCopy( - StringRef(reinterpret_cast<const char *>(ProgData), ProgDataSize)); - - auto DeviceImage = __tgt_device_image{ - const_cast<char *>(ImageData->getBuffer().data()), - const_cast<char *>(ImageData->getBuffer().data()) + ProgDataSize, nullptr, - nullptr}; - - ol_program_handle_t Prog = - new ol_program_impl_t(nullptr, std::move(ImageData), DeviceImage); - - auto Res = - Device->Device->loadBinary(Device->Device->Plugin, &Prog->DeviceImage); - if (!Res) { - delete Prog; + StringRef Buffer(reinterpret_cast<const char *>(ProgData), ProgDataSize); + Expected<plugin::DeviceImageTy *> Res = + Device->Device->loadBinary(Device->Device->Plugin, Buffer); + if (!Res) return Res.takeError(); - } - assert(*Res != nullptr && "loadBinary returned nullptr"); + assert(*Res && "loadBinary returned nullptr"); - Prog->Image = *Res; - *Program = Prog; + *Program = new ol_program_impl_t(*Res, (*Res)->getMemoryBuffer()); + return Error::success(); +} +Error olIsValidBinary_impl(ol_device_handle_t Device, const void *ProgData, + size_t ProgDataSize, bool *IsValid) { + StringRef Buffer(reinterpret_cast<const char *>(ProgData), ProgDataSize); + *IsValid = Device->Device ? Device->Device->Plugin.isDeviceCompatible( + Device->Device->getDeviceId(), Buffer) + : false; return Error::success(); } diff --git a/offload/libomptarget/OpenMP/InteropAPI.cpp b/offload/libomptarget/OpenMP/InteropAPI.cpp index eb5425ecbf06..c55ef2c2e672 100644 --- a/offload/libomptarget/OpenMP/InteropAPI.cpp +++ b/offload/libomptarget/OpenMP/InteropAPI.cpp @@ -124,7 +124,7 @@ void *getProperty<void *>(omp_interop_val_t &InteropVal, case omp_ipr_device_context: return InteropVal.device_info.Context; case omp_ipr_targetsync: - return InteropVal.async_info->Queue; + return InteropVal.async_info ? InteropVal.async_info->Queue : nullptr; default:; } getTypeMismatch(Property, Err); @@ -167,7 +167,6 @@ bool getPropertyCheck(omp_interop_val_t **InteropPtr, omp_interop_property_t property_id, \ int *err) { \ omp_interop_val_t *interop_val = (omp_interop_val_t *)interop; \ - assert((interop_val)->interop_type == kmp_interop_type_targetsync); \ if (!getPropertyCheck(&interop_val, property_id, err)) { \ return (RETURN_TYPE)(0); \ } \ @@ -275,8 +274,8 @@ omp_interop_val_t *__tgt_interop_get(ident_t *LocRef, int32_t InteropType, return Interop; } -int __tgt_interop_use(ident_t *LocRef, omp_interop_val_t *Interop, - interop_ctx_t *Ctx, dep_pack_t *Deps) { +int __tgt_interop_use60(ident_t *LocRef, omp_interop_val_t *Interop, + interop_ctx_t *Ctx, dep_pack_t *Deps) { bool Nowait = Ctx->flags.nowait; DP("Call to %s with interop " DPxMOD ", nowait %" PRId32 "\n", __func__, DPxPTR(Interop), Nowait); @@ -359,6 +358,40 @@ EXTERN int ompx_interop_add_completion_callback(omp_interop_val_t *Interop, return omp_irc_success; } +// Backwards compatibility wrappers +void __tgt_interop_init(ident_t *LocRef, int32_t Gtid, + omp_interop_val_t *&InteropPtr, int32_t InteropType, + int32_t DeviceId, int32_t Ndeps, + kmp_depend_info_t *DepList, int32_t HaveNowait) { + constexpr int32_t old_kmp_interop_type_targetsync = 2; + interop_ctx_t Ctx = {0, {false, (bool)HaveNowait, 0}, Gtid}; + dep_pack_t Deps = {Ndeps, 0, DepList, nullptr}; + InteropPtr = + __tgt_interop_get(LocRef, + InteropType == old_kmp_interop_type_targetsync + ? kmp_interop_type_targetsync + : kmp_interop_type_target, + DeviceId, 0, nullptr, &Ctx, Ndeps ? &Deps : nullptr); +} + +void __tgt_interop_use(ident_t *LocRef, int32_t Gtid, + omp_interop_val_t *&InteropPtr, int32_t DeviceId, + int32_t Ndeps, kmp_depend_info_t *DepList, + int32_t HaveNowait) { + interop_ctx_t Ctx = {0, {false, (bool)HaveNowait, 0}, Gtid}; + dep_pack_t Deps = {Ndeps, 0, DepList, nullptr}; + __tgt_interop_use60(LocRef, InteropPtr, &Ctx, Ndeps ? &Deps : nullptr); +} + +void __tgt_interop_destroy(ident_t *LocRef, int32_t Gtid, + omp_interop_val_t *&InteropPtr, int32_t DeviceId, + int32_t Ndeps, kmp_depend_info_t *DepList, + int32_t HaveNowait) { + interop_ctx_t Ctx = {0, {false, (bool)HaveNowait, 0}, Gtid}; + dep_pack_t Deps = {Ndeps, 0, DepList, nullptr}; + __tgt_interop_release(LocRef, InteropPtr, &Ctx, Ndeps ? &Deps : nullptr); +} + } // extern "C" llvm::Expected<DeviceTy &> omp_interop_val_t::getDevice() const { diff --git a/offload/libomptarget/PluginManager.cpp b/offload/libomptarget/PluginManager.cpp index b57a2f815cba..c8d6b42114d0 100644 --- a/offload/libomptarget/PluginManager.cpp +++ b/offload/libomptarget/PluginManager.cpp @@ -219,7 +219,10 @@ void PluginManager::registerLib(__tgt_bin_desc *Desc) { // Scan the RTLs that have associated images until we find one that supports // the current image. for (auto &R : plugins()) { - if (!R.is_plugin_compatible(Img)) + StringRef Buffer(reinterpret_cast<const char *>(Img->ImageStart), + utils::getPtrDiff(Img->ImageEnd, Img->ImageStart)); + + if (!R.isPluginCompatible(Buffer)) continue; if (!initializePlugin(R)) @@ -242,7 +245,7 @@ void PluginManager::registerLib(__tgt_bin_desc *Desc) { continue; } - if (!R.is_device_compatible(DeviceId, Img)) + if (!R.isDeviceCompatible(DeviceId, Buffer)) continue; DP("Image " DPxMOD " is compatible with RTL %s device %d!\n", diff --git a/offload/libomptarget/exports b/offload/libomptarget/exports index 8e2db6ba8bba..1374bfea8151 100644 --- a/offload/libomptarget/exports +++ b/offload/libomptarget/exports @@ -68,8 +68,11 @@ VERS1.0 { omp_get_interop_int; omp_get_interop_name; omp_get_interop_type_desc; - __tgt_interop_get; + __tgt_interop_init; __tgt_interop_use; + __tgt_interop_destroy; + __tgt_interop_get; + __tgt_interop_use60; __tgt_interop_release; __tgt_target_sync; __llvmPushCallConfiguration; diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index 4c8eba1e7180..a1950cbb6290 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -330,6 +330,54 @@ int targetDataMapper(ident_t *Loc, DeviceTy &Device, void *ArgBase, void *Arg, return Rc; } +/// Returns a buffer of the requested \p Size, to be used as the source for +/// `submitData`. +/// +/// For small buffers (`Size <= sizeof(void*)`), uses \p AsyncInfo's +/// getVoidPtrLocation(). +/// For larger buffers, creates a dynamic buffer which will be eventually +/// deleted by \p AsyncInfo's post-processing callback. +static char *getOrCreateSourceBufferForSubmitData(AsyncInfoTy &AsyncInfo, + int64_t Size) { + constexpr int64_t VoidPtrSize = sizeof(void *); + + if (Size <= VoidPtrSize) { + void *&BufferElement = AsyncInfo.getVoidPtrLocation(); + return reinterpret_cast<char *>(&BufferElement); + } + + // Create a dynamic buffer for larger data and schedule its deletion. + char *DataBuffer = new char[Size]; + AsyncInfo.addPostProcessingFunction([DataBuffer]() { + delete[] DataBuffer; + return OFFLOAD_SUCCESS; + }); + return DataBuffer; +} + +/// Calculates the target pointee base by applying the host +/// pointee begin/base delta to the target pointee begin. +/// +/// ``` +/// TgtPteeBase = TgtPteeBegin - (HstPteeBegin - HstPteeBase) +/// ``` +static void *calculateTargetPointeeBase(void *HstPteeBase, void *HstPteeBegin, + void *TgtPteeBegin) { + uint64_t Delta = reinterpret_cast<uint64_t>(HstPteeBegin) - + reinterpret_cast<uint64_t>(HstPteeBase); + void *TgtPteeBase = reinterpret_cast<void *>( + reinterpret_cast<uint64_t>(TgtPteeBegin) - Delta); + + DP("HstPteeBase: " DPxMOD ", HstPteeBegin: " DPxMOD + ", Delta (HstPteeBegin - HstPteeBase): %" PRIu64 ".\n", + DPxPTR(HstPteeBase), DPxPTR(HstPteeBegin), Delta); + DP("TgtPteeBase (TgtPteeBegin - Delta): " DPxMOD ", TgtPteeBegin : " DPxMOD + "\n", + DPxPTR(TgtPteeBase), DPxPTR(TgtPteeBegin)); + + return TgtPteeBase; +} + /// Utility function to perform a pointer attachment operation. /// /// For something like: @@ -396,38 +444,15 @@ static int performPointerAttachment(DeviceTy &Device, AsyncInfoTy &AsyncInfo, assert(PtrTPR.getEntry() && "Need a valid pointer entry to perform pointer-attachment"); - int64_t VoidPtrSize = sizeof(void *); + constexpr int64_t VoidPtrSize = sizeof(void *); assert(HstPtrSize >= VoidPtrSize && "PointerSize is too small"); - uint64_t Delta = reinterpret_cast<uint64_t>(HstPteeBegin) - - reinterpret_cast<uint64_t>(HstPteeBase); - void *TgtPteeBase = reinterpret_cast<void *>( - reinterpret_cast<uint64_t>(TgtPteeBegin) - Delta); - DP("HstPteeBase: " DPxMOD ", HstPteeBegin: " DPxMOD - ", Delta (HstPteeBegin - HstPteeBase): %" PRIu64 ".\n", - DPxPTR(HstPteeBase), DPxPTR(HstPteeBegin), Delta); - DP("TgtPteeBase (TgtPteeBegin - Delta): " DPxMOD ", TgtPteeBegin : " DPxMOD - "\n", - DPxPTR(TgtPteeBase), DPxPTR(TgtPteeBegin)); + void *TgtPteeBase = + calculateTargetPointeeBase(HstPteeBase, HstPteeBegin, TgtPteeBegin); // Add shadow pointer tracking - // TODO: Support shadow-tracking of larger than VoidPtrSize pointers, - // to support restoration of Fortran descriptors. Currently, this check - // would return false, even if the host Fortran descriptor had been - // updated since its previous map, and we should have updated its - // device counterpart. e.g. - // - // !$omp target enter data map(x(1:100)) ! (1) - // p => x(10: 19) - // !$omp target enter data map(p, p(:)) ! (2) - // p => x(5: 9) - // !$omp target enter data map(attach(always): p(:)) ! (3) - // - // While PtrAddr(&desc_p) and PteeBase(&p(1)) are same for (2) and (3), the - // pointer attachment for (3) needs to update the bounds information - // in the descriptor of p on device. if (!PtrTPR.getEntry()->addShadowPointer( - ShadowPtrInfoTy{HstPtrAddr, HstPteeBase, TgtPtrAddr, TgtPteeBase})) { + ShadowPtrInfoTy{HstPtrAddr, TgtPtrAddr, TgtPteeBase, HstPtrSize})) { DP("Pointer " DPxMOD " is already attached to " DPxMOD "\n", DPxPTR(TgtPtrAddr), DPxPTR(TgtPteeBase)); return OFFLOAD_SUCCESS; @@ -450,48 +475,32 @@ static int performPointerAttachment(DeviceTy &Device, AsyncInfoTy &AsyncInfo, return OFFLOAD_SUCCESS; }; - bool IsPtrAFortranDescriptor = HstPtrSize > VoidPtrSize; - if (!IsPtrAFortranDescriptor) { - // For "regular" pointers, we can use the VoidPtrLocation from AsyncInfo as - // the buffer space for the submission. - void *&BufferElement = AsyncInfo.getVoidPtrLocation(); - BufferElement = TgtPteeBase; - - // Submit the updated pointer value to device - return HandleSubmitResult(Device.submitData( - TgtPtrAddr, &BufferElement, VoidPtrSize, AsyncInfo, PtrTPR.getEntry())); + // Get a buffer to be used as the source for data submission. + char *SrcBuffer = getOrCreateSourceBufferForSubmitData(AsyncInfo, HstPtrSize); + + // The pointee's address should occupy the first VoidPtrSize bytes + // irrespective of HstPtrSize. + std::memcpy(SrcBuffer, &TgtPteeBase, VoidPtrSize); + + // For larger "pointers" (e.g., Fortran descriptors), copy remaining + // descriptor fields from the host descriptor into the buffer. + if (HstPtrSize > VoidPtrSize) { + uint64_t HstDescriptorFieldsSize = HstPtrSize - VoidPtrSize; + void *HstDescriptorFieldsAddr = + reinterpret_cast<char *>(HstPtrAddr) + VoidPtrSize; + std::memcpy(SrcBuffer + VoidPtrSize, HstDescriptorFieldsAddr, + HstDescriptorFieldsSize); + + DP("Updating %" PRId64 " bytes of descriptor (" DPxMOD + ") (pointer + %" PRId64 " additional bytes from host descriptor " DPxMOD + ")\n", + HstPtrSize, DPxPTR(TgtPtrAddr), HstDescriptorFieldsSize, + DPxPTR(HstDescriptorFieldsAddr)); } - // For larger "pointers" (like Fortran's descriptors), we create a dynamic - // buffer, which will be eventually destroyed by AsyncInfo's post-processing - // callback. - char *DataBuffer = new char[HstPtrSize]; - - // For such descriptors, to the first VoidPtrSize bytes, we store the - // pointee's device address. - std::memcpy(DataBuffer, &TgtPteeBase, sizeof(void *)); - - // And to the remaining bytes, we copy the remaining contents of the host - // descriptor after the initial VoidPtrSize bytes. - uint64_t HstDescriptorFieldsSize = HstPtrSize - VoidPtrSize; - void *HstDescriptorFieldsAddr = - reinterpret_cast<char *>(HstPtrAddr) + VoidPtrSize; - std::memcpy(DataBuffer + VoidPtrSize, HstDescriptorFieldsAddr, - HstDescriptorFieldsSize); - - DP("Updating %" PRId64 " bytes of descriptor (" DPxMOD ") (pointer + %" PRId64 - " additional bytes from host descriptor " DPxMOD ")\n", - HstPtrSize, DPxPTR(TgtPtrAddr), HstDescriptorFieldsSize, - DPxPTR(HstDescriptorFieldsAddr)); - - // Submit the entire buffer to device - int SubmitResult = Device.submitData(TgtPtrAddr, DataBuffer, HstPtrSize, + // Submit the populated source buffer to device. + int SubmitResult = Device.submitData(TgtPtrAddr, SrcBuffer, HstPtrSize, AsyncInfo, PtrTPR.getEntry()); - - AsyncInfo.addPostProcessingFunction([DataBuffer]() -> int { - delete[] DataBuffer; - return OFFLOAD_SUCCESS; - }); return HandleSubmitResult(SubmitResult); } @@ -540,10 +549,17 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, // ATTACH map-types are supposed to be handled after all mapping for the // construct is done. Defer their processing. if (ArgTypes[I] & OMP_TGT_MAPTYPE_ATTACH) { - AttachInfo->AttachEntries.emplace_back( - /*PointerBase=*/HstPtrBase, /*PointeeBegin=*/HstPtrBegin, - /*PointerSize=*/DataSize, /*MapType=*/ArgTypes[I], - /*PointeeName=*/HstPtrName); + const bool IsCorrespondingPointerInit = + (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE); + // We don't need to keep track of PRIVATE | ATTACH entries. They + // represent corresponding-pointer-initialization, and are handled + // similar to firstprivate (PRIVATE | TO) entries by + // PrivateArgumentManager. + if (!IsCorrespondingPointerInit) + AttachInfo->AttachEntries.emplace_back( + /*PointerBase=*/HstPtrBase, /*PointeeBegin=*/HstPtrBegin, + /*PointerSize=*/DataSize, /*MapType=*/ArgTypes[I], + /*PointeeName=*/HstPtrName); DP("Deferring ATTACH map-type processing for argument %d\n", I); continue; @@ -940,17 +956,29 @@ postProcessingTargetDataEnd(DeviceTy *Device, DelEntry = false; } - // If we copied back to the host a struct/array containing pointers, - // we need to restore the original host pointer values from their - // shadow copies. If the struct is going to be deallocated, remove any - // remaining shadow pointer entries for this struct. + // If we copied back to the host a struct/array containing pointers, or + // Fortran descriptors (which are larger than a "void *"), we need to + // restore the original host pointer/descriptor values from their shadow + // copies. If the struct is going to be deallocated, remove any remaining + // shadow pointer entries for this struct. const bool HasFrom = ArgType & OMP_TGT_MAPTYPE_FROM; if (HasFrom) { Entry->foreachShadowPointerInfo([&](const ShadowPtrInfoTy &ShadowPtr) { - *ShadowPtr.HstPtrAddr = ShadowPtr.HstPtrVal; - DP("Restoring original host pointer value " DPxMOD " for host " - "pointer " DPxMOD "\n", - DPxPTR(ShadowPtr.HstPtrVal), DPxPTR(ShadowPtr.HstPtrAddr)); + constexpr int64_t VoidPtrSize = sizeof(void *); + if (ShadowPtr.PtrSize > VoidPtrSize) { + DP("Restoring host descriptor " DPxMOD + " to its original content (%" PRId64 + " bytes), containing pointee address " DPxMOD "\n", + DPxPTR(ShadowPtr.HstPtrAddr), ShadowPtr.PtrSize, + DPxPTR(ShadowPtr.HstPtrContent.data())); + } else { + DP("Restoring host pointer " DPxMOD " to its original value " DPxMOD + "\n", + DPxPTR(ShadowPtr.HstPtrAddr), + DPxPTR(ShadowPtr.HstPtrContent.data())); + } + std::memcpy(ShadowPtr.HstPtrAddr, ShadowPtr.HstPtrContent.data(), + ShadowPtr.PtrSize); return OFFLOAD_SUCCESS; }); } @@ -1163,12 +1191,22 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase, if (TPR.getEntry()) { int Ret = TPR.getEntry()->foreachShadowPointerInfo( [&](ShadowPtrInfoTy &ShadowPtr) { - DP("Restoring original target pointer value " DPxMOD " for target " - "pointer " DPxMOD "\n", - DPxPTR(ShadowPtr.TgtPtrVal), DPxPTR(ShadowPtr.TgtPtrAddr)); + constexpr int64_t VoidPtrSize = sizeof(void *); + if (ShadowPtr.PtrSize > VoidPtrSize) { + DP("Restoring target descriptor " DPxMOD + " to its original content (%" PRId64 + " bytes), containing pointee address " DPxMOD "\n", + DPxPTR(ShadowPtr.TgtPtrAddr), ShadowPtr.PtrSize, + DPxPTR(ShadowPtr.TgtPtrContent.data())); + } else { + DP("Restoring target pointer " DPxMOD + " to its original value " DPxMOD "\n", + DPxPTR(ShadowPtr.TgtPtrAddr), + DPxPTR(ShadowPtr.TgtPtrContent.data())); + } Ret = Device.submitData(ShadowPtr.TgtPtrAddr, - (void *)&ShadowPtr.TgtPtrVal, - sizeof(void *), AsyncInfo); + ShadowPtr.TgtPtrContent.data(), + ShadowPtr.PtrSize, AsyncInfo); if (Ret != OFFLOAD_SUCCESS) { REPORT("Copying data to device failed.\n"); return OFFLOAD_FAIL; @@ -1193,15 +1231,26 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase, } // Wait for device-to-host memcopies for whole struct to complete, - // before restoring the correct host pointer. + // before restoring the correct host pointer/descriptor. if (auto *Entry = TPR.getEntry()) { AsyncInfo.addPostProcessingFunction([=]() -> int { int Ret = Entry->foreachShadowPointerInfo( [&](const ShadowPtrInfoTy &ShadowPtr) { - *ShadowPtr.HstPtrAddr = ShadowPtr.HstPtrVal; - DP("Restoring original host pointer value " DPxMOD - " for host pointer " DPxMOD "\n", - DPxPTR(ShadowPtr.HstPtrVal), DPxPTR(ShadowPtr.HstPtrAddr)); + constexpr int64_t VoidPtrSize = sizeof(void *); + if (ShadowPtr.PtrSize > VoidPtrSize) { + DP("Restoring host descriptor " DPxMOD + " to its original content (%" PRId64 + " bytes), containing pointee address " DPxMOD "\n", + DPxPTR(ShadowPtr.HstPtrAddr), ShadowPtr.PtrSize, + DPxPTR(ShadowPtr.HstPtrContent.data())); + } else { + DP("Restoring host pointer " DPxMOD + " to its original value " DPxMOD "\n", + DPxPTR(ShadowPtr.HstPtrAddr), + DPxPTR(ShadowPtr.HstPtrContent.data())); + } + std::memcpy(ShadowPtr.HstPtrAddr, ShadowPtr.HstPtrContent.data(), + ShadowPtr.PtrSize); return OFFLOAD_SUCCESS; }); Entry->unlock(); @@ -1379,13 +1428,24 @@ class PrivateArgumentManagerTy { uint32_t Padding; /// Host pointer name map_var_info_t HstPtrName = nullptr; + /// For corresponding-pointer-initialization: host pointee base address. + void *HstPteeBase = nullptr; + /// For corresponding-pointer-initialization: host pointee begin address. + void *HstPteeBegin = nullptr; + /// Whether this argument needs corresponding-pointer-initialization. + bool IsCorrespondingPointerInit = false; FirstPrivateArgInfoTy(int Index, void *HstPtr, uint32_t Size, uint32_t Alignment, uint32_t Padding, - map_var_info_t HstPtrName = nullptr) + map_var_info_t HstPtrName = nullptr, + void *HstPteeBase = nullptr, + void *HstPteeBegin = nullptr, + bool IsCorrespondingPointerInit = false) : HstPtrBegin(reinterpret_cast<char *>(HstPtr)), HstPtrEnd(HstPtrBegin + Size), Index(Index), Alignment(Alignment), - Size(Size), Padding(Padding), HstPtrName(HstPtrName) {} + Size(Size), Padding(Padding), HstPtrName(HstPtrName), + HstPteeBase(HstPteeBase), HstPteeBegin(HstPteeBegin), + IsCorrespondingPointerInit(IsCorrespondingPointerInit) {} }; /// A vector of target pointers for all private arguments @@ -1403,6 +1463,153 @@ class PrivateArgumentManagerTy { /// A pointer to a \p AsyncInfoTy object AsyncInfoTy &AsyncInfo; + /// \returns the value of the target pointee's base to be used for + /// corresponding-pointer-initialization. + void *getTargetPointeeBaseForCorrespondingPointerInitialization( + void *HstPteeBase, void *HstPteeBegin) { + // See if the pointee's begin address has corresponding storage on device. + void *TgtPteeBegin = [&]() -> void * { + if (!HstPteeBegin) { + DP("Corresponding-pointer-initialization: pointee begin address is " + "null\n"); + return nullptr; + } + + return Device.getMappingInfo() + .getTgtPtrBegin(HstPteeBegin, /*Size=*/0, /*UpdateRefCount=*/false, + /*UseHoldRefCount=*/false) + .TargetPointer; + }(); + + // If it does, we calculate target pointee base using it, and return it. + // Otherwise, we retain the host pointee's base as the target pointee base + // of the initialized pointer. It's the user's responsibility to ensure + // that if a lookup fails, the host pointee is accessible on the device. + return TgtPteeBegin ? calculateTargetPointeeBase(HstPteeBase, HstPteeBegin, + TgtPteeBegin) + : HstPteeBase; + } + + /// Initialize the source buffer for corresponding-pointer-initialization. + /// + /// It computes and stores the target pointee base address (or the host + /// pointee's base address, if lookup of target pointee fails) to the first + /// `sizeof(void*)` bytes of \p Buffer, and for larger pointers + /// (Fortran descriptors), the remaining fields of the host descriptor + /// \p HstPtr after those `sizeof(void*)` bytes. + /// + /// Corresponding-pointer-initialization represents the initialization of the + /// private version of a base-pointer/referring-pointer on a target construct. + /// + /// For example, for the following test: + /// ```cpp + /// int x[10]; + /// int *px = &x[0]; + /// ... + /// #pragma omp target data map(tofrom:px) + /// { + /// int **ppx = omp_get_mapped_ptr(&px, omp_get_default_device()); + /// #pragma omp target map(tofrom:px[1]) is_device_ptr(ppx) + /// { + /// foo(px, ppx); + /// } + /// } + /// ``` + /// The following shows a possible way to implement the mapping of `px`, + /// which is pre-determined firstprivate and should get initialized + /// via corresponding-pointer-initialization: + /// + /// (A) Possible way to implement the above with PRIVATE | ATTACH: + /// ```llvm + /// ; maps for px: + /// ; &px[0], &px[1], sizeof(px[1]), TO | FROM // (1) + /// ; &px, &px[1], sizeof(px), ATTACH // (2) + /// ; &px, &px[1], sizeof(px), PRIVATE | ATTACH | PARAM // (3) + /// call... @__omp_outlined...(ptr %px, ptr %ppx) + /// define ... @__omp_outlined(ptr %px, ptr %ppx) {... + /// foo(%px, %ppx) + /// ...} + /// ``` + /// `(1)` maps the pointee `px[1]. + /// `(2)` attaches it to the mapped version of `px`. It can be controlled by + /// the user based on the `attach(auto/always/never)` map-type modifier. + /// `(3)` privatizes and initializes the private pointer `px`, and passes it + /// into the kernel as the argument `%px`. Can be skipped if `px` is not + /// referenced in the target construct. + /// + /// While this method is not too beneficial compared to just doing the + /// initialization in the body of the kernel, like: + /// (B) Possible way to implement the above without PRIVATE | ATTACH: + /// ```llvm + /// ; maps for px: + /// ; &px[0], &px[1], sizeof(px[1]), TO | FROM | PARAM // (4) + /// ; &px, &px[1], sizeof(px), ATTACH // (5) + /// call... @__omp_outlined...(ptr %px0, ptr %ppx) + /// define ... __omp_outlined...(ptr %px0, ptr %ppx) { + /// %px = alloca ptr; + /// store ptr %px0, ptr %px + /// foo(%px, %ppx) + /// } + /// ``` + /// + /// (B) is not so convenient for Fortran descriptors, because in + /// addition to the lookup, the remaining fields of the descriptor have + /// to be passed into the kernel to initialize the private copy, which + /// makes (A) a cleaner option for them. e.g. + /// ```f90 + /// integer, pointer :: p(:) + /// !$omp target map(p(1)) + /// ``` + /// + /// (C) Possible mapping for the above Fortran test using PRIVATE | ATTACH: + /// ```llvm + /// ; maps for p: + /// ; &p(1), &p(1), sizeof(p(1)), TO | FROM + /// ; &ref_ptr(p), &p(1), sizeof(ref_ptr(p)), ATTACH + /// ; &ref_ptr(p), &p(1), sizeof(ref_ptr(p)), PRIVATE | ATTACH | PARAM + /// call... @__omp_outlined...(ptr %ref_ptr_of_p) + void initBufferForCorrespondingPointerInitialization(char *Buffer, + void *HstPtr, + int64_t HstPtrSize, + void *HstPteeBase, + void *HstPteeBegin) { + constexpr int64_t VoidPtrSize = sizeof(void *); + assert(HstPtrSize >= VoidPtrSize && + "corresponding-pointer-initialization: pointer size is too small"); + + void *TgtPteeBase = + getTargetPointeeBaseForCorrespondingPointerInitialization(HstPteeBase, + HstPteeBegin); + + // Store the target pointee base address to the first VoidPtrSize bytes + DP("Initializing corresponding-pointer-initialization source buffer " + "for " DPxMOD ", with pointee base " DPxMOD "\n", + DPxPTR(HstPtr), DPxPTR(TgtPteeBase)); + std::memcpy(Buffer, &TgtPteeBase, VoidPtrSize); + if (HstPtrSize <= VoidPtrSize) + return; + + // For Fortran descriptors, copy the remaining descriptor fields from host + uint64_t HstDescriptorFieldsSize = HstPtrSize - VoidPtrSize; + void *HstDescriptorFieldsAddr = static_cast<char *>(HstPtr) + VoidPtrSize; + DP("Copying %" PRId64 + " bytes of descriptor fields into corresponding-pointer-initialization " + "buffer at offset %" PRId64 ", from " DPxMOD "\n", + HstDescriptorFieldsSize, VoidPtrSize, DPxPTR(HstDescriptorFieldsAddr)); + std::memcpy(Buffer + VoidPtrSize, HstDescriptorFieldsAddr, + HstDescriptorFieldsSize); + } + + /// Helper function to create and initialize a buffer to be used as the source + /// for corresponding-pointer-initialization. + void *createAndInitSourceBufferForCorrespondingPointerInitialization( + void *HstPtr, int64_t HstPtrSize, void *HstPteeBase, void *HstPteeBegin) { + char *Buffer = getOrCreateSourceBufferForSubmitData(AsyncInfo, HstPtrSize); + initBufferForCorrespondingPointerInitialization(Buffer, HstPtr, HstPtrSize, + HstPteeBase, HstPteeBegin); + return Buffer; + } + // TODO: What would be the best value here? Should we make it configurable? // If the size is larger than this threshold, we will allocate and transfer it // immediately instead of packing it. @@ -1417,7 +1624,9 @@ public: int addArg(void *HstPtr, int64_t ArgSize, int64_t ArgOffset, bool IsFirstPrivate, void *&TgtPtr, int TgtArgsIndex, map_var_info_t HstPtrName = nullptr, - const bool AllocImmediately = false) { + const bool AllocImmediately = false, void *HstPteeBase = nullptr, + void *HstPteeBegin = nullptr, + bool IsCorrespondingPointerInit = false) { // If the argument is not first-private, or its size is greater than a // predefined threshold, we will allocate memory and issue the transfer // immediately. @@ -1440,9 +1649,19 @@ public: // If first-private, copy data from host if (IsFirstPrivate) { DP("Submitting firstprivate data to the device.\n"); - int Ret = Device.submitData(TgtPtr, HstPtr, ArgSize, AsyncInfo); + + // The source value used for corresponding-pointer-initialization + // is different vs regular firstprivates. + void *DataSource = + IsCorrespondingPointerInit + ? createAndInitSourceBufferForCorrespondingPointerInitialization( + HstPtr, ArgSize, HstPteeBase, HstPteeBegin) + : HstPtr; + int Ret = Device.submitData(TgtPtr, DataSource, ArgSize, AsyncInfo); if (Ret != OFFLOAD_SUCCESS) { - DP("Copying data to device failed, failed.\n"); + DP("Copying %s data to device failed.\n", + IsCorrespondingPointerInit ? "corresponding-pointer-initialization" + : "firstprivate"); return OFFLOAD_FAIL; } } @@ -1488,8 +1707,10 @@ public: } } - FirstPrivateArgInfo.emplace_back(TgtArgsIndex, HstPtr, ArgSize, - StartAlignment, Padding, HstPtrName); + FirstPrivateArgInfo.emplace_back( + TgtArgsIndex, HstPtr, ArgSize, StartAlignment, Padding, HstPtrName, + HstPteeBase, HstPteeBegin, IsCorrespondingPointerInit); + FirstPrivateArgSize += Padding + ArgSize; } @@ -1508,7 +1729,13 @@ public: for (FirstPrivateArgInfoTy &Info : FirstPrivateArgInfo) { // First pad the pointer as we (have to) pad it on the device too. Itr = std::next(Itr, Info.Padding); - std::copy(Info.HstPtrBegin, Info.HstPtrEnd, Itr); + + if (Info.IsCorrespondingPointerInit) + initBufferForCorrespondingPointerInitialization( + &*Itr, Info.HstPtrBegin, Info.Size, Info.HstPteeBase, + Info.HstPteeBegin); + else + std::copy(Info.HstPtrBegin, Info.HstPtrEnd, Itr); Itr = std::next(Itr, Info.Size); } // Allocate target memory @@ -1664,8 +1891,40 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr, TgtPtrBegin = HstPtrBase; TgtBaseOffset = 0; } else if (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE) { + // For cases like: + // ``` + // int *p = ...; + // #pragma omp target map(p[0:10]) + // ``` + // `p` is predetermined firstprivate on the target construct, and the + // method to determine the initial value of the private copy on the + // device is called "corresponding-pointer-initialization". + // + // Such firstprivate pointers that need + // corresponding-pointer-initialization are represented using the + // `PRIVATE | ATTACH` map-types, in contrast to regular firstprivate + // entries, which use `PRIVATE | TO`. The structure of these + // `PRIVATE | ATTACH` entries is the same as the non-private + // `ATTACH` entries used to represent pointer-attachments, i.e.: + // ``` + // &hst_ptr_base/begin, &hst_ptee_begin, sizeof(hst_ptr) + // ``` + const bool IsAttach = (ArgTypes[I] & OMP_TGT_MAPTYPE_ATTACH); + void *HstPteeBase = nullptr; + void *HstPteeBegin = nullptr; + if (IsAttach) { + // For corresponding-pointer-initialization, Args[I] is HstPteeBegin, + // and ArgBases[I] is both HstPtrBase/HstPtrBegin. + HstPteeBase = *reinterpret_cast<void **>(HstPtrBase); + HstPteeBegin = Args[I]; + HstPtrBegin = ArgBases[I]; + } TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin; - const bool IsFirstPrivate = (ArgTypes[I] & OMP_TGT_MAPTYPE_TO); + // Corresponding-pointer-initialization is a special case of firstprivate, + // since it also involves initializing the private pointer. + const bool IsFirstPrivate = + (ArgTypes[I] & OMP_TGT_MAPTYPE_TO) || IsAttach; + // If there is a next argument and it depends on the current one, we need // to allocate the private memory immediately. If this is not the case, // then the argument can be marked for optimization and packed with the @@ -1674,9 +1933,11 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr, (I < ArgNum - 1 && (ArgTypes[I + 1] & OMP_TGT_MAPTYPE_MEMBER_OF)); Ret = PrivateArgumentManager.addArg( HstPtrBegin, ArgSizes[I], TgtBaseOffset, IsFirstPrivate, TgtPtrBegin, - TgtArgs.size(), HstPtrName, AllocImmediately); + /*TgtArgsIndex=*/TgtArgs.size(), HstPtrName, AllocImmediately, + HstPteeBase, HstPteeBegin, /*IsCorrespondingPointerInit=*/IsAttach); if (Ret != OFFLOAD_SUCCESS) { - REPORT("Failed to process %sprivate argument " DPxMOD "\n", + REPORT("Failed to process %s%sprivate argument " DPxMOD "\n", + IsAttach ? "corresponding-pointer-initialization " : "", (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtrBegin)); return OFFLOAD_FAIL; } diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp index c26cfe961aa0..8d2f9755e035 100644 --- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp +++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp @@ -99,7 +99,7 @@ namespace hsa_utils { /// Iterate elements using an HSA iterate function. Do not use this function /// directly but the specialized ones below instead. template <typename ElemTy, typename IterFuncTy, typename CallbackTy> -hsa_status_t iterate(IterFuncTy Func, CallbackTy Cb) { +static hsa_status_t iterate(IterFuncTy Func, CallbackTy Cb) { auto L = [](ElemTy Elem, void *Data) -> hsa_status_t { CallbackTy *Unwrapped = static_cast<CallbackTy *>(Data); return (*Unwrapped)(Elem); @@ -111,7 +111,8 @@ hsa_status_t iterate(IterFuncTy Func, CallbackTy Cb) { /// use this function directly but the specialized ones below instead. template <typename ElemTy, typename IterFuncTy, typename IterFuncArgTy, typename CallbackTy> -hsa_status_t iterate(IterFuncTy Func, IterFuncArgTy FuncArg, CallbackTy Cb) { +static hsa_status_t iterate(IterFuncTy Func, IterFuncArgTy FuncArg, + CallbackTy Cb) { auto L = [](ElemTy Elem, void *Data) -> hsa_status_t { CallbackTy *Unwrapped = static_cast<CallbackTy *>(Data); return (*Unwrapped)(Elem); @@ -123,7 +124,8 @@ hsa_status_t iterate(IterFuncTy Func, IterFuncArgTy FuncArg, CallbackTy Cb) { /// use this function directly but the specialized ones below instead. template <typename Elem1Ty, typename Elem2Ty, typename IterFuncTy, typename IterFuncArgTy, typename CallbackTy> -hsa_status_t iterate(IterFuncTy Func, IterFuncArgTy FuncArg, CallbackTy Cb) { +static hsa_status_t iterate(IterFuncTy Func, IterFuncArgTy FuncArg, + CallbackTy Cb) { auto L = [](Elem1Ty Elem1, Elem2Ty Elem2, void *Data) -> hsa_status_t { CallbackTy *Unwrapped = static_cast<CallbackTy *>(Data); return (*Unwrapped)(Elem1, Elem2); @@ -132,21 +134,21 @@ hsa_status_t iterate(IterFuncTy Func, IterFuncArgTy FuncArg, CallbackTy Cb) { } /// Iterate agents. -template <typename CallbackTy> Error iterateAgents(CallbackTy Callback) { +template <typename CallbackTy> static Error iterateAgents(CallbackTy Callback) { hsa_status_t Status = iterate<hsa_agent_t>(hsa_iterate_agents, Callback); return Plugin::check(Status, "error in hsa_iterate_agents: %s"); } /// Iterate ISAs of an agent. template <typename CallbackTy> -Error iterateAgentISAs(hsa_agent_t Agent, CallbackTy Cb) { +static Error iterateAgentISAs(hsa_agent_t Agent, CallbackTy Cb) { hsa_status_t Status = iterate<hsa_isa_t>(hsa_agent_iterate_isas, Agent, Cb); return Plugin::check(Status, "error in hsa_agent_iterate_isas: %s"); } /// Iterate memory pools of an agent. template <typename CallbackTy> -Error iterateAgentMemoryPools(hsa_agent_t Agent, CallbackTy Cb) { +static Error iterateAgentMemoryPools(hsa_agent_t Agent, CallbackTy Cb) { hsa_status_t Status = iterate<hsa_amd_memory_pool_t>( hsa_amd_agent_iterate_memory_pools, Agent, Cb); return Plugin::check(Status, @@ -155,10 +157,12 @@ Error iterateAgentMemoryPools(hsa_agent_t Agent, CallbackTy Cb) { /// Dispatches an asynchronous memory copy. /// Enables different SDMA engines for the dispatch in a round-robin fashion. -Error asyncMemCopy(bool UseMultipleSdmaEngines, void *Dst, hsa_agent_t DstAgent, - const void *Src, hsa_agent_t SrcAgent, size_t Size, - uint32_t NumDepSignals, const hsa_signal_t *DepSignals, - hsa_signal_t CompletionSignal) { +static Error asyncMemCopy(bool UseMultipleSdmaEngines, void *Dst, + hsa_agent_t DstAgent, const void *Src, + hsa_agent_t SrcAgent, size_t Size, + uint32_t NumDepSignals, + const hsa_signal_t *DepSignals, + hsa_signal_t CompletionSignal) { if (!UseMultipleSdmaEngines) { hsa_status_t S = hsa_amd_memory_async_copy(Dst, DstAgent, Src, SrcAgent, Size, @@ -193,8 +197,8 @@ Error asyncMemCopy(bool UseMultipleSdmaEngines, void *Dst, hsa_agent_t DstAgent, #endif } -Error getTargetTripleAndFeatures(hsa_agent_t Agent, - SmallVector<SmallString<32>> &Targets) { +static Error getTargetTripleAndFeatures(hsa_agent_t Agent, + SmallVector<SmallString<32>> &Targets) { auto Err = hsa_utils::iterateAgentISAs(Agent, [&](hsa_isa_t ISA) { uint32_t Length; hsa_status_t Status; @@ -419,7 +423,11 @@ struct AMDGPUMemoryManagerTy : public DeviceAllocatorTy { assert(MemoryManager && "Invalid memory manager"); assert(PtrStorage && "Invalid pointer storage"); - *PtrStorage = MemoryManager->allocate(Size, nullptr); + auto PtrStorageOrErr = MemoryManager->allocate(Size, nullptr); + if (!PtrStorageOrErr) + return PtrStorageOrErr.takeError(); + + *PtrStorage = *PtrStorageOrErr; if (Size && *PtrStorage == nullptr) return Plugin::error(ErrorCode::OUT_OF_RESOURCES, "failure to allocate from AMDGPU memory manager"); @@ -439,15 +447,12 @@ struct AMDGPUMemoryManagerTy : public DeviceAllocatorTy { private: /// Allocation callback that will be called once the memory manager does not /// have more previously allocated buffers. - void *allocate(size_t Size, void *HstPtr, TargetAllocTy Kind) override; + Expected<void *> allocate(size_t Size, void *HstPtr, + TargetAllocTy Kind) override; /// Deallocation callback that will be called by the memory manager. - int free(void *TgtPtr, TargetAllocTy Kind) override { - if (auto Err = MemoryPool->deallocate(TgtPtr)) { - consumeError(std::move(Err)); - return OFFLOAD_FAIL; - } - return OFFLOAD_SUCCESS; + Error free(void *TgtPtr, TargetAllocTy Kind) override { + return MemoryPool->deallocate(TgtPtr); } /// The underlying plugin that owns this memory manager. @@ -464,8 +469,8 @@ private: struct AMDGPUDeviceImageTy : public DeviceImageTy { /// Create the AMDGPU image with the id and the target image pointer. AMDGPUDeviceImageTy(int32_t ImageId, GenericDeviceTy &Device, - const __tgt_device_image *TgtImage) - : DeviceImageTy(ImageId, Device, TgtImage) {} + std::unique_ptr<MemoryBuffer> &&TgtImage) + : DeviceImageTy(ImageId, Device, std::move(TgtImage)) {} /// Prepare and load the executable corresponding to the image. Error loadExecutable(const AMDGPUDeviceTy &Device); @@ -1222,7 +1227,7 @@ private: assert(Args->Dst && "Invalid destination buffer"); assert(Args->Src && "Invalid source buffer"); - auto BasePtr = Args->Dst; + auto *BasePtr = Args->Dst; for (size_t I = 0; I < Args->NumTimes; I++) { std::memcpy(BasePtr, Args->Src, Args->Size); BasePtr = reinterpret_cast<uint8_t *>(BasePtr) + Args->Size; @@ -2160,7 +2165,12 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { AMDGPUDeviceImageTy &AMDImage = static_cast<AMDGPUDeviceImageTy &>(*Image); // Unload the executable of the image. - return AMDImage.unloadExecutable(); + if (auto Err = AMDImage.unloadExecutable()) + return Err; + + // Destroy the associated memory and invalidate the object. + Plugin.free(Image); + return Error::success(); } /// Deinitialize the device and release its resources. @@ -2183,18 +2193,12 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { virtual Error callGlobalConstructors(GenericPluginTy &Plugin, DeviceImageTy &Image) override { - GenericGlobalHandlerTy &Handler = Plugin.getGlobalHandler(); - if (Handler.isSymbolInImage(*this, Image, "amdgcn.device.fini")) - Image.setPendingGlobalDtors(); - return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/true); } virtual Error callGlobalDestructors(GenericPluginTy &Plugin, DeviceImageTy &Image) override { - if (Image.hasPendingGlobalDtors()) - return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/false); - return Plugin::success(); + return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/false); } uint64_t getStreamBusyWaitMicroseconds() const { return OMPX_StreamBusyWait; } @@ -2321,11 +2325,12 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { } /// Load the binary image into the device and allocate an image object. - Expected<DeviceImageTy *> loadBinaryImpl(const __tgt_device_image *TgtImage, - int32_t ImageId) override { + Expected<DeviceImageTy *> + loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, + int32_t ImageId) override { // Allocate and initialize the image object. AMDGPUDeviceImageTy *AMDImage = Plugin.allocate<AMDGPUDeviceImageTy>(); - new (AMDImage) AMDGPUDeviceImageTy(ImageId, *this, TgtImage); + new (AMDImage) AMDGPUDeviceImageTy(ImageId, *this, std::move(TgtImage)); // Load the HSA executable. if (Error Err = AMDImage->loadExecutable(*this)) @@ -2335,18 +2340,17 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { } /// Allocate memory on the device or related to the device. - void *allocate(size_t Size, void *, TargetAllocTy Kind) override; + Expected<void *> allocate(size_t Size, void *, TargetAllocTy Kind) override; /// Deallocate memory on the device or related to the device. - int free(void *TgtPtr, TargetAllocTy Kind) override { + Error free(void *TgtPtr, TargetAllocTy Kind) override { if (TgtPtr == nullptr) - return OFFLOAD_SUCCESS; + return Plugin::success(); AMDGPUMemoryPoolTy *MemoryPool = nullptr; switch (Kind) { case TARGET_ALLOC_DEFAULT: case TARGET_ALLOC_DEVICE: - case TARGET_ALLOC_DEVICE_NON_BLOCKING: MemoryPool = CoarseGrainedMemoryPools[0]; break; case TARGET_ALLOC_HOST: @@ -2357,17 +2361,14 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { break; } - if (!MemoryPool) { - REPORT("No memory pool for the specified allocation kind\n"); - return OFFLOAD_FAIL; - } + if (!MemoryPool) + return Plugin::error(ErrorCode::OUT_OF_RESOURCES, + "no memory pool for the specified allocation kind"); - if (Error Err = MemoryPool->deallocate(TgtPtr)) { - REPORT("%s\n", toString(std::move(Err)).data()); - return OFFLOAD_FAIL; - } + if (auto Err = MemoryPool->deallocate(TgtPtr)) + return Err; - return OFFLOAD_SUCCESS; + return Plugin::success(); } /// Synchronize current thread with the pending operations on the async info. @@ -2673,11 +2674,10 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { // hsa_amd_memory_fill doesn't signal completion using a signal, so use // the existing host callback logic to handle that instead return Stream->pushHostCallback(Fill, Args); - } else { - // If there is no pending work, do the fill synchronously - auto Status = hsa_amd_memory_fill(TgtPtr, Pattern, Size / 4); - return Plugin::check(Status, "error in hsa_amd_memory_fill: %s\n"); } + // If there is no pending work, do the fill synchronously + auto Status = hsa_amd_memory_fill(TgtPtr, Pattern, Size / 4); + return Plugin::check(Status, "error in hsa_amd_memory_fill: %s\n"); } // Slow case; allocate an appropriate memory size and enqueue copies @@ -2712,6 +2712,37 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { return Plugin::success(); } + interop_spec_t selectInteropPreference(int32_t InteropType, + int32_t NumPrefers, + interop_spec_t *Prefers) override { + // TODO: update once targetsync is supported + if (InteropType == kmp_interop_type_target) + return interop_spec_t{tgt_fr_hsa, {false, 0}, 0}; + return interop_spec_t{tgt_fr_none, {false, 0}, 0}; + } + + Expected<omp_interop_val_t *> + createInterop(int32_t InteropType, interop_spec_t &InteropSpec) override { + auto *Ret = new omp_interop_val_t( + DeviceId, static_cast<kmp_interop_type_t>(InteropType)); + Ret->fr_id = tgt_fr_hsa; + Ret->vendor_id = omp_vendor_amd; + + // TODO: implement targetsync support + + Ret->device_info.Platform = nullptr; + Ret->device_info.Device = reinterpret_cast<void *>(Agent.handle); + Ret->device_info.Context = nullptr; + + return Ret; + } + + Error releaseInterop(omp_interop_val_t *Interop) override { + if (Interop) + delete Interop; + return Plugin::success(); + } + Error enqueueHostCallImpl(void (*Callback)(void *), void *UserData, AsyncInfoWrapperTy &AsyncInfo) override { AMDGPUStreamTy *Stream = nullptr; @@ -2759,7 +2790,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { } Expected<bool> hasPendingWorkImpl(AsyncInfoWrapperTy &AsyncInfo) override { - auto Stream = AsyncInfo.getQueueAs<AMDGPUStreamTy *>(); + auto *Stream = AsyncInfo.getQueueAs<AMDGPUStreamTy *>(); if (!Stream) return false; @@ -2772,7 +2803,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { Expected<bool> isEventCompleteImpl(void *EventPtr, AsyncInfoWrapperTy &AsyncInfo) override { AMDGPUEventTy *Event = reinterpret_cast<AMDGPUEventTy *>(EventPtr); - auto Stream = AsyncInfo.getQueueAs<AMDGPUStreamTy *>(); + auto *Stream = AsyncInfo.getQueueAs<AMDGPUStreamTy *>(); return Stream && Stream->isEventComplete(*Event); } @@ -2829,7 +2860,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { hsa_device_type_t DevType; Status = getDeviceAttrRaw(HSA_AGENT_INFO_DEVICE, DevType); if (Status == HSA_STATUS_SUCCESS) { - switch (DevType) { + switch (static_cast<int>(DevType)) { case HSA_DEVICE_TYPE_CPU: TmpCharPtr = "CPU"; break; @@ -3105,7 +3136,7 @@ private: // Perform a quick check for the named kernel in the image. The kernel // should be created by the 'amdgpu-lower-ctor-dtor' pass. GenericGlobalHandlerTy &Handler = Plugin.getGlobalHandler(); - if (IsCtor && !Handler.isSymbolInImage(*this, Image, KernelName)) + if (!Handler.isSymbolInImage(*this, Image, KernelName)) return Plugin::success(); // Allocate and construct the AMDGPU kernel. @@ -3656,11 +3687,6 @@ Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice, KernelArgsTy &KernelArgs, KernelLaunchParamsTy LaunchParams, AsyncInfoWrapperTy &AsyncInfoWrapper) const { - if (ArgsSize != LaunchParams.Size && - ArgsSize > LaunchParams.Size + getImplicitArgsSize()) - return Plugin::error(ErrorCode::INVALID_ARGUMENT, - "invalid kernel arguments size"); - AMDGPUPluginTy &AMDGPUPlugin = static_cast<AMDGPUPluginTy &>(GenericDevice.Plugin); AMDHostDeviceTy &HostDevice = AMDGPUPlugin.getHostDevice(); @@ -3746,8 +3772,8 @@ Error AMDGPUKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice, return Plugin::success(); // General Info - auto NumGroups = NumBlocks; - auto ThreadsPerGroup = NumThreads; + auto *NumGroups = NumBlocks; + auto *ThreadsPerGroup = NumThreads; // Kernel Arguments Info auto ArgNum = KernelArgs.NumArgs; @@ -3811,14 +3837,13 @@ static Error Plugin::check(int32_t Code, const char *ErrFmt, ArgsTy... Args) { return Plugin::error(OffloadErrCode, ErrFmt, Args..., Desc); } -void *AMDGPUMemoryManagerTy::allocate(size_t Size, void *HstPtr, - TargetAllocTy Kind) { +Expected<void *> AMDGPUMemoryManagerTy::allocate(size_t Size, void *HstPtr, + TargetAllocTy Kind) { // Allocate memory from the pool. void *Ptr = nullptr; - if (auto Err = MemoryPool->allocate(Size, &Ptr)) { - consumeError(std::move(Err)); - return nullptr; - } + if (auto Err = MemoryPool->allocate(Size, &Ptr)) + return std::move(Err); + assert(Ptr && "Invalid pointer"); // Get a list of agents that can access this memory pool. @@ -3828,14 +3853,13 @@ void *AMDGPUMemoryManagerTy::allocate(size_t Size, void *HstPtr, [&](hsa_agent_t Agent) { return MemoryPool->canAccess(Agent); }); // Allow all valid kernel agents to access the allocation. - if (auto Err = MemoryPool->enableAccess(Ptr, Size, Agents)) { - REPORT("%s\n", toString(std::move(Err)).data()); - return nullptr; - } + if (auto Err = MemoryPool->enableAccess(Ptr, Size, Agents)) + return std::move(Err); return Ptr; } -void *AMDGPUDeviceTy::allocate(size_t Size, void *, TargetAllocTy Kind) { +Expected<void *> AMDGPUDeviceTy::allocate(size_t Size, void *, + TargetAllocTy Kind) { if (Size == 0) return nullptr; @@ -3844,7 +3868,6 @@ void *AMDGPUDeviceTy::allocate(size_t Size, void *, TargetAllocTy Kind) { switch (Kind) { case TARGET_ALLOC_DEFAULT: case TARGET_ALLOC_DEVICE: - case TARGET_ALLOC_DEVICE_NON_BLOCKING: MemoryPool = CoarseGrainedMemoryPools[0]; break; case TARGET_ALLOC_HOST: @@ -3855,17 +3878,14 @@ void *AMDGPUDeviceTy::allocate(size_t Size, void *, TargetAllocTy Kind) { break; } - if (!MemoryPool) { - REPORT("No memory pool for the specified allocation kind\n"); - return nullptr; - } + if (!MemoryPool) + return Plugin::error(ErrorCode::UNSUPPORTED, + "no memory pool for the specified allocation kind"); // Allocate from the corresponding memory pool. void *Alloc = nullptr; - if (Error Err = MemoryPool->allocate(Size, &Alloc)) { - REPORT("%s\n", toString(std::move(Err)).data()); - return nullptr; - } + if (auto Err = MemoryPool->allocate(Size, &Alloc)) + return std::move(Err); if (Alloc) { // Get a list of agents that can access this memory pool. Inherently @@ -3878,10 +3898,8 @@ void *AMDGPUDeviceTy::allocate(size_t Size, void *, TargetAllocTy Kind) { }); // Enable all valid kernel agents to access the buffer. - if (auto Err = MemoryPool->enableAccess(Alloc, Size, Agents)) { - REPORT("%s\n", toString(std::move(Err)).data()); - return nullptr; - } + if (auto Err = MemoryPool->enableAccess(Alloc, Size, Agents)) + return std::move(Err); } return Alloc; diff --git a/offload/plugins-nextgen/common/include/ErrorReporting.h b/offload/plugins-nextgen/common/include/ErrorReporting.h index 2ad0f2b7dd6c..68d82cbea0f3 100644 --- a/offload/plugins-nextgen/common/include/ErrorReporting.h +++ b/offload/plugins-nextgen/common/include/ErrorReporting.h @@ -61,7 +61,6 @@ class ErrorReporter { /// Return a nice name for an TargetAllocTy. static StringRef getAllocTyName(TargetAllocTy Kind) { switch (Kind) { - case TARGET_ALLOC_DEVICE_NON_BLOCKING: case TARGET_ALLOC_DEFAULT: case TARGET_ALLOC_DEVICE: return "device memory"; diff --git a/offload/plugins-nextgen/common/include/JIT.h b/offload/plugins-nextgen/common/include/JIT.h index d62516d20764..b4e3712d9c98 100644 --- a/offload/plugins-nextgen/common/include/JIT.h +++ b/offload/plugins-nextgen/common/include/JIT.h @@ -51,27 +51,22 @@ struct JITEngine { /// Run jit compilation if \p Image is a bitcode image, otherwise simply /// return \p Image. It is expected to return a memory buffer containing the /// generated device image that could be loaded to the device directly. - Expected<const __tgt_device_image *> - process(const __tgt_device_image &Image, - target::plugin::GenericDeviceTy &Device); - - /// Remove \p Image from the jit engine's cache - void erase(const __tgt_device_image &Image, - target::plugin::GenericDeviceTy &Device); + Expected<std::unique_ptr<MemoryBuffer>> + process(StringRef Image, target::plugin::GenericDeviceTy &Device); private: /// Compile the bitcode image \p Image and generate the binary image that can /// be loaded to the target device of the triple \p Triple architecture \p /// MCpu. \p PostProcessing will be called after codegen to handle cases such /// as assembler as an external tool. - Expected<const __tgt_device_image *> - compile(const __tgt_device_image &Image, const std::string &ComputeUnitKind, + Expected<std::unique_ptr<MemoryBuffer>> + compile(StringRef Image, const std::string &ComputeUnitKind, PostProcessingFn PostProcessing); /// Create or retrieve the object image file from the file system or via /// compilation of the \p Image. Expected<std::unique_ptr<MemoryBuffer>> - getOrCreateObjFile(const __tgt_device_image &Image, LLVMContext &Ctx, + getOrCreateObjFile(StringRef Image, LLVMContext &Ctx, const std::string &ComputeUnitKind); /// Run backend, which contains optimization and code generation. @@ -92,14 +87,6 @@ private: struct ComputeUnitInfo { /// LLVM Context in which the modules will be constructed. LLVMContext Context; - - /// A map of embedded IR images to the buffer used to store JITed code - DenseMap<const __tgt_device_image *, std::unique_ptr<MemoryBuffer>> - JITImages; - - /// A map of embedded IR images to JITed images. - DenseMap<const __tgt_device_image *, std::unique_ptr<__tgt_device_image>> - TgtImageMap; }; /// Map from (march) "CPUs" (e.g., sm_80, or gfx90a), which we call compute diff --git a/offload/plugins-nextgen/common/include/MemoryManager.h b/offload/plugins-nextgen/common/include/MemoryManager.h index a4f6e628c403..8f6c1adcdaa5 100644 --- a/offload/plugins-nextgen/common/include/MemoryManager.h +++ b/offload/plugins-nextgen/common/include/MemoryManager.h @@ -25,6 +25,10 @@ #include "Shared/Utils.h" #include "omptarget.h" +#include "llvm/Support/Error.h" + +namespace llvm { + /// Base class of per-device allocator. class DeviceAllocatorTy { public: @@ -32,11 +36,13 @@ public: /// Allocate a memory of size \p Size . \p HstPtr is used to assist the /// allocation. - virtual void *allocate(size_t Size, void *HstPtr, - TargetAllocTy Kind = TARGET_ALLOC_DEFAULT) = 0; + virtual Expected<void *> + allocate(size_t Size, void *HstPtr, + TargetAllocTy Kind = TARGET_ALLOC_DEFAULT) = 0; /// Delete the pointer \p TgtPtr on the device - virtual int free(void *TgtPtr, TargetAllocTy Kind = TARGET_ALLOC_DEFAULT) = 0; + virtual Error free(void *TgtPtr, + TargetAllocTy Kind = TARGET_ALLOC_DEFAULT) = 0; }; /// Class of memory manager. The memory manager is per-device by using @@ -134,17 +140,17 @@ class MemoryManagerTy { size_t SizeThreshold = 1U << 13; /// Request memory from target device - void *allocateOnDevice(size_t Size, void *HstPtr) const { + Expected<void *> allocateOnDevice(size_t Size, void *HstPtr) const { return DeviceAllocator.allocate(Size, HstPtr, TARGET_ALLOC_DEVICE); } /// Deallocate data on device - int deleteOnDevice(void *Ptr) const { return DeviceAllocator.free(Ptr); } + Error deleteOnDevice(void *Ptr) const { return DeviceAllocator.free(Ptr); } /// This function is called when it tries to allocate memory on device but the /// device returns out of memory. It will first free all memory in the /// FreeList and try to allocate again. - void *freeAndAllocate(size_t Size, void *HstPtr) { + Expected<void *> freeAndAllocate(size_t Size, void *HstPtr) { std::vector<void *> RemoveList; // Deallocate all memory in FreeList @@ -154,7 +160,8 @@ class MemoryManagerTy { if (List.empty()) continue; for (const NodeTy &N : List) { - deleteOnDevice(N.Ptr); + if (auto Err = deleteOnDevice(N.Ptr)) + return Err; RemoveList.push_back(N.Ptr); } FreeLists[I].clear(); @@ -175,14 +182,22 @@ class MemoryManagerTy { /// allocate directly on the device. If a \p nullptr is returned, it might /// be because the device is OOM. In that case, it will free all unused /// memory and then try again. - void *allocateOrFreeAndAllocateOnDevice(size_t Size, void *HstPtr) { - void *TgtPtr = allocateOnDevice(Size, HstPtr); + Expected<void *> allocateOrFreeAndAllocateOnDevice(size_t Size, + void *HstPtr) { + auto TgtPtrOrErr = allocateOnDevice(Size, HstPtr); + if (!TgtPtrOrErr) + return TgtPtrOrErr.takeError(); + + void *TgtPtr = *TgtPtrOrErr; // We cannot get memory from the device. It might be due to OOM. Let's // free all memory in FreeLists and try again. if (TgtPtr == nullptr) { DP("Failed to get memory on device. Free all memory in FreeLists and " "try again.\n"); - TgtPtr = freeAndAllocate(Size, HstPtr); + TgtPtrOrErr = freeAndAllocate(Size, HstPtr); + if (!TgtPtrOrErr) + return TgtPtrOrErr.takeError(); + TgtPtr = *TgtPtrOrErr; } if (TgtPtr == nullptr) @@ -204,16 +219,17 @@ public: /// Destructor ~MemoryManagerTy() { - for (auto Itr = PtrToNodeTable.begin(); Itr != PtrToNodeTable.end(); - ++Itr) { - assert(Itr->second.Ptr && "nullptr in map table"); - deleteOnDevice(Itr->second.Ptr); + for (auto &PtrToNode : PtrToNodeTable) { + assert(PtrToNode.second.Ptr && "nullptr in map table"); + if (auto Err = deleteOnDevice(PtrToNode.second.Ptr)) + REPORT("Failure to delete memory: %s\n", + toString(std::move(Err)).data()); } } /// Allocate memory of size \p Size from target device. \p HstPtr is used to /// assist the allocation. - void *allocate(size_t Size, void *HstPtr) { + Expected<void *> allocate(size_t Size, void *HstPtr) { // If the size is zero, we will not bother the target device. Just return // nullptr directly. if (Size == 0) @@ -228,11 +244,14 @@ public: DP("%zu is greater than the threshold %zu. Allocate it directly from " "device\n", Size, SizeThreshold); - void *TgtPtr = allocateOrFreeAndAllocateOnDevice(Size, HstPtr); + auto TgtPtrOrErr = allocateOrFreeAndAllocateOnDevice(Size, HstPtr); + if (!TgtPtrOrErr) + return TgtPtrOrErr.takeError(); - DP("Got target pointer " DPxMOD ". Return directly.\n", DPxPTR(TgtPtr)); + DP("Got target pointer " DPxMOD ". Return directly.\n", + DPxPTR(*TgtPtrOrErr)); - return TgtPtr; + return *TgtPtrOrErr; } NodeTy *NodePtr = nullptr; @@ -260,8 +279,11 @@ public: if (NodePtr == nullptr) { DP("Cannot find a node in the FreeLists. Allocate on device.\n"); // Allocate one on device - void *TgtPtr = allocateOrFreeAndAllocateOnDevice(Size, HstPtr); + auto TgtPtrOrErr = allocateOrFreeAndAllocateOnDevice(Size, HstPtr); + if (!TgtPtrOrErr) + return TgtPtrOrErr.takeError(); + void *TgtPtr = *TgtPtrOrErr; if (TgtPtr == nullptr) return nullptr; @@ -282,7 +304,7 @@ public: } /// Deallocate memory pointed by \p TgtPtr - int free(void *TgtPtr) { + Error free(void *TgtPtr) { DP("MemoryManagerTy::free: target memory " DPxMOD ".\n", DPxPTR(TgtPtr)); NodeTy *P = nullptr; @@ -314,7 +336,7 @@ public: FreeLists[B].insert(*P); } - return OFFLOAD_SUCCESS; + return Error::success(); } /// Get the size threshold from the environment variable @@ -344,4 +366,6 @@ public: constexpr const size_t MemoryManagerTy::BucketSize[]; constexpr const int MemoryManagerTy::NumBuckets; +} // namespace llvm + #endif // LLVM_OPENMP_LIBOMPTARGET_PLUGINS_COMMON_MEMORYMANAGER_H diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h index 6ff3ef8cda17..5620437716b3 100644 --- a/offload/plugins-nextgen/common/include/PluginInterface.h +++ b/offload/plugins-nextgen/common/include/PluginInterface.h @@ -193,7 +193,7 @@ struct InfoTreeNode { InfoTreeNode() : InfoTreeNode("", std::monostate{}, "") {} InfoTreeNode(std::string Key, VariantType Value, std::string Units) - : Key(Key), Value(Value), Units(Units) {} + : Key(std::move(Key)), Value(Value), Units(std::move(Units)) {} /// Add a new info entry as a child of this node. The entry requires at least /// a key string in \p Key. The value in \p Value is optional and can be any @@ -202,7 +202,7 @@ struct InfoTreeNode { /// use that value for an appropriate olGetDeviceInfo query template <typename T = std::monostate> InfoTreeNode *add(std::string Key, T Value = T(), - const std::string &Units = std::string(), + std::string Units = std::string(), std::optional<DeviceInfo> DeviceInfoKey = std::nullopt) { assert(!Key.empty() && "Invalid info key"); @@ -217,7 +217,8 @@ struct InfoTreeNode { else ValueVariant = std::string{Value}; - auto Ptr = &Children->emplace_back(Key, ValueVariant, Units); + auto Ptr = + &Children->emplace_back(std::move(Key), ValueVariant, std::move(Units)); if (DeviceInfoKey) DeviceInfoMap[*DeviceInfoKey] = Children->size() - 1; @@ -306,26 +307,18 @@ class DeviceImageTy { /// not unique between different device; they may overlap. int32_t ImageId; - /// The pointer to the raw __tgt_device_image. - const __tgt_device_image *TgtImage; - const __tgt_device_image *TgtImageBitcode; + /// The managed image data. + std::unique_ptr<MemoryBuffer> Image; /// Reference to the device this image is loaded on. GenericDeviceTy &Device; - /// If this image has any global destructors that much be called. - /// FIXME: This is only required because we currently have no invariants - /// towards the lifetime of the underlying image. We should either copy - /// the image into memory locally or erase the pointers after init. - bool PendingGlobalDtors; - public: + virtual ~DeviceImageTy() = default; + DeviceImageTy(int32_t Id, GenericDeviceTy &Device, - const __tgt_device_image *Image) - : ImageId(Id), TgtImage(Image), TgtImageBitcode(nullptr), Device(Device), - PendingGlobalDtors(false) { - assert(TgtImage && "Invalid target image"); - } + std::unique_ptr<MemoryBuffer> &&Image) + : ImageId(Id), Image(std::move(Image)), Device(Device) {} /// Get the image identifier within the device. int32_t getId() const { return ImageId; } @@ -333,33 +326,17 @@ public: /// Get the device that this image is loaded onto. GenericDeviceTy &getDevice() const { return Device; } - /// Get the pointer to the raw __tgt_device_image. - const __tgt_device_image *getTgtImage() const { return TgtImage; } - - void setTgtImageBitcode(const __tgt_device_image *TgtImageBitcode) { - this->TgtImageBitcode = TgtImageBitcode; - } - - const __tgt_device_image *getTgtImageBitcode() const { - return TgtImageBitcode; - } - /// Get the image starting address. - void *getStart() const { return TgtImage->ImageStart; } + const void *getStart() const { return Image->getBufferStart(); } /// Get the image size. - size_t getSize() const { - return utils::getPtrDiff(TgtImage->ImageEnd, TgtImage->ImageStart); - } + size_t getSize() const { return Image->getBufferSize(); } /// Get a memory buffer reference to the whole image. MemoryBufferRef getMemoryBuffer() const { return MemoryBufferRef(StringRef((const char *)getStart(), getSize()), "Image"); } - /// Accessors to the boolean value - bool setPendingGlobalDtors() { return PendingGlobalDtors = true; } - bool hasPendingGlobalDtors() const { return PendingGlobalDtors; } }; /// Class implementing common functionalities of offload kernels. Each plugin @@ -831,9 +808,9 @@ struct GenericDeviceTy : public DeviceAllocatorTy { /// Load the binary image into the device and return the target table. Expected<DeviceImageTy *> loadBinary(GenericPluginTy &Plugin, - const __tgt_device_image *TgtImage); + StringRef TgtImage); virtual Expected<DeviceImageTy *> - loadBinaryImpl(const __tgt_device_image *TgtImage, int32_t ImageId) = 0; + loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, int32_t ImageId) = 0; /// Unload a previously loaded Image from the device Error unloadBinary(DeviceImageTy *Image); @@ -1402,10 +1379,10 @@ public: /// Returns non-zero if the \p Image is compatible with the plugin. This /// function does not require the plugin to be initialized before use. - int32_t is_plugin_compatible(__tgt_device_image *Image); + int32_t isPluginCompatible(StringRef Image); /// Returns non-zero if the \p Image is compatible with the device. - int32_t is_device_compatible(int32_t DeviceId, __tgt_device_image *Image); + int32_t isDeviceCompatible(int32_t DeviceId, StringRef Image); /// Returns non-zero if the plugin device has been initialized. int32_t is_device_initialized(int32_t DeviceId) const; diff --git a/offload/plugins-nextgen/common/src/JIT.cpp b/offload/plugins-nextgen/common/src/JIT.cpp index 00720fa2d810..881e27dad384 100644 --- a/offload/plugins-nextgen/common/src/JIT.cpp +++ b/offload/plugins-nextgen/common/src/JIT.cpp @@ -49,13 +49,6 @@ using namespace omp::target; namespace { -bool isImageBitcode(const __tgt_device_image &Image) { - StringRef Binary(reinterpret_cast<const char *>(Image.ImageStart), - utils::getPtrDiff(Image.ImageEnd, Image.ImageStart)); - - return identify_magic(Binary) == file_magic::bitcode; -} - Expected<std::unique_ptr<Module>> createModuleFromMemoryBuffer(std::unique_ptr<MemoryBuffer> &MB, LLVMContext &Context) { @@ -66,12 +59,10 @@ createModuleFromMemoryBuffer(std::unique_ptr<MemoryBuffer> &MB, "failed to create module"); return std::move(Mod); } -Expected<std::unique_ptr<Module>> -createModuleFromImage(const __tgt_device_image &Image, LLVMContext &Context) { - StringRef Data((const char *)Image.ImageStart, - utils::getPtrDiff(Image.ImageEnd, Image.ImageStart)); +Expected<std::unique_ptr<Module>> createModuleFromImage(StringRef Image, + LLVMContext &Context) { std::unique_ptr<MemoryBuffer> MB = MemoryBuffer::getMemBuffer( - Data, /*BufferName=*/"", /*RequiresNullTerminator=*/false); + Image, /*BufferName=*/"", /*RequiresNullTerminator=*/false); return createModuleFromMemoryBuffer(MB, Context); } @@ -189,9 +180,10 @@ Expected<std::unique_ptr<MemoryBuffer>> JITEngine::backend(Module &M, const std::string &ComputeUnitKind, unsigned OptLevel) { - auto RemarksFileOrErr = setupLLVMOptimizationRemarks( - M.getContext(), /*RemarksFilename=*/"", /*RemarksPasses=*/"", - /*RemarksFormat=*/"", /*RemarksWithHotness=*/false); + Expected<LLVMRemarkFileHandle> RemarksFileOrErr = + setupLLVMOptimizationRemarks( + M.getContext(), /*RemarksFilename=*/"", /*RemarksPasses=*/"", + /*RemarksFormat=*/"", /*RemarksWithHotness=*/false); if (Error E = RemarksFileOrErr.takeError()) return std::move(E); if (*RemarksFileOrErr) @@ -238,7 +230,7 @@ JITEngine::backend(Module &M, const std::string &ComputeUnitKind, } Expected<std::unique_ptr<MemoryBuffer>> -JITEngine::getOrCreateObjFile(const __tgt_device_image &Image, LLVMContext &Ctx, +JITEngine::getOrCreateObjFile(StringRef Image, LLVMContext &Ctx, const std::string &ComputeUnitKind) { // Check if the user replaces the module at runtime with a finished object. @@ -277,58 +269,28 @@ JITEngine::getOrCreateObjFile(const __tgt_device_image &Image, LLVMContext &Ctx, return backend(*Mod, ComputeUnitKind, JITOptLevel); } -Expected<const __tgt_device_image *> -JITEngine::compile(const __tgt_device_image &Image, - const std::string &ComputeUnitKind, +Expected<std::unique_ptr<MemoryBuffer>> +JITEngine::compile(StringRef Image, const std::string &ComputeUnitKind, PostProcessingFn PostProcessing) { std::lock_guard<std::mutex> Lock(ComputeUnitMapMutex); - // Check if we JITed this image for the given compute unit kind before. - ComputeUnitInfo &CUI = ComputeUnitMap[ComputeUnitKind]; - if (CUI.TgtImageMap.contains(&Image)) - return CUI.TgtImageMap[&Image].get(); - - auto ObjMBOrErr = getOrCreateObjFile(Image, CUI.Context, ComputeUnitKind); + LLVMContext Ctz; + auto ObjMBOrErr = getOrCreateObjFile(Image, Ctz, ComputeUnitKind); if (!ObjMBOrErr) return ObjMBOrErr.takeError(); - auto ImageMBOrErr = PostProcessing(std::move(*ObjMBOrErr)); - if (!ImageMBOrErr) - return ImageMBOrErr.takeError(); - - CUI.JITImages.insert({&Image, std::move(*ImageMBOrErr)}); - auto &ImageMB = CUI.JITImages[&Image]; - CUI.TgtImageMap.insert({&Image, std::make_unique<__tgt_device_image>()}); - auto &JITedImage = CUI.TgtImageMap[&Image]; - *JITedImage = Image; - JITedImage->ImageStart = const_cast<char *>(ImageMB->getBufferStart()); - JITedImage->ImageEnd = const_cast<char *>(ImageMB->getBufferEnd()); - - return JITedImage.get(); + return PostProcessing(std::move(*ObjMBOrErr)); } -Expected<const __tgt_device_image *> -JITEngine::process(const __tgt_device_image &Image, - target::plugin::GenericDeviceTy &Device) { - const std::string &ComputeUnitKind = Device.getComputeUnitKind(); +Expected<std::unique_ptr<MemoryBuffer>> +JITEngine::process(StringRef Image, target::plugin::GenericDeviceTy &Device) { + assert(identify_magic(Image) == file_magic::bitcode && "Image not LLVM-IR"); + const std::string &ComputeUnitKind = Device.getComputeUnitKind(); PostProcessingFn PostProcessing = [&Device](std::unique_ptr<MemoryBuffer> MB) -> Expected<std::unique_ptr<MemoryBuffer>> { return Device.doJITPostProcessing(std::move(MB)); }; - if (isImageBitcode(Image)) - return compile(Image, ComputeUnitKind, PostProcessing); - - return &Image; -} - -void JITEngine::erase(const __tgt_device_image &Image, - target::plugin::GenericDeviceTy &Device) { - std::lock_guard<std::mutex> Lock(ComputeUnitMapMutex); - const std::string &ComputeUnitKind = Device.getComputeUnitKind(); - ComputeUnitInfo &CUI = ComputeUnitMap[ComputeUnitKind]; - - CUI.TgtImageMap.erase(&Image); - CUI.JITImages.erase(&Image); + return compile(Image, ComputeUnitKind, PostProcessing); } diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp index 36cdd6035e26..15b6b9866e5a 100644 --- a/offload/plugins-nextgen/common/src/PluginInterface.cpp +++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp @@ -73,11 +73,17 @@ private: }; llvm::SmallVector<GlobalEntry> GlobalEntries{}; - void *suggestAddress(uint64_t MaxMemoryAllocation) { + Expected<void *> suggestAddress(uint64_t MaxMemoryAllocation) { // Get a valid pointer address for this system - void *Addr = + auto AddrOrErr = Device->allocate(1024, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT); - Device->free(Addr); + if (!AddrOrErr) + return AddrOrErr.takeError(); + + void *Addr = *AddrOrErr; + if (auto Err = Device->free(Addr)) + return std::move(Err); + // Align Address to MaxMemoryAllocation Addr = (void *)utils::alignPtr((Addr), MaxMemoryAllocation); return Addr; @@ -86,8 +92,12 @@ private: Error preAllocateVAMemory(uint64_t MaxMemoryAllocation, void *VAddr) { size_t ASize = MaxMemoryAllocation; - if (!VAddr && isRecording()) - VAddr = suggestAddress(MaxMemoryAllocation); + if (!VAddr && isRecording()) { + auto VAddrOrErr = suggestAddress(MaxMemoryAllocation); + if (!VAddrOrErr) + return VAddrOrErr.takeError(); + VAddr = *VAddrOrErr; + } DP("Request %ld bytes allocated at %p\n", MaxMemoryAllocation, VAddr); @@ -117,8 +127,11 @@ private: constexpr size_t STEP = 1024 * 1024 * 1024ULL; MemoryStart = nullptr; for (TotalSize = MAX_MEMORY_ALLOCATION; TotalSize > 0; TotalSize -= STEP) { - MemoryStart = + auto MemoryStartOrErr = Device->allocate(TotalSize, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT); + if (!MemoryStartOrErr) + return MemoryStartOrErr.takeError(); + MemoryStart = *MemoryStartOrErr; if (MemoryStart) break; } @@ -214,15 +227,7 @@ public: raw_fd_ostream OS(ImageName, EC); if (EC) report_fatal_error("Error saving image : " + StringRef(EC.message())); - if (const auto *TgtImageBitcode = Image.getTgtImageBitcode()) { - size_t Size = utils::getPtrDiff(TgtImageBitcode->ImageEnd, - TgtImageBitcode->ImageStart); - MemoryBufferRef MBR = MemoryBufferRef( - StringRef((const char *)TgtImageBitcode->ImageStart, Size), ""); - OS << MBR.getBuffer(); - } else { - OS << Image.getMemoryBuffer().getBuffer(); - } + OS << Image.getMemoryBuffer().getBuffer(); OS.close(); } @@ -360,13 +365,15 @@ public: return Plugin::success(); } - void deinit() { + Error deinit() { if (UsedVAMap) { if (auto Err = Device->memoryVAUnMap(MemoryStart, TotalSize)) - report_fatal_error("Error on releasing virtual memory space"); + return Err; } else { - Device->free(MemoryStart); + if (auto Err = Device->free(MemoryStart)) + return Err; } + return Plugin::success(); } }; } // namespace llvm::omp::target::plugin @@ -813,9 +820,6 @@ Error GenericDeviceTy::unloadBinary(DeviceImageTy *Image) { return Err; } - if (Image->getTgtImageBitcode()) - Plugin.getJIT().erase(*Image->getTgtImageBitcode(), Image->getDevice()); - return unloadBinaryImpl(Image); } @@ -849,7 +853,8 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) { RecordReplayTy &RecordReplay = Plugin.getRecordReplay(); if (RecordReplay.isRecordingOrReplaying()) - RecordReplay.deinit(); + if (auto Err = RecordReplay.deinit()) + return Err; if (RPCServer) if (auto Err = RPCServer->deinitDevice(*this)) @@ -865,32 +870,29 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) { return deinitImpl(); } -Expected<DeviceImageTy *> -GenericDeviceTy::loadBinary(GenericPluginTy &Plugin, - const __tgt_device_image *InputTgtImage) { - assert(InputTgtImage && "Expected non-null target image"); - DP("Load data from image " DPxMOD "\n", DPxPTR(InputTgtImage->ImageStart)); - - auto PostJITImageOrErr = Plugin.getJIT().process(*InputTgtImage, *this); - if (!PostJITImageOrErr) { - auto Err = PostJITImageOrErr.takeError(); - REPORT("Failure to jit IR image %p on device %d: %s\n", InputTgtImage, - DeviceId, toStringWithoutConsuming(Err).data()); - return Plugin::error(ErrorCode::COMPILE_FAILURE, std::move(Err), - "failure to jit IR image"); +Expected<DeviceImageTy *> GenericDeviceTy::loadBinary(GenericPluginTy &Plugin, + StringRef InputTgtImage) { + DP("Load data from image " DPxMOD "\n", DPxPTR(InputTgtImage.bytes_begin())); + + std::unique_ptr<MemoryBuffer> Buffer; + if (identify_magic(InputTgtImage) == file_magic::bitcode) { + auto CompiledImageOrErr = Plugin.getJIT().process(InputTgtImage, *this); + if (!CompiledImageOrErr) { + return Plugin::error(ErrorCode::COMPILE_FAILURE, + CompiledImageOrErr.takeError(), + "failure to jit IR image"); + } + Buffer = std::move(*CompiledImageOrErr); + } else { + Buffer = MemoryBuffer::getMemBufferCopy(InputTgtImage); } // Load the binary and allocate the image object. Use the next available id // for the image id, which is the number of previously loaded images. - auto ImageOrErr = - loadBinaryImpl(PostJITImageOrErr.get(), LoadedImages.size()); + auto ImageOrErr = loadBinaryImpl(std::move(Buffer), LoadedImages.size()); if (!ImageOrErr) return ImageOrErr.takeError(); - DeviceImageTy *Image = *ImageOrErr; - assert(Image != nullptr && "Invalid image"); - if (InputTgtImage != PostJITImageOrErr.get()) - Image->setTgtImageBitcode(InputTgtImage); // Add the image to list. LoadedImages.push_back(Image); @@ -912,12 +914,12 @@ GenericDeviceTy::loadBinary(GenericPluginTy &Plugin, #ifdef OMPT_SUPPORT if (ompt::Initialized) { - size_t Bytes = - utils::getPtrDiff(InputTgtImage->ImageEnd, InputTgtImage->ImageStart); + size_t Bytes = InputTgtImage.size(); performOmptCallback( device_load, Plugin.getUserId(DeviceId), /*FileName=*/nullptr, /*FileOffset=*/0, /*VmaInFile=*/nullptr, - /*ImgSize=*/Bytes, /*HostAddr=*/InputTgtImage->ImageStart, + /*ImgSize=*/Bytes, + /*HostAddr=*/const_cast<unsigned char *>(InputTgtImage.bytes_begin()), /*DeviceAddr=*/nullptr, /* FIXME: ModuleId */ 0); } #endif @@ -1309,10 +1311,12 @@ Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr, switch (Kind) { case TARGET_ALLOC_DEFAULT: - case TARGET_ALLOC_DEVICE_NON_BLOCKING: case TARGET_ALLOC_DEVICE: if (MemoryManager) { - Alloc = MemoryManager->allocate(Size, HostPtr); + auto AllocOrErr = MemoryManager->allocate(Size, HostPtr); + if (!AllocOrErr) + return AllocOrErr.takeError(); + Alloc = *AllocOrErr; if (!Alloc) return Plugin::error(ErrorCode::OUT_OF_RESOURCES, "failed to allocate from memory manager"); @@ -1320,12 +1324,16 @@ Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr, } [[fallthrough]]; case TARGET_ALLOC_HOST: - case TARGET_ALLOC_SHARED: - Alloc = allocate(Size, HostPtr, Kind); + case TARGET_ALLOC_SHARED: { + auto AllocOrErr = allocate(Size, HostPtr, Kind); + if (!AllocOrErr) + return AllocOrErr.takeError(); + Alloc = *AllocOrErr; if (!Alloc) return Plugin::error(ErrorCode::OUT_OF_RESOURCES, "failed to allocate from device allocator"); } + } // Report error if the memory manager or the device allocator did not return // any memory buffer. @@ -1397,29 +1405,19 @@ Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) { #undef DEALLOCATION_ERROR } - int Res; switch (Kind) { case TARGET_ALLOC_DEFAULT: - case TARGET_ALLOC_DEVICE_NON_BLOCKING: case TARGET_ALLOC_DEVICE: if (MemoryManager) { - Res = MemoryManager->free(TgtPtr); - if (Res) - return Plugin::error( - ErrorCode::OUT_OF_RESOURCES, - "failure to deallocate device pointer %p via memory manager", - TgtPtr); + if (auto Err = MemoryManager->free(TgtPtr)) + return Err; break; } [[fallthrough]]; case TARGET_ALLOC_HOST: case TARGET_ALLOC_SHARED: - Res = free(TgtPtr, Kind); - if (Res) - return Plugin::error( - ErrorCode::UNKNOWN, - "failure to deallocate device pointer %p via device deallocator", - TgtPtr); + if (auto Err = free(TgtPtr, Kind)) + return Err; } // Unregister deallocated pinned memory buffer if the type is host memory. @@ -1727,28 +1725,26 @@ Expected<bool> GenericPluginTy::checkBitcodeImage(StringRef Image) const { int32_t GenericPluginTy::is_initialized() const { return Initialized; } -int32_t GenericPluginTy::is_plugin_compatible(__tgt_device_image *Image) { - StringRef Buffer(reinterpret_cast<const char *>(Image->ImageStart), - utils::getPtrDiff(Image->ImageEnd, Image->ImageStart)); - +int32_t GenericPluginTy::isPluginCompatible(StringRef Image) { auto HandleError = [&](Error Err) -> bool { [[maybe_unused]] std::string ErrStr = toString(std::move(Err)); - DP("Failure to check validity of image %p: %s", Image, ErrStr.c_str()); + DP("Failure to check validity of image %p: %s", Image.data(), + ErrStr.c_str()); return false; }; - switch (identify_magic(Buffer)) { + switch (identify_magic(Image)) { case file_magic::elf: case file_magic::elf_relocatable: case file_magic::elf_executable: case file_magic::elf_shared_object: case file_magic::elf_core: { - auto MatchOrErr = checkELFImage(Buffer); + auto MatchOrErr = checkELFImage(Image); if (Error Err = MatchOrErr.takeError()) return HandleError(std::move(Err)); return *MatchOrErr; } case file_magic::bitcode: { - auto MatchOrErr = checkBitcodeImage(Buffer); + auto MatchOrErr = checkBitcodeImage(Image); if (Error Err = MatchOrErr.takeError()) return HandleError(std::move(Err)); return *MatchOrErr; @@ -1758,36 +1754,33 @@ int32_t GenericPluginTy::is_plugin_compatible(__tgt_device_image *Image) { } } -int32_t GenericPluginTy::is_device_compatible(int32_t DeviceId, - __tgt_device_image *Image) { - StringRef Buffer(reinterpret_cast<const char *>(Image->ImageStart), - utils::getPtrDiff(Image->ImageEnd, Image->ImageStart)); - +int32_t GenericPluginTy::isDeviceCompatible(int32_t DeviceId, StringRef Image) { auto HandleError = [&](Error Err) -> bool { [[maybe_unused]] std::string ErrStr = toString(std::move(Err)); - DP("Failure to check validity of image %p: %s", Image, ErrStr.c_str()); + DP("Failure to check validity of image %p: %s", Image.data(), + ErrStr.c_str()); return false; }; - switch (identify_magic(Buffer)) { + switch (identify_magic(Image)) { case file_magic::elf: case file_magic::elf_relocatable: case file_magic::elf_executable: case file_magic::elf_shared_object: case file_magic::elf_core: { - auto MatchOrErr = checkELFImage(Buffer); + auto MatchOrErr = checkELFImage(Image); if (Error Err = MatchOrErr.takeError()) return HandleError(std::move(Err)); if (!*MatchOrErr) return false; // Perform plugin-dependent checks for the specific architecture if needed. - auto CompatibleOrErr = isELFCompatible(DeviceId, Buffer); + auto CompatibleOrErr = isELFCompatible(DeviceId, Image); if (Error Err = CompatibleOrErr.takeError()) return HandleError(std::move(Err)); return *CompatibleOrErr; } case file_magic::bitcode: { - auto MatchOrErr = checkBitcodeImage(Buffer); + auto MatchOrErr = checkBitcodeImage(Image); if (Error Err = MatchOrErr.takeError()) return HandleError(std::move(Err)); return *MatchOrErr; @@ -1848,7 +1841,9 @@ int32_t GenericPluginTy::load_binary(int32_t DeviceId, __tgt_device_binary *Binary) { GenericDeviceTy &Device = getDevice(DeviceId); - auto ImageOrErr = Device.loadBinary(*this, TgtImage); + StringRef Buffer(reinterpret_cast<const char *>(TgtImage->ImageStart), + utils::getPtrDiff(TgtImage->ImageEnd, TgtImage->ImageStart)); + auto ImageOrErr = Device.loadBinary(*this, Buffer); if (!ImageOrErr) { auto Err = ImageOrErr.takeError(); REPORT("Failure to load binary image %p on device %d: %s\n", TgtImage, diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp index 678be78b56af..e19f2ef94de6 100644 --- a/offload/plugins-nextgen/common/src/RPC.cpp +++ b/offload/plugins-nextgen/common/src/RPC.cpp @@ -28,15 +28,22 @@ rpc::Status handleOffloadOpcodes(plugin::GenericDeviceTy &Device, switch (Port.get_opcode()) { case LIBC_MALLOC: { Port.recv_and_send([&](rpc::Buffer *Buffer, uint32_t) { - Buffer->data[0] = reinterpret_cast<uintptr_t>(Device.allocate( - Buffer->data[0], nullptr, TARGET_ALLOC_DEVICE_NON_BLOCKING)); + auto PtrOrErr = + Device.allocate(Buffer->data[0], nullptr, TARGET_ALLOC_DEVICE); + void *Ptr = nullptr; + if (!PtrOrErr) + llvm::consumeError(PtrOrErr.takeError()); + else + Ptr = *PtrOrErr; + Buffer->data[0] = reinterpret_cast<uintptr_t>(Ptr); }); break; } case LIBC_FREE: { Port.recv([&](rpc::Buffer *Buffer, uint32_t) { - Device.free(reinterpret_cast<void *>(Buffer->data[0]), - TARGET_ALLOC_DEVICE_NON_BLOCKING); + if (auto Err = Device.free(reinterpret_cast<void *>(Buffer->data[0]), + TARGET_ALLOC_DEVICE)) + llvm::consumeError(std::move(Err)); }); break; } @@ -171,9 +178,13 @@ Error RPCServerTy::initDevice(plugin::GenericDeviceTy &Device, plugin::DeviceImageTy &Image) { uint64_t NumPorts = std::min(Device.requestedRPCPortCount(), rpc::MAX_PORT_COUNT); - void *RPCBuffer = Device.allocate( + auto RPCBufferOrErr = Device.allocate( rpc::Server::allocation_size(Device.getWarpSize(), NumPorts), nullptr, TARGET_ALLOC_HOST); + if (!RPCBufferOrErr) + return RPCBufferOrErr.takeError(); + + void *RPCBuffer = *RPCBufferOrErr; if (!RPCBuffer) return plugin::Plugin::error( error::ErrorCode::UNKNOWN, @@ -198,7 +209,8 @@ Error RPCServerTy::initDevice(plugin::GenericDeviceTy &Device, Error RPCServerTy::deinitDevice(plugin::GenericDeviceTy &Device) { std::lock_guard<decltype(BufferMutex)> Lock(BufferMutex); - Device.free(Buffers[Device.getDeviceId()], TARGET_ALLOC_HOST); + if (auto Err = Device.free(Buffers[Device.getDeviceId()], TARGET_ALLOC_HOST)) + return Err; Buffers[Device.getDeviceId()] = nullptr; Devices[Device.getDeviceId()] = nullptr; return Error::success(); diff --git a/offload/plugins-nextgen/cuda/src/rtl.cpp b/offload/plugins-nextgen/cuda/src/rtl.cpp index af3c74636bff..e5c4a1bfa985 100644 --- a/offload/plugins-nextgen/cuda/src/rtl.cpp +++ b/offload/plugins-nextgen/cuda/src/rtl.cpp @@ -81,8 +81,8 @@ CUresult cuMemFreeAsync(CUdeviceptr dptr, CUstream hStream) {} struct CUDADeviceImageTy : public DeviceImageTy { /// Create the CUDA image with the id and the target image pointer. CUDADeviceImageTy(int32_t ImageId, GenericDeviceTy &Device, - const __tgt_device_image *TgtImage) - : DeviceImageTy(ImageId, Device, TgtImage), Module(nullptr) {} + std::unique_ptr<MemoryBuffer> &&TgtImage) + : DeviceImageTy(ImageId, Device, std::move(TgtImage)), Module(nullptr) {} /// Load the image as a CUDA module. Error loadModule() { @@ -160,15 +160,15 @@ struct CUDAKernelTy : public GenericKernelTy { /// Return maximum block size for maximum occupancy Expected<uint64_t> maxGroupSize(GenericDeviceTy &, uint64_t DynamicMemSize) const override { - int minGridSize; - int maxBlockSize; + int MinGridSize; + int MaxBlockSize; auto Res = cuOccupancyMaxPotentialBlockSize( - &minGridSize, &maxBlockSize, Func, NULL, DynamicMemSize, INT_MAX); + &MinGridSize, &MaxBlockSize, Func, NULL, DynamicMemSize, INT_MAX); if (auto Err = Plugin::check( Res, "error in cuOccupancyMaxPotentialBlockSize: %s")) { return Err; } - return maxBlockSize; + return MaxBlockSize; } private: @@ -385,6 +385,8 @@ struct CUDADeviceTy : public GenericDeviceTy { if (auto Err = CUDAImage.unloadModule()) return Err; + // Destroy the associated memory and invalidate the object. + Plugin.free(Image); return Plugin::success(); } @@ -418,20 +420,12 @@ struct CUDADeviceTy : public GenericDeviceTy { virtual Error callGlobalConstructors(GenericPluginTy &Plugin, DeviceImageTy &Image) override { - // Check for the presence of global destructors at initialization time. This - // is required when the image may be deallocated before destructors are run. - GenericGlobalHandlerTy &Handler = Plugin.getGlobalHandler(); - if (Handler.isSymbolInImage(*this, Image, "nvptx$device$fini")) - Image.setPendingGlobalDtors(); - return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/true); } virtual Error callGlobalDestructors(GenericPluginTy &Plugin, DeviceImageTy &Image) override { - if (Image.hasPendingGlobalDtors()) - return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/false); - return Plugin::success(); + return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/false); } Expected<std::unique_ptr<MemoryBuffer>> @@ -549,14 +543,15 @@ struct CUDADeviceTy : public GenericDeviceTy { CUdevice getCUDADevice() const { return Device; } /// Load the binary image into the device and allocate an image object. - Expected<DeviceImageTy *> loadBinaryImpl(const __tgt_device_image *TgtImage, - int32_t ImageId) override { + Expected<DeviceImageTy *> + loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, + int32_t ImageId) override { if (auto Err = setContext()) return std::move(Err); // Allocate and initialize the image object. CUDADeviceImageTy *CUDAImage = Plugin.allocate<CUDADeviceImageTy>(); - new (CUDAImage) CUDADeviceImageTy(ImageId, *this, TgtImage); + new (CUDAImage) CUDADeviceImageTy(ImageId, *this, std::move(TgtImage)); // Load the CUDA module. if (auto Err = CUDAImage->loadModule()) @@ -566,14 +561,12 @@ struct CUDADeviceTy : public GenericDeviceTy { } /// Allocate memory on the device or related to the device. - void *allocate(size_t Size, void *, TargetAllocTy Kind) override { + Expected<void *> allocate(size_t Size, void *, TargetAllocTy Kind) override { if (Size == 0) return nullptr; - if (auto Err = setContext()) { - REPORT("Failure to alloc memory: %s\n", toString(std::move(Err)).data()); - return nullptr; - } + if (auto Err = setContext()) + return std::move(Err); void *MemAlloc = nullptr; CUdeviceptr DevicePtr; @@ -592,35 +585,20 @@ struct CUDADeviceTy : public GenericDeviceTy { Res = cuMemAllocManaged(&DevicePtr, Size, CU_MEM_ATTACH_GLOBAL); MemAlloc = (void *)DevicePtr; break; - case TARGET_ALLOC_DEVICE_NON_BLOCKING: { - CUstream Stream; - if ((Res = cuStreamCreate(&Stream, CU_STREAM_NON_BLOCKING))) - break; - if ((Res = cuMemAllocAsync(&DevicePtr, Size, Stream))) - break; - cuStreamSynchronize(Stream); - Res = cuStreamDestroy(Stream); - MemAlloc = (void *)DevicePtr; - } } - if (auto Err = - Plugin::check(Res, "error in cuMemAlloc[Host|Managed]: %s")) { - REPORT("Failure to alloc memory: %s\n", toString(std::move(Err)).data()); - return nullptr; - } + if (auto Err = Plugin::check(Res, "error in cuMemAlloc[Host|Managed]: %s")) + return std::move(Err); return MemAlloc; } /// Deallocate memory on the device or related to the device. - int free(void *TgtPtr, TargetAllocTy Kind) override { + Error free(void *TgtPtr, TargetAllocTy Kind) override { if (TgtPtr == nullptr) - return OFFLOAD_SUCCESS; + return Plugin::success(); - if (auto Err = setContext()) { - REPORT("Failure to free memory: %s\n", toString(std::move(Err)).data()); - return OFFLOAD_FAIL; - } + if (auto Err = setContext()) + return Err; CUresult Res; switch (Kind) { @@ -632,22 +610,9 @@ struct CUDADeviceTy : public GenericDeviceTy { case TARGET_ALLOC_HOST: Res = cuMemFreeHost(TgtPtr); break; - case TARGET_ALLOC_DEVICE_NON_BLOCKING: { - CUstream Stream; - if ((Res = cuStreamCreate(&Stream, CU_STREAM_NON_BLOCKING))) - break; - cuMemFreeAsync(reinterpret_cast<CUdeviceptr>(TgtPtr), Stream); - cuStreamSynchronize(Stream); - if ((Res = cuStreamDestroy(Stream))) - break; - } } - if (auto Err = Plugin::check(Res, "error in cuMemFree[Host]: %s")) { - REPORT("Failure to free memory: %s\n", toString(std::move(Err)).data()); - return OFFLOAD_FAIL; - } - return OFFLOAD_SUCCESS; + return Plugin::check(Res, "error in cuMemFree[Host]: %s"); } /// Synchronize current thread with the pending operations on the async info. @@ -952,6 +917,50 @@ struct CUDADeviceTy : public GenericDeviceTy { return Plugin::success(); } + interop_spec_t selectInteropPreference(int32_t InteropType, + int32_t NumPrefers, + interop_spec_t *Prefers) override { + return interop_spec_t{tgt_fr_cuda, {true, 0}, 0}; + } + + Expected<omp_interop_val_t *> + createInterop(int32_t InteropType, interop_spec_t &InteropSpec) override { + auto *Ret = new omp_interop_val_t( + DeviceId, static_cast<kmp_interop_type_t>(InteropType)); + Ret->fr_id = tgt_fr_cuda; + Ret->vendor_id = omp_vendor_nvidia; + + if (InteropType == kmp_interop_type_target || + InteropType == kmp_interop_type_targetsync) { + Ret->device_info.Platform = nullptr; + Ret->device_info.Device = reinterpret_cast<void *>(Device); + Ret->device_info.Context = Context; + } + + if (InteropType == kmp_interop_type_targetsync) { + Ret->async_info = new __tgt_async_info(); + if (auto Err = setContext()) + return Err; + CUstream Stream; + if (auto Err = CUDAStreamManager.getResource(Stream)) + return Err; + + Ret->async_info->Queue = Stream; + } + return Ret; + } + + Error releaseInterop(omp_interop_val_t *Interop) override { + if (!Interop) + return Plugin::success(); + + if (Interop->async_info) + delete Interop->async_info; + + delete Interop; + return Plugin::success(); + } + Error enqueueHostCallImpl(void (*Callback)(void *), void *UserData, AsyncInfoWrapperTy &AsyncInfo) override { if (auto Err = setContext()) @@ -1299,7 +1308,7 @@ private: // Perform a quick check for the named kernel in the image. The kernel // should be created by the 'nvptx-lower-ctor-dtor' pass. GenericGlobalHandlerTy &Handler = Plugin.getGlobalHandler(); - if (IsCtor && !Handler.isSymbolInImage(*this, Image, KernelName)) + if (!Handler.isSymbolInImage(*this, Image, KernelName)) return Plugin::success(); // The Nvidia backend cannot handle creating the ctor / dtor array @@ -1334,8 +1343,12 @@ private: // Allocate a buffer to store all of the known constructor / destructor // functions in so we can iterate them on the device. - void *Buffer = + auto BufferOrErr = allocate(Funcs.size() * sizeof(void *), nullptr, TARGET_ALLOC_DEVICE); + if (!BufferOrErr) + return BufferOrErr.takeError(); + + void *Buffer = *BufferOrErr; if (!Buffer) return Plugin::error(ErrorCode::OUT_OF_RESOURCES, "failed to allocate memory for global buffer"); @@ -1384,12 +1397,10 @@ private: Error Err = Plugin::success(); AsyncInfoWrapper.finalize(Err); + if (Err) + return Err; - if (free(Buffer, TARGET_ALLOC_DEVICE) != OFFLOAD_SUCCESS) - return Plugin::error(ErrorCode::UNKNOWN, - "failed to free memory for global buffer"); - - return Err; + return free(Buffer, TARGET_ALLOC_DEVICE); } /// Stream manager for CUDA streams. @@ -1586,7 +1597,7 @@ struct CUDAPluginTy final : public GenericPluginTy { unsigned SM = Header.e_ident[ELF::EI_ABIVERSION] == ELF::ELFABIVERSION_CUDA_V1 ? Header.e_flags & ELF::EF_CUDA_SM - : (Header.e_flags & ELF::EF_CUDA_SM_MASK) >> 8; + : (Header.e_flags & ELF::EF_CUDA_SM_MASK) >> ELF::EF_CUDA_SM_OFFSET; CUdevice Device; CUresult Res = cuDeviceGet(&Device, DeviceId); diff --git a/offload/plugins-nextgen/host/src/rtl.cpp b/offload/plugins-nextgen/host/src/rtl.cpp index 5436cae3b029..0845032d0aae 100644 --- a/offload/plugins-nextgen/host/src/rtl.cpp +++ b/offload/plugins-nextgen/host/src/rtl.cpp @@ -131,8 +131,8 @@ private: struct GenELF64DeviceImageTy : public DeviceImageTy { /// Create the GenELF64 image with the id and the target image pointer. GenELF64DeviceImageTy(int32_t ImageId, GenericDeviceTy &Device, - const __tgt_device_image *TgtImage) - : DeviceImageTy(ImageId, Device, TgtImage), DynLib() {} + std::unique_ptr<MemoryBuffer> &&TgtImage) + : DeviceImageTy(ImageId, Device, std::move(TgtImage)), DynLib() {} /// Getter and setter for the dynamic library. DynamicLibrary &getDynamicLibrary() { return DynLib; } @@ -189,11 +189,12 @@ struct GenELF64DeviceTy : public GenericDeviceTy { Error setContext() override { return Plugin::success(); } /// Load the binary image into the device and allocate an image object. - Expected<DeviceImageTy *> loadBinaryImpl(const __tgt_device_image *TgtImage, - int32_t ImageId) override { + Expected<DeviceImageTy *> + loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage, + int32_t ImageId) override { // Allocate and initialize the image object. GenELF64DeviceImageTy *Image = Plugin.allocate<GenELF64DeviceImageTy>(); - new (Image) GenELF64DeviceImageTy(ImageId, *this, TgtImage); + new (Image) GenELF64DeviceImageTy(ImageId, *this, std::move(TgtImage)); // Create a temporary file. char TmpFileName[] = "/tmp/tmpfile_XXXXXX"; @@ -239,7 +240,7 @@ struct GenELF64DeviceTy : public GenericDeviceTy { } /// Allocate memory. Use std::malloc in all cases. - void *allocate(size_t Size, void *, TargetAllocTy Kind) override { + Expected<void *> allocate(size_t Size, void *, TargetAllocTy Kind) override { if (Size == 0) return nullptr; @@ -249,7 +250,6 @@ struct GenELF64DeviceTy : public GenericDeviceTy { case TARGET_ALLOC_DEVICE: case TARGET_ALLOC_HOST: case TARGET_ALLOC_SHARED: - case TARGET_ALLOC_DEVICE_NON_BLOCKING: MemAlloc = std::malloc(Size); break; } @@ -257,9 +257,9 @@ struct GenELF64DeviceTy : public GenericDeviceTy { } /// Free the memory. Use std::free in all cases. - int free(void *TgtPtr, TargetAllocTy Kind) override { + Error free(void *TgtPtr, TargetAllocTy Kind) override { std::free(TgtPtr); - return OFFLOAD_SUCCESS; + return Plugin::success(); } /// This plugin does nothing to lock buffers. Do not return an error, just diff --git a/offload/test/mapping/lambda_by_value.cpp b/offload/test/mapping/lambda_by_value.cpp index 5516dedd72a9..4c0278d40592 100644 --- a/offload/test/mapping/lambda_by_value.cpp +++ b/offload/test/mapping/lambda_by_value.cpp @@ -1,4 +1,5 @@ -// RUN: %libomptarget-compilexx-run-and-check-generic +// RUN: %libomptarget-compileopt-generic -fno-exceptions +// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic #include <stdint.h> #include <stdio.h> diff --git a/offload/test/mapping/map_back_race.cpp b/offload/test/mapping/map_back_race.cpp index 8a988d3be3b4..49bbe87e2449 100644 --- a/offload/test/mapping/map_back_race.cpp +++ b/offload/test/mapping/map_back_race.cpp @@ -2,6 +2,9 @@ // Taken from https://github.com/llvm/llvm-project/issues/54216 +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// UNSUPPORTED: gpu + #include <algorithm> #include <cstdlib> #include <iostream> diff --git a/offload/test/mapping/map_both_pointer_pointee.c b/offload/test/mapping/map_both_pointer_pointee.c index 7be1ba465e7d..1934b702dbba 100644 --- a/offload/test/mapping/map_both_pointer_pointee.c +++ b/offload/test/mapping/map_both_pointer_pointee.c @@ -1,11 +1,10 @@ -// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu -// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu -// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu -// RUN: %libomptarget-compile-run-and-check-x86_64-unknown-linux-gnu -// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda +// RUN: %libomptarget-compile-run-and-check-generic // REQUIRES: unified_shared_memory // UNSUPPORTED: amdgcn-amd-amdhsa +// +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// XFAIL: nvidiagpu #pragma omp declare target int *ptr1; diff --git a/offload/test/mapping/map_ptr_and_star_local.c b/offload/test/mapping/map_ptr_and_star_local.c index cc826b3c0290..97fa7cd53715 100644 --- a/offload/test/mapping/map_ptr_and_star_local.c +++ b/offload/test/mapping/map_ptr_and_star_local.c @@ -1,6 +1,9 @@ -// RUN: %libomptarget-compilexx-run-and-check-generic +// RUN: %libomptarget-compile-run-and-check-generic // REQUIRES: libc +// +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// XFAIL: gpu #include <omp.h> #include <stdio.h> diff --git a/offload/test/mapping/map_structptr_and_member_global.c b/offload/test/mapping/map_structptr_and_member_global.c index 960eea419964..f855e87d7218 100644 --- a/offload/test/mapping/map_structptr_and_member_global.c +++ b/offload/test/mapping/map_structptr_and_member_global.c @@ -1,6 +1,9 @@ -// RUN: %libomptarget-compilexx-run-and-check-generic +// RUN: %libomptarget-compile-run-and-check-generic // REQUIRES: libc +// +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// XFAIL: gpu #include <omp.h> #include <stdio.h> diff --git a/offload/test/mapping/map_structptr_and_member_local.c b/offload/test/mapping/map_structptr_and_member_local.c index bd759407ef09..bd9e2a89eb6f 100644 --- a/offload/test/mapping/map_structptr_and_member_local.c +++ b/offload/test/mapping/map_structptr_and_member_local.c @@ -1,6 +1,9 @@ -// RUN: %libomptarget-compilexx-run-and-check-generic +// RUN: %libomptarget-compile-run-and-check-generic // REQUIRES: libc +// +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// XFAIL: gpu #include <omp.h> #include <stdio.h> diff --git a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu index 1f84a0e1288d..b2e1edf51e17 100644 --- a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu +++ b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu @@ -5,10 +5,10 @@ // RUN: %t | %fcheck-generic // clang-format on -// UNSUPPORTED: aarch64-unknown-linux-gnu -// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO -// UNSUPPORTED: x86_64-unknown-linux-gnu -// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO +// REQUIRES: gpu +// +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// XFAIL: gpu #include <stdio.h> diff --git a/offload/test/offloading/bug51781.c b/offload/test/offloading/bug51781.c index 2f30b035afbb..ff7fa51aafc2 100644 --- a/offload/test/offloading/bug51781.c +++ b/offload/test/offloading/bug51781.c @@ -16,6 +16,7 @@ // the generic state machine. // // RUN: %libomptarget-compile-generic -O2 -foffload-lto -Rpass=openmp-opt \ +// RUN: -Xoffload-linker -mllvm=-openmp-opt-disable-spmdization \ // RUN: -mllvm -openmp-opt-disable-spmdization > %t.custom 2>&1 // RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=CUSTOM -input-file=%t.custom // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom @@ -24,7 +25,9 @@ // Repeat with reduction clause, which has managed to break the custom state // machine in the past. // -// RUN: %libomptarget-compile-generic -O2 -foffload-lto -Rpass=openmp-opt -DADD_REDUCTION \ +// RUN: %libomptarget-compile-generic -O2 -foffload-lto -Rpass=openmp-opt \ +// RUN: -DADD_REDUCTION \ +// RUN: -Xoffload-linker -mllvm=-openmp-opt-disable-spmdization \ // RUN: -mllvm -openmp-opt-disable-spmdization > %t.custom 2>&1 // RUN: %fcheck-nvptx64-nvidia-cuda -check-prefix=CUSTOM -input-file=%t.custom // RUN: %fcheck-amdgcn-amd-amdhsa -check-prefix=CUSTOM -input-file=%t.custom diff --git a/offload/test/offloading/force-usm.cpp b/offload/test/offloading/force-usm.cpp index a043ba47f54a..9988c3dc4e9e 100644 --- a/offload/test/offloading/force-usm.cpp +++ b/offload/test/offloading/force-usm.cpp @@ -48,7 +48,7 @@ int main(void) { // clang-format off // NO-USM: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4 -// NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=12 +// NO-USM: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=12 // NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4 // NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=8, Name=pGI // NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4 diff --git a/offload/test/offloading/fortran/declare-target-automap.f90 b/offload/test/offloading/fortran/declare-target-automap.f90 index b9c2d34c834f..b44c0b281527 100644 --- a/offload/test/offloading/fortran/declare-target-automap.f90 +++ b/offload/test/offloading/fortran/declare-target-automap.f90 @@ -1,6 +1,9 @@ !Offloading test for AUTOMAP modifier in declare target enter ! REQUIRES: flang, amdgpu +! FIXME: https://github.com/llvm/llvm-project/issues/161265 +! XFAIL: amdgpu + ! RUN: %libomptarget-compile-fortran-run-and-check-generic program automap_program use iso_c_binding, only: c_loc diff --git a/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy-2d.f90 b/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy-2d.f90 new file mode 100644 index 000000000000..c6f576acb90b --- /dev/null +++ b/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy-2d.f90 @@ -0,0 +1,53 @@ +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-generic -fdo-concurrent-to-openmp=device +! RUN: env LIBOMPTARGET_INFO=16 %libomptarget-run-generic 2>&1 | %fcheck-generic +module saxpymod + use iso_fortran_env + public :: saxpy +contains + +subroutine saxpy(a, x, y, n, m) + use iso_fortran_env + implicit none + integer,intent(in) :: n, m + real(kind=real32),intent(in) :: a + real(kind=real32), dimension(:,:),intent(in) :: x + real(kind=real32), dimension(:,:),intent(inout) :: y + integer :: i, j + + do concurrent(i=1:n, j=1:m) + y(i,j) = a * x(i,j) + y(i,j) + end do + + write(*,*) "plausibility check:" + write(*,'("y(1,1) ",f8.6)') y(1,1) + write(*,'("y(n,m) ",f8.6)') y(n,m) +end subroutine saxpy + +end module saxpymod + +program main + use iso_fortran_env + use saxpymod, ONLY:saxpy + implicit none + + integer,parameter :: n = 1000, m=10000 + real(kind=real32), allocatable, dimension(:,:) :: x, y + real(kind=real32) :: a + integer :: i + + allocate(x(1:n,1:m), y(1:n,1:m)) + a = 2.0_real32 + x(:,:) = 1.0_real32 + y(:,:) = 2.0_real32 + + call saxpy(a, x, y, n, m) + + deallocate(x,y) +end program main + +! CHECK: "PluginInterface" device {{[0-9]+}} info: Launching kernel {{.*}} +! CHECK: plausibility check: +! CHECK: y(1,1) 4.0 +! CHECK: y(n,m) 4.0 diff --git a/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy.f90 b/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy.f90 new file mode 100644 index 000000000000..e094a1d7459e --- /dev/null +++ b/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy.f90 @@ -0,0 +1,53 @@ +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-generic -fdo-concurrent-to-openmp=device +! RUN: env LIBOMPTARGET_INFO=16 %libomptarget-run-generic 2>&1 | %fcheck-generic +module saxpymod + use iso_fortran_env + public :: saxpy +contains + +subroutine saxpy(a, x, y, n) + use iso_fortran_env + implicit none + integer,intent(in) :: n + real(kind=real32),intent(in) :: a + real(kind=real32), dimension(:),intent(in) :: x + real(kind=real32), dimension(:),intent(inout) :: y + integer :: i + + do concurrent(i=1:n) + y(i) = a * x(i) + y(i) + end do + + write(*,*) "plausibility check:" + write(*,'("y(1) ",f8.6)') y(1) + write(*,'("y(n) ",f8.6)') y(n) +end subroutine saxpy + +end module saxpymod + +program main + use iso_fortran_env + use saxpymod, ONLY:saxpy + implicit none + + integer,parameter :: n = 10000000 + real(kind=real32), allocatable, dimension(:) :: x, y + real(kind=real32) :: a + integer :: i + + allocate(x(1:n), y(1:n)) + a = 2.0_real32 + x(:) = 1.0_real32 + y(:) = 2.0_real32 + + call saxpy(a, x, y, n) + + deallocate(x,y) +end program main + +! CHECK: "PluginInterface" device {{[0-9]+}} info: Launching kernel {{.*}} +! CHECK: plausibility check: +! CHECK: y(1) 4.0 +! CHECK: y(n) 4.0 diff --git a/offload/test/offloading/fortran/target-declare-mapper-allocatable.f90 b/offload/test/offloading/fortran/target-declare-mapper-allocatable.f90 new file mode 100644 index 000000000000..d8d5e1b5631a --- /dev/null +++ b/offload/test/offloading/fortran/target-declare-mapper-allocatable.f90 @@ -0,0 +1,48 @@ +! This test validates that declare mapper for a derived type with an +! allocatable component preserves TO/FROM semantics for the component, +! ensuring the payload is copied back to the host on target exit. + +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic + +program target_declare_mapper_allocatable + implicit none + + type :: real_t + real, allocatable :: real_arr(:) + end type real_t + + ! Map the allocatable array payload via a named mapper. + !$omp declare mapper (xyz : real_t :: t) map(tofrom: t%real_arr) + + type(real_t) :: r + integer :: i + logical :: ok + + allocate(r%real_arr(10)) + r%real_arr = 1.0 + + !$omp target map(mapper(xyz), tofrom: r) + do i = 1, size(r%real_arr) + r%real_arr(i) = 3.0 + end do + !$omp end target + + ok = .true. + do i = 1, size(r%real_arr) + if (r%real_arr(i) /= 3.0) ok = .false. + end do + if (ok) then + print *, "Test passed!" + else + print *, "Test failed!" + do i = 1, size(r%real_arr) + print *, r%real_arr(i) + end do + end if + + deallocate(r%real_arr) +end program target_declare_mapper_allocatable + +! CHECK: Test passed! diff --git a/offload/test/offloading/fortran/target-declare-mapper-parent-allocatable.f90 b/offload/test/offloading/fortran/target-declare-mapper-parent-allocatable.f90 new file mode 100644 index 000000000000..65e04af66e02 --- /dev/null +++ b/offload/test/offloading/fortran/target-declare-mapper-parent-allocatable.f90 @@ -0,0 +1,43 @@ +! This test validates that declare mapper for a derived type that extends +! a parent type with an allocatable component correctly maps the nested +! allocatable payload via the mapper when the whole object is mapped on +! target. + +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic + +program target_declare_mapper_parent_allocatable + implicit none + + type, abstract :: base_t + real, allocatable :: base_arr(:) + end type base_t + + type, extends(base_t) :: real_t + real, allocatable :: real_arr(:) + end type real_t + !$omp declare mapper(custommapper: real_t :: t) map(t%base_arr, t%real_arr) + + type(real_t) :: r + integer :: i + allocate(r%base_arr(10), source=1.0) + allocate(r%real_arr(10), source=1.0) + + !$omp target map(mapper(custommapper), tofrom: r) + do i = 1, size(r%base_arr) + r%base_arr(i) = 2.0 + r%real_arr(i) = 3.0 + r%real_arr(i) = r%base_arr(1) + end do + !$omp end target + + + !CHECK: base_arr: 2. 2. 2. 2. 2. 2. 2. 2. 2. 2. + print*, "base_arr: ", r%base_arr + !CHECK: real_arr: 2. 2. 2. 2. 2. 2. 2. 2. 2. 2. + print*, "real_arr: ", r%real_arr + + deallocate(r%real_arr) + deallocate(r%base_arr) +end program target_declare_mapper_parent_allocatable diff --git a/offload/test/offloading/fortran/target-no-loop.f90 b/offload/test/offloading/fortran/target-no-loop.f90 new file mode 100644 index 000000000000..3c88b00a5354 --- /dev/null +++ b/offload/test/offloading/fortran/target-no-loop.f90 @@ -0,0 +1,97 @@ +! REQUIRES: flang +! REQUIRES: gpu + +! RUN: %libomptarget-compile-fortran-generic -O3 -fopenmp-assume-threads-oversubscription -fopenmp-assume-teams-oversubscription +! RUN: env LIBOMPTARGET_INFO=16 OMP_NUM_TEAMS=16 OMP_TEAMS_THREAD_LIMIT=16 %libomptarget-run-generic 2>&1 | %fcheck-generic +function check_errors(array) result (errors) + integer, intent(in) :: array(1024) + integer :: errors + integer :: i + errors = 0 + do i = 1, 1024 + if ( array( i) .ne. (i) ) then + errors = errors + 1 + end if + end do +end function + +program main + use omp_lib + implicit none + integer :: i,j,red + integer :: array(1024), errors = 0 + array = 1 + + ! No-loop kernel + !$omp target teams distribute parallel do + do i = 1, 1024 + array(i) = i + end do + errors = errors + check_errors(array) + + ! SPMD kernel (num_teams clause blocks promotion to no-loop) + array = 1 + !$omp target teams distribute parallel do num_teams(3) + do i = 1, 1024 + array(i) = i + end do + + errors = errors + check_errors(array) + + ! No-loop kernel + array = 1 + !$omp target teams distribute parallel do num_threads(64) + do i = 1, 1024 + array(i) = i + end do + + errors = errors + check_errors(array) + + ! SPMD kernel + array = 1 + !$omp target parallel do + do i = 1, 1024 + array(i) = i + end do + + errors = errors + check_errors(array) + + ! Generic kernel + array = 1 + !$omp target teams distribute + do i = 1, 1024 + array(i) = i + end do + + errors = errors + check_errors(array) + + ! SPMD kernel (reduction clause blocks promotion to no-loop) + array = 1 + red =0 + !$omp target teams distribute parallel do reduction(+:red) + do i = 1, 1024 + red = red + array(i) + end do + + if (red .ne. 1024) then + errors = errors + 1 + end if + + print *,"number of errors: ", errors + +end program main + +! CHECK: "PluginInterface" device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode +! CHECK: info: #Args: 3 Teams x Thrds: 64x 16 +! CHECK: "PluginInterface" device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode +! CHECK: info: #Args: 3 Teams x Thrds: 3x 16 {{.*}} +! CHECK: "PluginInterface" device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode +! CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}} +! CHECK: "PluginInterface" device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode +! CHECK: info: #Args: 3 Teams x Thrds: 1x 16 +! CHECK: "PluginInterface" device {{[0-9]+}} info: Launching kernel {{.*}} Generic mode +! CHECK: info: #Args: 3 Teams x Thrds: 16x 16 {{.*}} +! CHECK: "PluginInterface" device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode +! CHECK: info: #Args: 4 Teams x Thrds: 16x 16 {{.*}} +! CHECK: number of errors: 0 + diff --git a/offload/test/offloading/interop.c b/offload/test/offloading/interop.c index 26287e3ec533..d9fa2ef883b9 100644 --- a/offload/test/offloading/interop.c +++ b/offload/test/offloading/interop.c @@ -1,5 +1,6 @@ // RUN: %libomptarget-compile-run-and-check-generic -// REQUIRES: nvptx64-nvidia-cuda + +// XFAIL: * #include <assert.h> #include <omp.h> diff --git a/offload/test/offloading/single_threaded_for_barrier_hang_1.c b/offload/test/offloading/single_threaded_for_barrier_hang_1.c index 8ee6b51fb681..a007521a5c74 100644 --- a/offload/test/offloading/single_threaded_for_barrier_hang_1.c +++ b/offload/test/offloading/single_threaded_for_barrier_hang_1.c @@ -1,6 +1,9 @@ // RUN: %libomptarget-compile-run-and-check-generic // RUN: %libomptarget-compileopt-run-and-check-generic +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// UNSUPPORTED: gpu + #include <omp.h> #include <stdio.h> diff --git a/offload/test/offloading/single_threaded_for_barrier_hang_2.c b/offload/test/offloading/single_threaded_for_barrier_hang_2.c index a98abd6922da..cabd2ed3dde7 100644 --- a/offload/test/offloading/single_threaded_for_barrier_hang_2.c +++ b/offload/test/offloading/single_threaded_for_barrier_hang_2.c @@ -1,6 +1,7 @@ // RUN: %libomptarget-compile-run-and-check-generic -// FIXME: This fails with optimization enabled and prints b: 0 -// FIXME: RUN: %libomptarget-compileopt-run-and-check-generic + +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// UNSUPPORTED: gpu #include <omp.h> #include <stdio.h> diff --git a/offload/test/offloading/spmdization.c b/offload/test/offloading/spmdization.c index 7f3f47d9ef32..48627cd7dae1 100644 --- a/offload/test/offloading/spmdization.c +++ b/offload/test/offloading/spmdization.c @@ -2,7 +2,8 @@ // RUN: %libomptarget-compileopt-generic // RUN: env LIBOMPTARGET_INFO=16 \ // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,SPMD -// RUN: %libomptarget-compileopt-generic -mllvm --openmp-opt-disable-spmdization +// RUN: %libomptarget-compileopt-generic -mllvm --openmp-opt-disable-spmdization \ +// RUN: -Xoffload-linker -mllvm=--openmp-opt-disable-spmdization // RUN: env LIBOMPTARGET_INFO=16 \ // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,GENERIC // clang-format on diff --git a/offload/test/sanitizer/ptr_outside_alloc_1.c b/offload/test/sanitizer/ptr_outside_alloc_1.c index bdd028352e40..b30ce12ef1ea 100644 --- a/offload/test/sanitizer/ptr_outside_alloc_1.c +++ b/offload/test/sanitizer/ptr_outside_alloc_1.c @@ -5,12 +5,10 @@ // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,TRACE // clang-format on -// UNSUPPORTED: aarch64-unknown-linux-gnu -// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO -// UNSUPPORTED: x86_64-unknown-linux-gnu -// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO -// UNSUPPORTED: s390x-ibm-linux-gnu -// UNSUPPORTED: s390x-ibm-linux-gnu-LTO +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// UNSUPPORTED: nvidiagpu +// +// REQUIRES: gpu #include <omp.h> diff --git a/offload/test/sanitizer/ptr_outside_alloc_2.c b/offload/test/sanitizer/ptr_outside_alloc_2.c index 6a67962f9eb3..3bb8bdaca8b4 100644 --- a/offload/test/sanitizer/ptr_outside_alloc_2.c +++ b/offload/test/sanitizer/ptr_outside_alloc_2.c @@ -3,12 +3,10 @@ // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK // clang-format on -// UNSUPPORTED: aarch64-unknown-linux-gnu -// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO -// UNSUPPORTED: x86_64-unknown-linux-gnu -// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO -// UNSUPPORTED: s390x-ibm-linux-gnu -// UNSUPPORTED: s390x-ibm-linux-gnu-LTO +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// UNSUPPORTED: nvidiagpu +// +// REQUIRES: gpu #include <omp.h> diff --git a/offload/test/sanitizer/use_after_free_1.c b/offload/test/sanitizer/use_after_free_1.c index c4783c5c36df..acc1de373f9e 100644 --- a/offload/test/sanitizer/use_after_free_1.c +++ b/offload/test/sanitizer/use_after_free_1.c @@ -5,12 +5,10 @@ // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,TRACE // clang-format on -// UNSUPPORTED: aarch64-unknown-linux-gnu -// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO -// UNSUPPORTED: x86_64-unknown-linux-gnu -// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO -// UNSUPPORTED: s390x-ibm-linux-gnu -// UNSUPPORTED: s390x-ibm-linux-gnu-LTO +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// UNSUPPORTED: nvidiagpu +// +// REQUIRES: gpu #include <omp.h> diff --git a/offload/test/sanitizer/use_after_free_2.c b/offload/test/sanitizer/use_after_free_2.c index 1c1e09744a75..3d70fb7b3a3f 100644 --- a/offload/test/sanitizer/use_after_free_2.c +++ b/offload/test/sanitizer/use_after_free_2.c @@ -3,12 +3,10 @@ // RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK // clang-format on -// UNSUPPORTED: aarch64-unknown-linux-gnu -// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO -// UNSUPPORTED: x86_64-unknown-linux-gnu -// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO -// UNSUPPORTED: s390x-ibm-linux-gnu -// UNSUPPORTED: s390x-ibm-linux-gnu-LTO +// FIXME: https://github.com/llvm/llvm-project/issues/161265 +// UNSUPPORTED: nvidiagpu +// +// REQUIRES: gpu // If offload memory pooling is enabled for a large allocation, reuse error is // not detected. UNSUPPORTED: large_allocation_memory_pool diff --git a/offload/tools/deviceinfo/llvm-offload-device-info.cpp b/offload/tools/deviceinfo/llvm-offload-device-info.cpp index a2955d49d396..9b58d67f017c 100644 --- a/offload/tools/deviceinfo/llvm-offload-device-info.cpp +++ b/offload/tools/deviceinfo/llvm-offload-device-info.cpp @@ -137,7 +137,7 @@ ol_result_t printDeviceValue(std::ostream &S, ol_device_handle_t Dev, size_t Size; OFFLOAD_ERR(olGetDeviceInfoSize(Dev, Info, &Size)); Val.resize(Size); - OFFLOAD_ERR(olGetDeviceInfo(Dev, Info, sizeof(Val), Val.data())); + OFFLOAD_ERR(olGetDeviceInfo(Dev, Info, Size, Val.data())); doWrite<T, PK>(S, reinterpret_cast<T>(Val.data())); } else { T Val; @@ -174,6 +174,8 @@ ol_result_t printDevice(std::ostream &S, ol_device_handle_t D) { OFFLOAD_ERR( printDeviceValue<const char *>(S, D, OL_DEVICE_INFO_NAME, "Name")); + OFFLOAD_ERR(printDeviceValue<const char *>(S, D, OL_DEVICE_INFO_PRODUCT_NAME, + "Product Name")); OFFLOAD_ERR( printDeviceValue<ol_device_type_t>(S, D, OL_DEVICE_INFO_TYPE, "Type")); OFFLOAD_ERR(printDeviceValue<const char *>( diff --git a/offload/unittests/OffloadAPI/CMakeLists.txt b/offload/unittests/OffloadAPI/CMakeLists.txt index b2d514423a6e..50c99a5d5b63 100644 --- a/offload/unittests/OffloadAPI/CMakeLists.txt +++ b/offload/unittests/OffloadAPI/CMakeLists.txt @@ -27,7 +27,9 @@ add_offload_unittest("memory" memory/olMemAlloc.cpp memory/olMemFill.cpp memory/olMemFree.cpp - memory/olMemcpy.cpp) + memory/olMemcpy.cpp + memory/olGetMemInfo.cpp + memory/olGetMemInfoSize.cpp) add_offload_unittest("platform" platform/olGetPlatformInfo.cpp @@ -35,6 +37,7 @@ add_offload_unittest("platform" add_offload_unittest("program" program/olCreateProgram.cpp + program/olIsValidBinary.cpp program/olDestroyProgram.cpp) add_offload_unittest("queue" diff --git a/offload/unittests/OffloadAPI/common/Fixtures.hpp b/offload/unittests/OffloadAPI/common/Fixtures.hpp index c5a35faba7a2..6f9961e2c6d5 100644 --- a/offload/unittests/OffloadAPI/common/Fixtures.hpp +++ b/offload/unittests/OffloadAPI/common/Fixtures.hpp @@ -40,12 +40,16 @@ } while (0) #endif -// TODO: rework this so the EXPECTED/ACTUAL results are readable #ifndef ASSERT_ERROR #define ASSERT_ERROR(EXPECTED, ACTUAL) \ do { \ ol_result_t Res = ACTUAL; \ - ASSERT_TRUE(Res && (Res->Code == EXPECTED)); \ + if (!Res) \ + GTEST_FAIL() << #ACTUAL " succeeded when we expected it to fail"; \ + if (Res->Code != EXPECTED) \ + GTEST_FAIL() << #ACTUAL " was expected to return " \ + << #EXPECTED " but instead returned " << Res->Code << ": " \ + << Res->Details; \ } while (0) #endif diff --git a/offload/unittests/OffloadAPI/device_code/CMakeLists.txt b/offload/unittests/OffloadAPI/device_code/CMakeLists.txt index 50e430597e64..1a042e1b3831 100644 --- a/offload/unittests/OffloadAPI/device_code/CMakeLists.txt +++ b/offload/unittests/OffloadAPI/device_code/CMakeLists.txt @@ -2,6 +2,7 @@ add_offload_test_device_code(foo.cpp foo) add_offload_test_device_code(bar.cpp bar) # Compile with optimizations to eliminate AMDGPU implicit arguments. add_offload_test_device_code(noargs.cpp noargs -O3) +add_offload_test_device_code(byte.cpp byte) add_offload_test_device_code(localmem.cpp localmem) add_offload_test_device_code(localmem_reduction.cpp localmem_reduction) add_offload_test_device_code(localmem_static.cpp localmem_static) @@ -14,6 +15,7 @@ add_custom_target(offload_device_binaries DEPENDS foo.bin bar.bin noargs.bin + byte.bin localmem.bin localmem_reduction.bin localmem_static.bin diff --git a/offload/unittests/OffloadAPI/device_code/byte.cpp b/offload/unittests/OffloadAPI/device_code/byte.cpp new file mode 100644 index 000000000000..779d120fefca --- /dev/null +++ b/offload/unittests/OffloadAPI/device_code/byte.cpp @@ -0,0 +1,3 @@ +#include <gpuintrin.h> + +extern "C" __gpu_kernel void byte(unsigned char c) { (void)c; } diff --git a/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp b/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp index 1dac8c50271b..c9eca36a4d44 100644 --- a/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp +++ b/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp @@ -55,6 +55,7 @@ struct LaunchSingleKernelTestBase : LaunchKernelTestBase { KERNEL_TEST(Foo, foo) KERNEL_TEST(NoArgs, noargs) +KERNEL_TEST(Byte, byte) KERNEL_TEST(LocalMem, localmem) KERNEL_TEST(LocalMemReduction, localmem_reduction) KERNEL_TEST(LocalMemStatic, localmem_static) diff --git a/offload/unittests/OffloadAPI/memory/olGetMemInfo.cpp b/offload/unittests/OffloadAPI/memory/olGetMemInfo.cpp new file mode 100644 index 000000000000..a4b382ff298a --- /dev/null +++ b/offload/unittests/OffloadAPI/memory/olGetMemInfo.cpp @@ -0,0 +1,130 @@ +//===------- Offload API tests - olGetMemInfo -----------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "../common/Fixtures.hpp" +#include <OffloadAPI.h> +#include <gtest/gtest.h> + +constexpr size_t SIZE = 1024; + +struct olGetMemInfoBaseTest : OffloadDeviceTest { + void *OffsetPtr() { return &reinterpret_cast<char *>(Ptr)[123]; } + + void *Ptr; +}; + +template <ol_alloc_type_t AllocType> +struct olGetMemInfoTest : olGetMemInfoBaseTest { + void SetUp() override { + RETURN_ON_FATAL_FAILURE(OffloadDeviceTest::SetUp()); + ASSERT_SUCCESS(olMemAlloc(Device, AllocType, SIZE, &Ptr)); + } + + void TearDown() override { + ASSERT_SUCCESS(olMemFree(Ptr)); + RETURN_ON_FATAL_FAILURE(OffloadDeviceTest::TearDown()); + } +}; +using olGetMemInfoDeviceTest = olGetMemInfoTest<OL_ALLOC_TYPE_DEVICE>; +OFFLOAD_TESTS_INSTANTIATE_DEVICE_FIXTURE(olGetMemInfoDeviceTest); +using olGetMemInfoManagedTest = olGetMemInfoTest<OL_ALLOC_TYPE_MANAGED>; +OFFLOAD_TESTS_INSTANTIATE_DEVICE_FIXTURE(olGetMemInfoManagedTest); +using olGetMemInfoHostTest = olGetMemInfoTest<OL_ALLOC_TYPE_HOST>; +OFFLOAD_TESTS_INSTANTIATE_DEVICE_FIXTURE(olGetMemInfoHostTest); + +#define PER_ALLOC_TEST(FUNCTION) \ + TEST_P(olGetMemInfoDeviceTest, FUNCTION) { \ + FUNCTION(this, Ptr, OL_ALLOC_TYPE_DEVICE); \ + } \ + TEST_P(olGetMemInfoManagedTest, FUNCTION) { \ + FUNCTION(this, Ptr, OL_ALLOC_TYPE_MANAGED); \ + } \ + TEST_P(olGetMemInfoHostTest, FUNCTION) { \ + FUNCTION(this, OffsetPtr(), OL_ALLOC_TYPE_HOST); \ + } \ + TEST_P(olGetMemInfoDeviceTest, FUNCTION##Offset) { \ + FUNCTION(this, Ptr, OL_ALLOC_TYPE_DEVICE); \ + } \ + TEST_P(olGetMemInfoManagedTest, FUNCTION##Offset) { \ + FUNCTION(this, OffsetPtr(), OL_ALLOC_TYPE_MANAGED); \ + } \ + TEST_P(olGetMemInfoHostTest, FUNCTION##Offset) { \ + FUNCTION(this, OffsetPtr(), OL_ALLOC_TYPE_HOST); \ + } + +void SuccessDevice(olGetMemInfoBaseTest *Fixture, void *Ptr, + ol_alloc_type_t Type) { + ol_device_handle_t RetrievedDevice; + ASSERT_SUCCESS(olGetMemInfo(Fixture->Ptr, OL_MEM_INFO_DEVICE, + sizeof(RetrievedDevice), &RetrievedDevice)); + ASSERT_EQ(RetrievedDevice, Fixture->Device); +} +PER_ALLOC_TEST(SuccessDevice); + +void SuccessBase(olGetMemInfoBaseTest *Fixture, void *Ptr, + ol_alloc_type_t Type) { + void *RetrievedBase; + ASSERT_SUCCESS(olGetMemInfo(Fixture->Ptr, OL_MEM_INFO_BASE, + sizeof(RetrievedBase), &RetrievedBase)); + ASSERT_EQ(RetrievedBase, Fixture->Ptr); +} +PER_ALLOC_TEST(SuccessBase); + +void SuccessSize(olGetMemInfoBaseTest *Fixture, void *Ptr, + ol_alloc_type_t Type) { + size_t RetrievedSize; + ASSERT_SUCCESS(olGetMemInfo(Fixture->Ptr, OL_MEM_INFO_SIZE, + sizeof(RetrievedSize), &RetrievedSize)); + ASSERT_EQ(RetrievedSize, SIZE); +} +PER_ALLOC_TEST(SuccessSize); + +void SuccessType(olGetMemInfoBaseTest *Fixture, void *Ptr, + ol_alloc_type_t Type) { + ol_alloc_type_t RetrievedType; + ASSERT_SUCCESS(olGetMemInfo(Fixture->Ptr, OL_MEM_INFO_TYPE, + sizeof(RetrievedType), &RetrievedType)); + ASSERT_EQ(RetrievedType, Type); +} +PER_ALLOC_TEST(SuccessType); + +TEST_P(olGetMemInfoDeviceTest, InvalidNotFound) { + // Assuming that we aren't unlucky and happen to get 0x1234 as a random + // pointer + void *RetrievedBase; + ASSERT_ERROR(OL_ERRC_NOT_FOUND, + olGetMemInfo(reinterpret_cast<void *>(0x1234), OL_MEM_INFO_BASE, + sizeof(RetrievedBase), &RetrievedBase)); +} + +TEST_P(olGetMemInfoDeviceTest, InvalidNullPtr) { + ol_device_handle_t RetrievedDevice; + ASSERT_ERROR(OL_ERRC_INVALID_NULL_POINTER, + olGetMemInfo(nullptr, OL_MEM_INFO_DEVICE, + sizeof(RetrievedDevice), &RetrievedDevice)); +} + +TEST_P(olGetMemInfoDeviceTest, InvalidSizeZero) { + ol_device_handle_t RetrievedDevice; + ASSERT_ERROR(OL_ERRC_INVALID_SIZE, + olGetMemInfo(Ptr, OL_MEM_INFO_DEVICE, 0, &RetrievedDevice)); +} + +TEST_P(olGetMemInfoDeviceTest, InvalidSizeSmall) { + ol_device_handle_t RetrievedDevice; + ASSERT_ERROR(OL_ERRC_INVALID_SIZE, + olGetMemInfo(Ptr, OL_MEM_INFO_DEVICE, + sizeof(RetrievedDevice) - 1, &RetrievedDevice)); +} + +TEST_P(olGetMemInfoDeviceTest, InvalidNullPointerPropValue) { + ol_device_handle_t RetrievedDevice; + ASSERT_ERROR( + OL_ERRC_INVALID_NULL_POINTER, + olGetMemInfo(Ptr, OL_MEM_INFO_DEVICE, sizeof(RetrievedDevice), nullptr)); +} diff --git a/offload/unittests/OffloadAPI/memory/olGetMemInfoSize.cpp b/offload/unittests/OffloadAPI/memory/olGetMemInfoSize.cpp new file mode 100644 index 000000000000..f1a1e790fb22 --- /dev/null +++ b/offload/unittests/OffloadAPI/memory/olGetMemInfoSize.cpp @@ -0,0 +1,63 @@ +//===------- Offload API tests - olGetMemInfoSize -------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include <OffloadAPI.h> + +#include "../common/Fixtures.hpp" + +struct olGetMemInfoSizeTest : OffloadDeviceTest { + void *OffsetPtr() { return &reinterpret_cast<char *>(Ptr)[123]; } + + void SetUp() override { + RETURN_ON_FATAL_FAILURE(OffloadDeviceTest::SetUp()); + ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_DEVICE, 0x1024, &Ptr)); + } + + void TearDown() override { + ASSERT_SUCCESS(olMemFree(Ptr)); + RETURN_ON_FATAL_FAILURE(OffloadDeviceTest::TearDown()); + } + + void *Ptr; +}; +OFFLOAD_TESTS_INSTANTIATE_DEVICE_FIXTURE(olGetMemInfoSizeTest); + +TEST_P(olGetMemInfoSizeTest, SuccessDevice) { + size_t Size = 0; + ASSERT_SUCCESS(olGetMemInfoSize(Ptr, OL_MEM_INFO_DEVICE, &Size)); + ASSERT_EQ(Size, sizeof(ol_device_handle_t)); +} + +TEST_P(olGetMemInfoSizeTest, SuccessBase) { + size_t Size = 0; + ASSERT_SUCCESS(olGetMemInfoSize(Ptr, OL_MEM_INFO_BASE, &Size)); + ASSERT_EQ(Size, sizeof(void *)); +} + +TEST_P(olGetMemInfoSizeTest, SuccessSize) { + size_t Size = 0; + ASSERT_SUCCESS(olGetMemInfoSize(Ptr, OL_MEM_INFO_SIZE, &Size)); + ASSERT_EQ(Size, sizeof(size_t)); +} + +TEST_P(olGetMemInfoSizeTest, SuccessType) { + size_t Size = 0; + ASSERT_SUCCESS(olGetMemInfoSize(Ptr, OL_MEM_INFO_TYPE, &Size)); + ASSERT_EQ(Size, sizeof(ol_alloc_type_t)); +} + +TEST_P(olGetMemInfoSizeTest, InvalidSymbolInfoEnumeration) { + size_t Size = 0; + ASSERT_ERROR(OL_ERRC_INVALID_ENUMERATION, + olGetMemInfoSize(Ptr, OL_MEM_INFO_FORCE_UINT32, &Size)); +} + +TEST_P(olGetMemInfoSizeTest, InvalidNullPointer) { + ASSERT_ERROR(OL_ERRC_INVALID_NULL_POINTER, + olGetMemInfoSize(Ptr, OL_MEM_INFO_DEVICE, nullptr)); +} diff --git a/offload/unittests/OffloadAPI/memory/olMemAlloc.cpp b/offload/unittests/OffloadAPI/memory/olMemAlloc.cpp index 00e428ec2abc..445262aa0c58 100644 --- a/offload/unittests/OffloadAPI/memory/olMemAlloc.cpp +++ b/offload/unittests/OffloadAPI/memory/olMemAlloc.cpp @@ -34,6 +34,26 @@ TEST_P(olMemAllocTest, SuccessAllocDevice) { olMemFree(Alloc); } +TEST_P(olMemAllocTest, SuccessAllocMany) { + std::vector<void *> Allocs; + Allocs.reserve(1000); + + constexpr ol_alloc_type_t TYPES[3] = { + OL_ALLOC_TYPE_DEVICE, OL_ALLOC_TYPE_MANAGED, OL_ALLOC_TYPE_HOST}; + + for (size_t I = 1; I < 1000; I++) { + void *Alloc = nullptr; + ASSERT_SUCCESS(olMemAlloc(Device, TYPES[I % 3], 1024 * I, &Alloc)); + ASSERT_NE(Alloc, nullptr); + + Allocs.push_back(Alloc); + } + + for (auto *A : Allocs) { + olMemFree(A); + } +} + TEST_P(olMemAllocTest, InvalidNullDevice) { void *Alloc = nullptr; ASSERT_ERROR(OL_ERRC_INVALID_NULL_HANDLE, diff --git a/offload/unittests/OffloadAPI/program/olIsValidBinary.cpp b/offload/unittests/OffloadAPI/program/olIsValidBinary.cpp new file mode 100644 index 000000000000..02e805dd1135 --- /dev/null +++ b/offload/unittests/OffloadAPI/program/olIsValidBinary.cpp @@ -0,0 +1,49 @@ +//===------- Offload API tests - olIsValidBinary --------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "../common/Fixtures.hpp" +#include <OffloadAPI.h> +#include <gtest/gtest.h> + +using olIsValidBinaryTest = OffloadDeviceTest; +OFFLOAD_TESTS_INSTANTIATE_DEVICE_FIXTURE(olIsValidBinaryTest); + +TEST_P(olIsValidBinaryTest, Success) { + + std::unique_ptr<llvm::MemoryBuffer> DeviceBin; + ASSERT_TRUE(TestEnvironment::loadDeviceBinary("foo", Device, DeviceBin)); + ASSERT_GE(DeviceBin->getBufferSize(), 0lu); + + bool IsValid = false; + ASSERT_SUCCESS(olIsValidBinary(Device, DeviceBin->getBufferStart(), + DeviceBin->getBufferSize(), &IsValid)); + ASSERT_TRUE(IsValid); + + ASSERT_SUCCESS( + olIsValidBinary(Device, DeviceBin->getBufferStart(), 0, &IsValid)); + ASSERT_FALSE(IsValid); +} + +TEST_P(olIsValidBinaryTest, Invalid) { + + std::unique_ptr<llvm::MemoryBuffer> DeviceBin; + ASSERT_TRUE(TestEnvironment::loadDeviceBinary("foo", Device, DeviceBin)); + ASSERT_GE(DeviceBin->getBufferSize(), 0lu); + + bool IsValid = false; + ASSERT_SUCCESS( + olIsValidBinary(Device, DeviceBin->getBufferStart(), 0, &IsValid)); + ASSERT_FALSE(IsValid); +} + +TEST_P(olIsValidBinaryTest, NullPointer) { + bool IsValid = false; + ASSERT_ERROR(OL_ERRC_INVALID_NULL_POINTER, + olIsValidBinary(Device, nullptr, 42, &IsValid)); + ASSERT_FALSE(IsValid); +} |