diff --git a/openmp/libomptarget/DeviceRTL/include/State.h b/openmp/libomptarget/DeviceRTL/include/State.h index c93de4191f83f..1a3490394458f 100644 --- a/openmp/libomptarget/DeviceRTL/include/State.h +++ b/openmp/libomptarget/DeviceRTL/include/State.h @@ -240,29 +240,29 @@ lookupPtr(ValueKind Kind, bool IsReadonly, bool ForceTeamState) { /// update ICV values we can declare in global scope. template struct Value { [[gnu::flatten, gnu::always_inline]] operator Ty() { - return lookup(/* IsReadonly */ true, /* IdentTy */ nullptr, - /* ForceTeamState */ false); + return lookup(/*IsReadonly=*/true, /*IdentTy=*/nullptr, + /*ForceTeamState=*/false); } [[gnu::flatten, gnu::always_inline]] Value &operator=(const Ty &Other) { - set(Other, /* IdentTy */ nullptr); + set(Other, /*IdentTy=*/nullptr); return *this; } [[gnu::flatten, gnu::always_inline]] Value &operator++() { - inc(1, /* IdentTy */ nullptr); + inc(1, /*IdentTy=*/nullptr); return *this; } [[gnu::flatten, gnu::always_inline]] Value &operator--() { - inc(-1, /* IdentTy */ nullptr); + inc(-1, /*IdentTy=*/nullptr); return *this; } [[gnu::flatten, gnu::always_inline]] void assert_eq(const Ty &V, IdentTy *Ident = nullptr, bool ForceTeamState = false) { - ASSERT(lookup(/* IsReadonly */ true, Ident, ForceTeamState) == V, nullptr); + ASSERT(lookup(/*IsReadonly=*/true, Ident, ForceTeamState) == V, nullptr); } private: @@ -273,12 +273,12 @@ template struct Value { } [[gnu::flatten, gnu::always_inline]] Ty &inc(int UpdateVal, IdentTy *Ident) { - return (lookup(/* IsReadonly */ false, Ident, /* ForceTeamState */ false) += + return (lookup(/*IsReadonly=*/false, Ident, /*ForceTeamState=*/false) += UpdateVal); } [[gnu::flatten, gnu::always_inline]] Ty &set(Ty UpdateVal, IdentTy *Ident) { - return (lookup(/* IsReadonly */ false, Ident, /* ForceTeamState */ false) = + return (lookup(/*IsReadonly=*/false, Ident, /*ForceTeamState=*/false) = UpdateVal); } @@ -290,8 +290,8 @@ template struct Value { /// we can declare in global scope. template struct PtrValue { [[gnu::flatten, gnu::always_inline]] operator Ty() { - return lookup(/* IsReadonly */ true, /* IdentTy */ nullptr, - /* ForceTeamState */ false); + return lookup(/*IsReadonly=*/true, /*IdentTy=*/nullptr, + /*ForceTeamState=*/false); } [[gnu::flatten, gnu::always_inline]] PtrValue &operator=(const Ty Other) { @@ -305,8 +305,8 @@ template struct PtrValue { } Ty &set(Ty UpdateVal) { - return (lookup(/* IsReadonly */ false, /* IdentTy */ nullptr, - /* ForceTeamState */ false) = UpdateVal); + return (lookup(/*IsReadonly=*/false, /*IdentTy=*/nullptr, + /*ForceTeamState=*/false) = UpdateVal); } template friend struct ValueRAII; @@ -315,7 +315,7 @@ template struct PtrValue { template struct ValueRAII { ValueRAII(VTy &V, Ty NewValue, Ty OldValue, bool Active, IdentTy *Ident, bool ForceTeamState = false) - : Ptr(Active ? &V.lookup(/* IsReadonly */ false, Ident, ForceTeamState) + : Ptr(Active ? &V.lookup(/*IsReadonly=*/false, Ident, ForceTeamState) : (Ty *)utils::UndefPtr), Val(OldValue), Active(Active) { if (!Active) diff --git a/openmp/libomptarget/DeviceRTL/src/Kernel.cpp b/openmp/libomptarget/DeviceRTL/src/Kernel.cpp index 06b12fec21677..95d4c728016d2 100644 --- a/openmp/libomptarget/DeviceRTL/src/Kernel.cpp +++ b/openmp/libomptarget/DeviceRTL/src/Kernel.cpp @@ -78,11 +78,11 @@ int32_t __kmpc_target_init(KernelEnvironmentTy &KernelEnvironment, llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD; bool UseGenericStateMachine = Configuration.UseGenericStateMachine; if (IsSPMD) { - inititializeRuntime(/* IsSPMD */ true, KernelEnvironment, + inititializeRuntime(/*IsSPMD=*/true, KernelEnvironment, KernelLaunchEnvironment); synchronize::threadsAligned(atomic::relaxed); } else { - inititializeRuntime(/* IsSPMD */ false, KernelEnvironment, + inititializeRuntime(/*IsSPMD=*/false, KernelEnvironment, KernelLaunchEnvironment); // No need to wait since only the main threads will execute user // code and workers will run into a barrier right away. diff --git a/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp b/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp index 2c0701bd5358f..7005477bf4c79 100644 --- a/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp +++ b/openmp/libomptarget/DeviceRTL/src/Parallelism.cpp @@ -121,20 +121,20 @@ __kmpc_parallel_51(IdentTy *ident, int32_t, int32_t if_expr, // created. state::ValueRAII ParallelTeamSizeRAII(state::ParallelTeamSize, PTeamSize, 1u, TId == 0, ident, - /* ForceTeamState */ true); + /*ForceTeamState=*/true); state::ValueRAII ActiveLevelRAII(icv::ActiveLevel, 1u, 0u, TId == 0, - ident, /* ForceTeamState */ true); + ident, /*ForceTeamState=*/true); state::ValueRAII LevelRAII(icv::Level, 1u, 0u, TId == 0, ident, - /* ForceTeamState */ true); + /*ForceTeamState=*/true); // Synchronize all threads after the main thread (TId == 0) set up the // team state properly. synchronize::threadsAligned(atomic::acq_rel); state::ParallelTeamSize.assert_eq(PTeamSize, ident, - /* ForceTeamState */ true); - icv::ActiveLevel.assert_eq(1u, ident, /* ForceTeamState */ true); - icv::Level.assert_eq(1u, ident, /* ForceTeamState */ true); + /*ForceTeamState=*/true); + icv::ActiveLevel.assert_eq(1u, ident, /*ForceTeamState=*/true); + icv::Level.assert_eq(1u, ident, /*ForceTeamState=*/true); // Ensure we synchronize before we run user code to avoid invalidating the // assumptions above. @@ -152,9 +152,9 @@ __kmpc_parallel_51(IdentTy *ident, int32_t, int32_t if_expr, // __kmpc_target_deinit may not hold. synchronize::threadsAligned(atomic::acq_rel); - state::ParallelTeamSize.assert_eq(1u, ident, /* ForceTeamState */ true); - icv::ActiveLevel.assert_eq(0u, ident, /* ForceTeamState */ true); - icv::Level.assert_eq(0u, ident, /* ForceTeamState */ true); + state::ParallelTeamSize.assert_eq(1u, ident, /*ForceTeamState=*/true); + icv::ActiveLevel.assert_eq(0u, ident, /*ForceTeamState=*/true); + icv::Level.assert_eq(0u, ident, /*ForceTeamState=*/true); // Ensure we synchronize to create an aligned region around the assumptions. synchronize::threadsAligned(atomic::relaxed); @@ -242,14 +242,14 @@ __kmpc_parallel_51(IdentTy *ident, int32_t, int32_t if_expr, // created. state::ValueRAII ParallelTeamSizeRAII(state::ParallelTeamSize, PTeamSize, 1u, true, ident, - /* ForceTeamState */ true); + /*ForceTeamState=*/true); state::ValueRAII ParallelRegionFnRAII(state::ParallelRegionFn, wrapper_fn, (void *)nullptr, true, ident, - /* ForceTeamState */ true); + /*ForceTeamState=*/true); state::ValueRAII ActiveLevelRAII(icv::ActiveLevel, 1u, 0u, true, ident, - /* ForceTeamState */ true); + /*ForceTeamState=*/true); state::ValueRAII LevelRAII(icv::Level, 1u, 0u, true, ident, - /* ForceTeamState */ true); + /*ForceTeamState=*/true); // Master signals work to activate workers. synchronize::threads(atomic::seq_cst); diff --git a/openmp/libomptarget/DeviceRTL/src/Reduction.cpp b/openmp/libomptarget/DeviceRTL/src/Reduction.cpp index 0a9db3edabccb..744d1a3a231c8 100644 --- a/openmp/libomptarget/DeviceRTL/src/Reduction.cpp +++ b/openmp/libomptarget/DeviceRTL/src/Reduction.cpp @@ -69,7 +69,7 @@ static int32_t nvptx_parallel_reduce_nowait(void *reduce_data, ShuffleReductFnTy shflFct, InterWarpCopyFnTy cpyFct) { uint32_t BlockThreadId = mapping::getThreadIdInBlock(); - if (mapping::isMainThreadInGenericMode(/* IsSPMD */ false)) + if (mapping::isMainThreadInGenericMode(/*IsSPMD=*/false)) BlockThreadId = 0; uint32_t NumThreads = omp_get_num_threads(); if (NumThreads == 1) diff --git a/openmp/libomptarget/include/Shared/Profile.h b/openmp/libomptarget/include/Shared/Profile.h index 7e580988a39ba..39817bab36cbb 100644 --- a/openmp/libomptarget/include/Shared/Profile.h +++ b/openmp/libomptarget/include/Shared/Profile.h @@ -32,7 +32,7 @@ class Profiler { Int32Envar ProfileGranularity = Int32Envar("LIBOMPTARGET_PROFILE_GRANULARITY", 500); - llvm::timeTraceProfilerInitialize(ProfileGranularity /* us */, + llvm::timeTraceProfilerInitialize(ProfileGranularity /*us=*/, "libomptarget"); } diff --git a/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp index b67642e9e1bcb..8424d0f5df082 100644 --- a/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp +++ b/openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp @@ -2211,10 +2211,9 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { hsa_amd_pointer_info_t Info; Info.size = sizeof(hsa_amd_pointer_info_t); - hsa_status_t Status = - hsa_amd_pointer_info(HstPtr, &Info, /* Allocator */ nullptr, - /* Number of accessible agents (out) */ nullptr, - /* Accessible agents */ nullptr); + hsa_status_t Status = hsa_amd_pointer_info( + HstPtr, &Info, /*Allocator=*/nullptr, /*num_agents_accessible=*/nullptr, + /*accessible=*/nullptr); if (auto Err = Plugin::check(Status, "Error in hsa_amd_pointer_info: %s")) return std::move(Err); @@ -2789,7 +2788,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { AMDHostDeviceTy &HostDevice; /// The current size of the global device memory pool (managed by us). - uint64_t DeviceMemoryPoolSize = 1L << 29L /* 512MB */; + uint64_t DeviceMemoryPoolSize = 1L << 29L /*512MB=*/; /// The current size of the stack that will be used in cases where it could /// not be statically determined. @@ -3031,9 +3030,8 @@ struct AMDGPUPluginTy final : public GenericPluginTy { /// Check whether the image is compatible with an AMDGPU device. Expected isELFCompatible(StringRef Image) const override { // Get the associated architecture and flags from the ELF. - auto ElfOrErr = - ELF64LEObjectFile::create(MemoryBufferRef(Image, /*Identifier=*/""), - /*InitContent=*/false); + auto ElfOrErr = ELF64LEObjectFile::create( + MemoryBufferRef(Image, /*Identifier=*/""), /*InitContent=*/false); if (!ElfOrErr) return ElfOrErr.takeError(); std::optional Processor = ElfOrErr->tryGetCPUName(); diff --git a/openmp/libomptarget/plugins-nextgen/common/include/GlobalHandler.h b/openmp/libomptarget/plugins-nextgen/common/include/GlobalHandler.h index d9fe938790ca7..8707e7b4c504e 100644 --- a/openmp/libomptarget/plugins-nextgen/common/include/GlobalHandler.h +++ b/openmp/libomptarget/plugins-nextgen/common/include/GlobalHandler.h @@ -138,7 +138,7 @@ class GenericGlobalHandlerTy { const GlobalTy &HostGlobal, const GlobalTy &DeviceGlobal) { return moveGlobalBetweenDeviceAndHost(Device, HostGlobal, DeviceGlobal, - /* D2H */ true); + /*D2H=*/true); } /// Copy the memory associated with a global from the device to its @@ -147,7 +147,7 @@ class GenericGlobalHandlerTy { Error readGlobalFromDevice(GenericDeviceTy &Device, DeviceImageTy &Image, const GlobalTy &HostGlobal) { return moveGlobalBetweenDeviceAndHost(Device, Image, HostGlobal, - /* D2H */ true); + /*D2H=*/true); } /// Copy the memory associated with a global from the host to its counterpart @@ -156,7 +156,7 @@ class GenericGlobalHandlerTy { Error writeGlobalToDevice(GenericDeviceTy &Device, const GlobalTy &HostGlobal, const GlobalTy &DeviceGlobal) { return moveGlobalBetweenDeviceAndHost(Device, HostGlobal, DeviceGlobal, - /* D2H */ false); + /*D2H=*/false); } /// Copy the memory associated with a global from the host to its counterpart @@ -165,7 +165,7 @@ class GenericGlobalHandlerTy { Error writeGlobalToDevice(GenericDeviceTy &Device, DeviceImageTy &Image, const GlobalTy &HostGlobal) { return moveGlobalBetweenDeviceAndHost(Device, Image, HostGlobal, - /* D2H */ false); + /*D2H=*/false); } }; diff --git a/openmp/libomptarget/plugins-nextgen/common/src/JIT.cpp b/openmp/libomptarget/plugins-nextgen/common/src/JIT.cpp index 7275be4edfca5..9eb610cab4de6 100644 --- a/openmp/libomptarget/plugins-nextgen/common/src/JIT.cpp +++ b/openmp/libomptarget/plugins-nextgen/common/src/JIT.cpp @@ -93,7 +93,7 @@ createModuleFromImage(const __tgt_device_image &Image, LLVMContext &Context) { StringRef Data((const char *)Image.ImageStart, target::getPtrDiff(Image.ImageEnd, Image.ImageStart)); std::unique_ptr MB = MemoryBuffer::getMemBuffer( - Data, /* BufferName */ "", /* RequiresNullTerminator */ false); + Data, /*BufferName=*/"", /*RequiresNullTerminator=*/false); return createModuleFromMemoryBuffer(MB, Context); } @@ -186,7 +186,7 @@ void JITEngine::codegen(TargetMachine *TM, TargetLibraryInfoImpl *TLII, TM->addPassesToEmitFile(PM, OS, nullptr, TT.isNVPTX() ? CodeGenFileType::AssemblyFile : CodeGenFileType::ObjectFile, - /* DisableVerify */ false, MMIWP); + /*DisableVerify=*/false, MMIWP); PM.run(M); } @@ -196,8 +196,8 @@ JITEngine::backend(Module &M, const std::string &ComputeUnitKind, unsigned OptLevel) { auto RemarksFileOrErr = setupLLVMOptimizationRemarks( - M.getContext(), /* RemarksFilename */ "", /* RemarksPasses */ "", - /* RemarksFormat */ "", /* RemarksWithHotness */ false); + M.getContext(), /*RemarksFilename=*/"", /*RemarksPasses=*/"", + /*RemarksFormat=*/"", /*RemarksWithHotness=*/false); if (Error E = RemarksFileOrErr.takeError()) return std::move(E); if (*RemarksFileOrErr) diff --git a/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp b/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp index 9490e58fc669c..0db7910ec1056 100644 --- a/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp +++ b/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp @@ -64,7 +64,7 @@ struct RecordReplayTy { void *suggestAddress(uint64_t MaxMemoryAllocation) { // Get a valid pointer address for this system void *Addr = - Device->allocate(1024, /* HstPtr */ nullptr, TARGET_ALLOC_DEFAULT); + Device->allocate(1024, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT); Device->free(Addr); // Align Address to MaxMemoryAllocation Addr = (void *)alignPtr((Addr), MaxMemoryAllocation); @@ -104,8 +104,8 @@ struct RecordReplayTy { constexpr size_t STEP = 1024 * 1024 * 1024ULL; MemoryStart = nullptr; for (TotalSize = MAX_MEMORY_ALLOCATION; TotalSize > 0; TotalSize -= STEP) { - MemoryStart = Device->allocate(TotalSize, /* HstPtr */ nullptr, - TARGET_ALLOC_DEFAULT); + MemoryStart = + Device->allocate(TotalSize, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT); if (MemoryStart) break; } @@ -214,8 +214,8 @@ struct RecordReplayTy { for (auto &OffloadEntry : Image.getOffloadEntryTable()) { if (!OffloadEntry.size) continue; - Size += std::strlen(OffloadEntry.name) + /* '\0' */ 1 + - /* OffloadEntry.size value */ sizeof(uint32_t) + + // Get the total size of the string and entry including the null byte. + Size += std::strlen(OffloadEntry.name) + 1 + sizeof(uint32_t) + OffloadEntry.size; } @@ -735,13 +735,12 @@ Error GenericDeviceTy::init(GenericPluginTy &Plugin) { if (ompt::Initialized) { bool ExpectedStatus = false; if (OmptInitialized.compare_exchange_strong(ExpectedStatus, true)) - performOmptCallback(device_initialize, - /* device_num */ DeviceId + - Plugin.getDeviceIdStartIndex(), - /* type */ getComputeUnitKind().c_str(), - /* device */ reinterpret_cast(this), - /* lookup */ ompt::lookupCallbackByName, - /* documentation */ nullptr); + performOmptCallback(device_initialize, /*device_num=*/DeviceId + + Plugin.getDeviceIdStartIndex(), + /*type=*/getComputeUnitKind().c_str(), + /*device=*/reinterpret_cast(this), + /*lookup=*/ompt::lookupCallbackByName, + /*documentation=*/nullptr); } #endif @@ -835,7 +834,7 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) { bool ExpectedStatus = true; if (OmptInitialized.compare_exchange_strong(ExpectedStatus, false)) performOmptCallback(device_finalize, - /* device_num */ DeviceId + + /*device_num=*/DeviceId + Plugin.getDeviceIdStartIndex()); } #endif @@ -897,16 +896,11 @@ GenericDeviceTy::loadBinary(GenericPluginTy &Plugin, if (ompt::Initialized) { size_t Bytes = getPtrDiff(InputTgtImage->ImageEnd, InputTgtImage->ImageStart); - performOmptCallback(device_load, - /* device_num */ DeviceId + - Plugin.getDeviceIdStartIndex(), - /* FileName */ nullptr, - /* File Offset */ 0, - /* VmaInFile */ nullptr, - /* ImgSize */ Bytes, - /* HostAddr */ InputTgtImage->ImageStart, - /* DeviceAddr */ nullptr, - /* FIXME: ModuleId */ 0); + performOmptCallback( + device_load, /*device_num=*/DeviceId + Plugin.getDeviceIdStartIndex(), + /*FileName=*/nullptr, /*FileOffset=*/0, /*VmaInFile=*/nullptr, + /*ImgSize=*/Bytes, /*HostAddr=*/InputTgtImage->ImageStart, + /*DeviceAddr=*/nullptr, /* FIXME: ModuleId */ 0); } #endif @@ -1293,7 +1287,7 @@ Error PinnedAllocationMapTy::lockMappedHostBuffer(void *HstPtr, size_t Size) { // If pinned, just insert the entry representing the whole pinned buffer. if (*IsPinnedOrErr) return insertEntry(BaseHstPtr, BaseDevAccessiblePtr, BaseSize, - /* Externally locked */ true); + /*Externallylocked=*/true); // Not externally pinned. Do nothing if locking of mapped buffers is disabled. if (!LockMappedBuffers) @@ -1863,7 +1857,7 @@ int32_t __tgt_rtl_data_notify_unmapped(int32_t DeviceId, void *HstPtr) { int32_t __tgt_rtl_data_submit(int32_t DeviceId, void *TgtPtr, void *HstPtr, int64_t Size) { return __tgt_rtl_data_submit_async(DeviceId, TgtPtr, HstPtr, Size, - /* AsyncInfoPtr */ nullptr); + /*AsyncInfoPtr=*/nullptr); } int32_t __tgt_rtl_data_submit_async(int32_t DeviceId, void *TgtPtr, @@ -1885,7 +1879,7 @@ int32_t __tgt_rtl_data_submit_async(int32_t DeviceId, void *TgtPtr, int32_t __tgt_rtl_data_retrieve(int32_t DeviceId, void *HstPtr, void *TgtPtr, int64_t Size) { return __tgt_rtl_data_retrieve_async(DeviceId, HstPtr, TgtPtr, Size, - /* AsyncInfoPtr */ nullptr); + /*AsyncInfoPtr=*/nullptr); } int32_t __tgt_rtl_data_retrieve_async(int32_t DeviceId, void *HstPtr, @@ -1909,7 +1903,7 @@ int32_t __tgt_rtl_data_exchange(int32_t SrcDeviceId, void *SrcPtr, int64_t Size) { return __tgt_rtl_data_exchange_async(SrcDeviceId, SrcPtr, DstDeviceId, DstPtr, Size, - /* AsyncInfoPtr */ nullptr); + /*AsyncInfoPtr=*/nullptr); } int32_t __tgt_rtl_data_exchange_async(int32_t SrcDeviceId, void *SrcPtr, diff --git a/openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp b/openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp index 0005bff7a8035..fb51f96c07b74 100644 --- a/openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp +++ b/openmp/libomptarget/plugins-nextgen/cuda/src/rtl.cpp @@ -1216,10 +1216,10 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice, std::max(KernelArgs.DynCGroupMem, GenericDevice.getDynamicMemorySize()); CUresult Res = - cuLaunchKernel(Func, NumBlocks, /* gridDimY */ 1, - /* gridDimZ */ 1, NumThreads, - /* blockDimY */ 1, /* blockDimZ */ 1, MaxDynCGroupMem, - Stream, (void **)Args, nullptr); + cuLaunchKernel(Func, NumBlocks, /*gridDimY=*/1, + /*gridDimZ=*/1, NumThreads, + /*blockDimY=*/1, /*blockDimZ=*/1, MaxDynCGroupMem, Stream, + (void **)Args, nullptr); return Plugin::check(Res, "Error in cuLaunchKernel for '%s': %s", getName()); } diff --git a/openmp/libomptarget/plugins-nextgen/generic-elf-64bit/src/rtl.cpp b/openmp/libomptarget/plugins-nextgen/generic-elf-64bit/src/rtl.cpp index 7f66b6827bce0..6466afc543b56 100644 --- a/openmp/libomptarget/plugins-nextgen/generic-elf-64bit/src/rtl.cpp +++ b/openmp/libomptarget/plugins-nextgen/generic-elf-64bit/src/rtl.cpp @@ -73,8 +73,8 @@ struct GenELF64KernelTy : public GenericKernelTy { Func = (void (*)())Global.getPtr(); KernelEnvironment.Configuration.ExecMode = OMP_TGT_EXEC_MODE_GENERIC; - KernelEnvironment.Configuration.MayUseNestedParallelism = /* Unknown */ 2; - KernelEnvironment.Configuration.UseGenericStateMachine = /* Unknown */ 2; + KernelEnvironment.Configuration.MayUseNestedParallelism = /*Unknown=*/2; + KernelEnvironment.Configuration.UseGenericStateMachine = /*Unknown=*/2; // Set the maximum number of threads to a single. MaxNumThreads = 1; diff --git a/openmp/libomptarget/src/OpenMP/Mapping.cpp b/openmp/libomptarget/src/OpenMP/Mapping.cpp index a5c24810e0af9..833856f2abf29 100644 --- a/openmp/libomptarget/src/OpenMP/Mapping.cpp +++ b/openmp/libomptarget/src/OpenMP/Mapping.cpp @@ -303,9 +303,9 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer( // Notify the plugin about the new mapping. if (Device.notifyDataMapped(HstPtrBegin, Size)) - return {{false /* IsNewEntry */, false /* IsHostPointer */}, - nullptr /* Entry */, - nullptr /* TargetPointer */}; + return {{false /*IsNewEntry=*/, false /*IsHostPointer=*/}, + nullptr /*Entry=*/, + nullptr /*TargetPointer=*/}; } else { // This entry is not present and we did not create a new entry for it. LR.TPR.Flags.IsPresent = false; @@ -333,9 +333,9 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer( LR.TPR.TargetPointer = nullptr; } else if (LR.TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) != OFFLOAD_SUCCESS) - return {{false /* IsNewEntry */, false /* IsHostPointer */}, - nullptr /* Entry */, - nullptr /* TargetPointer */}; + return {{false /*IsNewEntry=*/, false /*IsHostPointer=*/}, + nullptr /*Entry=*/, + nullptr /*TargetPointer=*/}; } else { // If not a host pointer and no present modifier, we need to wait for the // event if it exists. @@ -349,9 +349,9 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer( // If it fails to wait for the event, we need to return nullptr in // case of any data race. REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event)); - return {{false /* IsNewEntry */, false /* IsHostPointer */}, - nullptr /* Entry */, - nullptr /* TargetPointer */}; + return {{false /*IsNewEntry=*/, false /*IsHostPointer=*/}, + nullptr /*Entry=*/, + nullptr /*TargetPointer=*/}; } } } diff --git a/openmp/libomptarget/src/OpenMP/OMPT/Callback.cpp b/openmp/libomptarget/src/OpenMP/OMPT/Callback.cpp index da955e101956f..82934f10486c5 100644 --- a/openmp/libomptarget/src/OpenMP/OMPT/Callback.cpp +++ b/openmp/libomptarget/src/OpenMP/OMPT/Callback.cpp @@ -88,16 +88,16 @@ void Interface::beginTargetDataAlloc(int64_t DeviceId, void *HstPtrBegin, ompt_callback_target_data_op_emi_fn( ompt_scope_begin, TargetTaskData, &TargetData, &HostOpId, ompt_target_data_alloc, HstPtrBegin, - /* SrcDeviceNum */ omp_get_initial_device(), *TgtPtrBegin, - /* TgtDeviceNum */ DeviceId, Size, Code); + /*SrcDeviceNum=*/omp_get_initial_device(), *TgtPtrBegin, + /*TgtDeviceNum=*/DeviceId, Size, Code); } else if (ompt_callback_target_data_op_fn) { // HostOpId is set by the runtime HostOpId = createOpId(); // Invoke the tool supplied data op callback ompt_callback_target_data_op_fn( TargetData.value, HostOpId, ompt_target_data_alloc, HstPtrBegin, - /* SrcDeviceNum */ omp_get_initial_device(), *TgtPtrBegin, - /* TgtDeviceNum */ DeviceId, Size, Code); + /*SrcDeviceNum=*/omp_get_initial_device(), *TgtPtrBegin, + /*TgtDeviceNum=*/DeviceId, Size, Code); } } @@ -111,8 +111,8 @@ void Interface::endTargetDataAlloc(int64_t DeviceId, void *HstPtrBegin, ompt_callback_target_data_op_emi_fn( ompt_scope_end, TargetTaskData, &TargetData, &HostOpId, ompt_target_data_alloc, HstPtrBegin, - /* SrcDeviceNum */ omp_get_initial_device(), *TgtPtrBegin, - /* TgtDeviceNum */ DeviceId, Size, Code); + /*SrcDeviceNum=*/omp_get_initial_device(), *TgtPtrBegin, + /*TgtDeviceNum=*/DeviceId, Size, Code); } endTargetDataOperation(); } @@ -127,15 +127,15 @@ void Interface::beginTargetDataSubmit(int64_t DeviceId, void *TgtPtrBegin, ompt_callback_target_data_op_emi_fn( ompt_scope_begin, TargetTaskData, &TargetData, &HostOpId, ompt_target_data_transfer_to_device, HstPtrBegin, - /* SrcDeviceNum */ omp_get_initial_device(), TgtPtrBegin, DeviceId, - Size, Code); + /*SrcDeviceNum=*/omp_get_initial_device(), TgtPtrBegin, DeviceId, Size, + Code); } else if (ompt_callback_target_data_op_fn) { // HostOpId is set by the runtime HostOpId = createOpId(); // Invoke the tool supplied data op callback ompt_callback_target_data_op_fn( TargetData.value, HostOpId, ompt_target_data_transfer_to_device, - HstPtrBegin, /* SrcDeviceNum */ omp_get_initial_device(), TgtPtrBegin, + HstPtrBegin, /*SrcDeviceNum=*/omp_get_initial_device(), TgtPtrBegin, DeviceId, Size, Code); } } @@ -150,8 +150,8 @@ void Interface::endTargetDataSubmit(int64_t DeviceId, void *TgtPtrBegin, ompt_callback_target_data_op_emi_fn( ompt_scope_end, TargetTaskData, &TargetData, &HostOpId, ompt_target_data_transfer_to_device, HstPtrBegin, - /* SrcDeviceNum */ omp_get_initial_device(), TgtPtrBegin, DeviceId, - Size, Code); + /*SrcDeviceNum=*/omp_get_initial_device(), TgtPtrBegin, DeviceId, Size, + Code); } endTargetDataOperation(); } @@ -165,15 +165,15 @@ void Interface::beginTargetDataDelete(int64_t DeviceId, void *TgtPtrBegin, ompt_callback_target_data_op_emi_fn( ompt_scope_begin, TargetTaskData, &TargetData, &HostOpId, ompt_target_data_delete, TgtPtrBegin, DeviceId, - /* TgtPtrBegin */ nullptr, /* TgtDeviceNum */ -1, /* Bytes */ 0, Code); + /*TgtPtrBegin=*/nullptr, /*TgtDeviceNum=*/-1, /*Bytes=*/0, Code); } else if (ompt_callback_target_data_op_fn) { // HostOpId is set by the runtime HostOpId = createOpId(); // Invoke the tool supplied data op callback ompt_callback_target_data_op_fn(TargetData.value, HostOpId, ompt_target_data_delete, TgtPtrBegin, - DeviceId, /* TgtPtrBegin */ nullptr, - /* TgtDeviceNum */ -1, /* Bytes */ 0, Code); + DeviceId, /*TgtPtrBegin=*/nullptr, + /*TgtDeviceNum=*/-1, /*Bytes=*/0, Code); } } @@ -186,7 +186,7 @@ void Interface::endTargetDataDelete(int64_t DeviceId, void *TgtPtrBegin, ompt_callback_target_data_op_emi_fn( ompt_scope_end, TargetTaskData, &TargetData, &HostOpId, ompt_target_data_delete, TgtPtrBegin, DeviceId, - /* TgtPtrBegin */ nullptr, /* TgtDeviceNum */ -1, /* Bytes */ 0, Code); + /*TgtPtrBegin=*/nullptr, /*TgtDeviceNum=*/-1, /*Bytes=*/0, Code); } endTargetDataOperation(); } @@ -202,7 +202,7 @@ void Interface::beginTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin, ompt_scope_begin, TargetTaskData, &TargetData, &HostOpId, ompt_target_data_transfer_from_device, TgtPtrBegin, DeviceId, HstPtrBegin, - /* TgtDeviceNum */ omp_get_initial_device(), Size, Code); + /*TgtDeviceNum=*/omp_get_initial_device(), Size, Code); } else if (ompt_callback_target_data_op_fn) { // HostOpId is set by the runtime HostOpId = createOpId(); @@ -210,7 +210,7 @@ void Interface::beginTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin, ompt_callback_target_data_op_fn( TargetData.value, HostOpId, ompt_target_data_transfer_from_device, TgtPtrBegin, DeviceId, HstPtrBegin, - /* TgtDeviceNum */ omp_get_initial_device(), Size, Code); + /*TgtDeviceNum=*/omp_get_initial_device(), Size, Code); } } @@ -225,7 +225,7 @@ void Interface::endTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin, ompt_scope_end, TargetTaskData, &TargetData, &HostOpId, ompt_target_data_transfer_from_device, TgtPtrBegin, DeviceId, HstPtrBegin, - /* TgtDeviceNum */ omp_get_initial_device(), Size, Code); + /*TgtDeviceNum=*/omp_get_initial_device(), Size, Code); } endTargetDataOperation(); } diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp index fa8932361a51c..654efd524024a 100644 --- a/openmp/libomptarget/src/device.cpp +++ b/openmp/libomptarget/src/device.cpp @@ -117,7 +117,7 @@ void *DeviceTy::allocData(int64_t Size, void *HstPtr, int32_t Kind) { OMPT_IF_BUILT(InterfaceRAII TargetDataAllocRAII( RegionInterface.getCallbacks(), DeviceID, HstPtr, &TargetPtr, Size, - /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) + /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));) TargetPtr = RTL->data_alloc(RTLDeviceID, Size, HstPtr, Kind); return TargetPtr; @@ -128,7 +128,7 @@ int32_t DeviceTy::deleteData(void *TgtAllocBegin, int32_t Kind) { OMPT_IF_BUILT(InterfaceRAII TargetDataDeleteRAII( RegionInterface.getCallbacks(), DeviceID, TgtAllocBegin, - /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) + /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));) return RTL->data_delete(RTLDeviceID, TgtAllocBegin, Kind); } @@ -146,7 +146,7 @@ int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size, InterfaceRAII TargetDataSubmitRAII( RegionInterface.getCallbacks(), DeviceID, TgtPtrBegin, HstPtrBegin, Size, - /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) + /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));) if (!AsyncInfo || !RTL->data_submit_async || !RTL->synchronize) return RTL->data_submit(RTLDeviceID, TgtPtrBegin, HstPtrBegin, Size); @@ -168,7 +168,7 @@ int32_t DeviceTy::retrieveData(void *HstPtrBegin, void *TgtPtrBegin, InterfaceRAII TargetDataRetrieveRAII( RegionInterface.getCallbacks(), DeviceID, HstPtrBegin, TgtPtrBegin, Size, - /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) + /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));) if (!RTL->data_retrieve_async || !RTL->synchronize) return RTL->data_retrieve(RTLDeviceID, HstPtrBegin, TgtPtrBegin, Size); diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp index 61d9db17f5100..49495ac266f1b 100644 --- a/openmp/libomptarget/src/interface.cpp +++ b/openmp/libomptarget/src/interface.cpp @@ -113,7 +113,7 @@ targetData(ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase, int Rc = OFFLOAD_SUCCESS; Rc = TargetDataFunction(Loc, *DeviceOrErr, ArgNum, ArgsBase, Args, ArgSizes, ArgTypes, ArgNames, ArgMappers, AsyncInfo, - false /* FromMapper */); + false /*FromMapper=*/); if (Rc == OFFLOAD_SUCCESS) Rc = AsyncInfo.synchronize(); @@ -293,7 +293,7 @@ static inline int targetKernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams, /// RAII to establish tool anchors before and after target region OMPT_IF_BUILT(InterfaceRAII TargetRAII( RegionInterface.getCallbacks(), DeviceId, - /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) + /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));) int Rc = OFFLOAD_SUCCESS; Rc = target(Loc, *DeviceOrErr, HostPtr, *KernelArgs, AsyncInfo); @@ -387,7 +387,7 @@ EXTERN int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId, /// RAII to establish tool anchors before and after target region OMPT_IF_BUILT(InterfaceRAII TargetRAII( RegionInterface.getCallbacks(), DeviceId, - /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) + /*CodePtr=*/OMPT_GET_RETURN_ADDRESS(0));) AsyncInfoTy AsyncInfo(*DeviceOrErr); int Rc = target_replay(Loc, *DeviceOrErr, HostPtr, DeviceMemory, diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index a7d55d7ebd539..ed34870e12d01 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -225,7 +225,7 @@ static int initLibrary(DeviceTy &Device) { AsyncInfoTy AsyncInfo(Device); void *DevPtr; Device.retrieveData(&DevPtr, CurrDeviceEntryAddr, sizeof(void *), - AsyncInfo, /* Entry */ nullptr, &HDTTMap); + AsyncInfo, /*Entry=*/nullptr, &HDTTMap); if (AsyncInfo.synchronize() != OFFLOAD_SUCCESS) return OFFLOAD_FAIL; CurrDeviceEntryAddr = DevPtr; @@ -617,7 +617,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, /*HstPtrName=*/nullptr, /*HasFlagTo=*/false, /*HasFlagAlways=*/false, IsImplicit, UpdateRef, HasCloseModifier, HasPresentModifier, HasHoldModifier, AsyncInfo, - /* OwnedTPR */ nullptr, /* ReleaseHDTTMap */ false); + /*OwnedTPR=*/nullptr, /*ReleaseHDTTMap=*/false); PointerTgtPtrBegin = PointerTpr.TargetPointer; IsHostPtr = PointerTpr.Flags.IsHostPointer; if (!PointerTgtPtrBegin) { @@ -1723,7 +1723,7 @@ int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr, DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n", TargetTable->EntriesBegin[TM->Index].name, DPxPTR(TgtEntryPtr), TM->Index); - void *TgtPtr = Device.allocData(DeviceMemorySize, /* HstPtr */ nullptr, + void *TgtPtr = Device.allocData(DeviceMemorySize, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT); Device.submitData(TgtPtr, DeviceMemory, DeviceMemorySize, AsyncInfo); diff --git a/openmp/libomptarget/tools/kernelreplay/llvm-omp-kernel-replay.cpp b/openmp/libomptarget/tools/kernelreplay/llvm-omp-kernel-replay.cpp index fc3ff078a81fb..761e04e4c7bbd 100644 --- a/openmp/libomptarget/tools/kernelreplay/llvm-omp-kernel-replay.cpp +++ b/openmp/libomptarget/tools/kernelreplay/llvm-omp-kernel-replay.cpp @@ -57,8 +57,8 @@ int main(int argc, char **argv) { cl::ParseCommandLineOptions(argc, argv, "llvm-omp-kernel-replay\n"); ErrorOr> KernelInfoMB = - MemoryBuffer::getFile(InputFilename, /* isText */ true, - /* RequiresNullTerminator */ true); + MemoryBuffer::getFile(InputFilename, /*isText=*/true, + /*RequiresNullTerminator=*/true); if (!KernelInfoMB) report_fatal_error("Error reading the kernel info json file"); Expected JsonKernelInfo = @@ -100,8 +100,8 @@ int main(int argc, char **argv) { KernelEntry.addr = (void *)0x1; ErrorOr> ImageMB = - MemoryBuffer::getFile(KernelEntryName + ".image", /* isText */ false, - /* RequiresNullTerminator */ false); + MemoryBuffer::getFile(KernelEntryName + ".image", /*isText=*/false, + /*RequiresNullTerminator=*/false); if (!ImageMB) report_fatal_error("Error reading the kernel image."); @@ -127,7 +127,7 @@ int main(int argc, char **argv) { int32_t DeviceId = (DeviceIdOpt > -1 ? DeviceIdOpt : DeviceIdJson.value()); // TODO: do we need requires? - //__tgt_register_requires(/* Flags */1); + //__tgt_register_requires(/*Flags=*/1); __tgt_register_lib(&Desc); @@ -140,8 +140,8 @@ int main(int argc, char **argv) { } ErrorOr> DeviceMemoryMB = - MemoryBuffer::getFile(KernelEntryName + ".memory", /* isText */ false, - /* RequiresNullTerminator */ false); + MemoryBuffer::getFile(KernelEntryName + ".memory", /*isText=*/false, + /*RequiresNullTerminator=*/false); if (!DeviceMemoryMB) report_fatal_error("Error reading the kernel input device memory."); @@ -166,7 +166,7 @@ int main(int argc, char **argv) { } __tgt_target_kernel_replay( - /* Loc */ nullptr, DeviceId, KernelEntry.addr, (char *)recored_data, + /*Loc=*/nullptr, DeviceId, KernelEntry.addr, (char *)recored_data, DeviceMemoryMB.get()->getBufferSize(), TgtArgs.data(), TgtArgOffsets.data(), NumArgs.value(), NumTeams, NumThreads, LoopTripCount.value()); @@ -174,15 +174,15 @@ int main(int argc, char **argv) { if (VerifyOpt) { ErrorOr> OriginalOutputMB = MemoryBuffer::getFile(KernelEntryName + ".original.output", - /* isText */ false, - /* RequiresNullTerminator */ false); + /*isText=*/false, + /*RequiresNullTerminator=*/false); if (!OriginalOutputMB) report_fatal_error("Error reading the kernel original output file, make " "sure LIBOMPTARGET_SAVE_OUTPUT is set when recording"); ErrorOr> ReplayOutputMB = MemoryBuffer::getFile(KernelEntryName + ".replay.output", - /* isText */ false, - /* RequiresNullTerminator */ false); + /*isText=*/false, + /*RequiresNullTerminator=*/false); if (!ReplayOutputMB) report_fatal_error("Error reading the kernel replay output file");