[Offload][NFC] Fix typos discovered by codespell

offload/DeviceRTL/include/Configuration.h

@@ -27,7 +27,7 @@ uint32_t getNumDevices();
 /// Return the device number in the system for omp_get_device_num.
 uint32_t getDeviceNum();
 
-/// Return the user choosen debug level.
+/// Return the user chosen debug level.
 uint32_t getDebugKind();
 
 /// Return if teams oversubscription is assumed

offload/DeviceRTL/include/Mapping.h

@@ -63,7 +63,7 @@ LaneMaskTy activemask();
 /// Return a mask describing all threads with a smaller Id in the warp.
 LaneMaskTy lanemaskLT();
 
-/// Return a mask describing all threads with a larget Id in the warp.
+/// Return a mask describing all threads with a larger Id in the warp.
 LaneMaskTy lanemaskGT();
 
 /// Return the thread Id in the warp, in [0, getWarpSize()).

offload/DeviceRTL/include/State.h

@@ -33,7 +33,7 @@ namespace memory {
 /// Note: See the restrictions on __kmpc_alloc_shared for proper usage.
 void *allocShared(uint64_t Size, const char *Reason);
 
-/// Free \p Ptr, alloated via allocShared, for \p Reason.
+/// Free \p Ptr, allocated via allocShared, for \p Reason.
 ///
 /// Note: See the restrictions on __kmpc_free_shared for proper usage.
 void freeShared(void *Ptr, uint64_t Bytes, const char *Reason);
@@ -44,7 +44,7 @@ void *allocGlobal(uint64_t Size, const char *Reason);
 /// Return a pointer to the dynamic shared memory buffer.
 void *getDynamicBuffer();
 
-/// Free \p Ptr, alloated via allocGlobal, for \p Reason.
+/// Free \p Ptr, allocated via allocGlobal, for \p Reason.
 void freeGlobal(void *Ptr, const char *Reason);
 
 } // namespace memory
@@ -365,7 +365,7 @@ inline state::Value<uint32_t, state::VK_Level> Level;
 /// The `active-level` describes which of the parallel level counted with the
 /// `level-var` is active. There can only be one.
 ///
-/// active-level-var is 1, if ActiveLevelVar is not 0, otherweise it is 0.
+/// active-level-var is 1, if ActiveLevelVar is not 0, otherwise it is 0.
 inline state::Value<uint32_t, state::VK_ActiveLevel> ActiveLevel;
 
 /// TODO

offload/DeviceRTL/include/Synchronization.h

@@ -22,7 +22,7 @@ namespace atomic {
 
 enum OrderingTy {
   relaxed = __ATOMIC_RELAXED,
-  aquire = __ATOMIC_ACQUIRE,
+  acquire = __ATOMIC_ACQUIRE,
   release = __ATOMIC_RELEASE,
   acq_rel = __ATOMIC_ACQ_REL,
   seq_cst = __ATOMIC_SEQ_CST,

offload/DeviceRTL/src/Configuration.cpp

@@ -27,7 +27,7 @@ using namespace ompx;
     0;
 [[gnu::weak]] extern const uint32_t __omp_rtl_assume_teams_oversubscription = 0;
 
-// This variable should be visibile to the plugin so we override the default
+// This variable should be visible to the plugin so we override the default
 // hidden visibility.
 [[gnu::used, gnu::retain, gnu::weak,
   gnu::visibility("protected")]] DeviceEnvironmentTy

offload/DeviceRTL/src/Misc.cpp

@@ -33,7 +33,7 @@ double getWTime();
 
 double getWTick() {
   // The number of ticks per second for the AMDGPU clock varies by card and can
-  // only be retrived by querying the driver. We rely on the device environment
+  // only be retrieved by querying the driver. We rely on the device environment
   // to inform us what the proper frequency is.
   return 1.0 / config::getClockFrequency();
 }

offload/DeviceRTL/src/Reduction.cpp

@@ -206,7 +206,7 @@ int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
   // to the number of slots in the buffer.
   bool IsMaster = (ThreadId == 0);
   while (IsMaster) {
-    Bound = atomic::load(&IterCnt, atomic::aquire);
+    Bound = atomic::load(&IterCnt, atomic::acquire);
     if (TeamId < Bound + num_of_records)
       break;
   }
@@ -259,7 +259,7 @@ int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
   unsigned NumRecs = kmpcMin(NumTeams, uint32_t(num_of_records));
   if (ChunkTeamCount == NumTeams - Bound - 1) {
     // Ensure we see the global memory writes by other teams
-    fence::kernel(atomic::aquire);
+    fence::kernel(atomic::acquire);
 
     //
     // Last team processing.

offload/DeviceRTL/src/Synchronization.cpp

@@ -84,7 +84,7 @@ uint32_t atomicInc(uint32_t *A, uint32_t V, atomic::OrderingTy Ordering,
   default:
     __builtin_unreachable();
     Case(atomic::relaxed);
-    Case(atomic::aquire);
+    Case(atomic::acquire);
     Case(atomic::release);
     Case(atomic::acq_rel);
     Case(atomic::seq_cst);
@@ -107,7 +107,7 @@ void namedBarrier() {
   uint32_t WarpSize = mapping::getWarpSize();
   uint32_t NumWaves = NumThreads / WarpSize;
 
-  fence::team(atomic::aquire);
+  fence::team(atomic::acquire);
 
   // named barrier implementation for amdgcn.
   // Uses two 16 bit unsigned counters. One for the number of waves to have
@@ -172,7 +172,7 @@ void syncThreads(atomic::OrderingTy Ordering) {
   __builtin_amdgcn_s_barrier();
 
   if (Ordering != atomic::relaxed)
-    fenceTeam(Ordering == atomic::acq_rel ? atomic::aquire : atomic::seq_cst);
+    fenceTeam(Ordering == atomic::acq_rel ? atomic::acquire : atomic::seq_cst);
 }
 void syncThreadsAligned(atomic::OrderingTy Ordering) { syncThreads(Ordering); }
 
@@ -198,7 +198,7 @@ void setCriticalLock(omp_lock_t *Lock) {
         !cas((uint32_t *)Lock, UNSET, SET, atomic::relaxed, atomic::relaxed)) {
       __builtin_amdgcn_s_sleep(32);
     }
-    fenceKernel(atomic::aquire);
+    fenceKernel(atomic::acquire);
   }
 }
 

offload/DeviceRTL/src/Workshare.cpp

@@ -79,7 +79,7 @@ template <typename T, typename ST> struct omptarget_nvptx_LoopSupport {
     lb = lb + entityId * chunk;
     T inputUb = ub;
     ub = lb + chunk - 1; // Clang uses i <= ub
-    // Say ub' is the begining of the last chunk. Then who ever has a
+    // Say ub' is the beginning of the last chunk. Then who ever has a
     // lower bound plus a multiple of the increment equal to ub' is
     // the last one.
     T beginingLastChunk = inputUb - (inputUb % chunk);
@@ -806,7 +806,7 @@ template <typename Ty> class StaticLoopChunker {
                                 NumIters, OneIterationPerThread);
   }
 
-  /// Worksharing `distrbute`-loop.
+  /// Worksharing `distribute`-loop.
   static void Distribute(IdentTy *Loc, void (*LoopBody)(Ty, void *), void *Arg,
                          Ty NumIters, Ty BlockChunk) {
     ASSERT(icv::Level == 0, "Bad distribute");
@@ -853,7 +853,7 @@ template <typename Ty> class StaticLoopChunker {
     ASSERT(state::ParallelTeamSize == 1, "Bad distribute");
   }
 
-  /// Worksharing `distrbute parallel for`-loop.
+  /// Worksharing `distribute parallel for`-loop.
   static void DistributeFor(IdentTy *Loc, void (*LoopBody)(Ty, void *),
                             void *Arg, Ty NumIters, Ty NumThreads,
                             Ty BlockChunk, Ty ThreadChunk) {

offload/include/OpenMP/OMPT/Callback.h

@@ -56,7 +56,7 @@ FOREACH_OMPT_EMI_EVENT(declareOmptCallback)
 
 /// This function will call an OpenMP API function. Which in turn will lookup a
 /// given enum value of type \p ompt_callbacks_t and copy the address of the
-/// corresponding callback funtion into the provided pointer.
+/// corresponding callback function into the provided pointer.
 /// The pointer to the runtime function is passed during 'initializeLibrary'.
 /// \p which the enum value of the requested callback function
 /// \p callback the destination pointer where the address shall be copied

offload/include/PluginManager.h

@@ -77,7 +77,7 @@ struct PluginManager {
   /// Iterate over all device images registered with this plugin.
   auto deviceImages() { return llvm::make_pointee_range(DeviceImages); }
 
-  /// Translation table retreived from the binary
+  /// Translation table retrieved from the binary
   HostEntriesBeginToTransTableTy HostEntriesBeginToTransTable;
   std::mutex TrlTblMtx; ///< For Translation Table
   /// Host offload entries in order of image registration
@@ -171,7 +171,7 @@ struct PluginManager {
   /// Devices associated with plugins, accesses to the container are exclusive.
   ProtectedObj<DeviceContainerTy> Devices;
 
-  /// References to upgraded legacy offloading entires.
+  /// References to upgraded legacy offloading entries.
   std::list<llvm::SmallVector<llvm::offloading::EntryTy, 0>> LegacyEntries;
   std::list<llvm::SmallVector<__tgt_device_image, 0>> LegacyImages;
   llvm::DenseMap<__tgt_bin_desc *, __tgt_bin_desc> UpgradedDescriptors;

offload/include/device.h

@@ -134,7 +134,7 @@ struct DeviceTy {
   int32_t recordEvent(void *Event, AsyncInfoTy &AsyncInfo);
 
   /// Wait for an event. This function can be blocking or non-blocking,
-  /// depending on the implmentation. It is expected to set a dependence on the
+  /// depending on the implementation. It is expected to set a dependence on the
   /// event such that corresponding operations shall only start once the event
   /// is fulfilled.
   int32_t waitEvent(void *Event, AsyncInfoTy &AsyncInfo);

offload/include/omptarget.h

@@ -42,7 +42,7 @@ enum __tgt_target_return_t : int {
   OMP_TGT_SUCCESS = 0,
   /// offload may not execute on the requested target device
   /// this scenario can be caused by the device not available or unsupported
-  /// as described in the Execution Model in the specifcation
+  /// as described in the Execution Model in the specification
   /// this status may not be used for target device execution failure
   /// which should be handled internally in libomptarget
   OMP_TGT_FAIL = ~0

offload/liboffload/API/APIDefs.td

@@ -36,7 +36,7 @@ class IsPointerType<string Type> {
   bit ret = !ne(!find(Type, "*", !sub(!size(Type), 1)), -1);
 }
 
-// Describes the valid range of a pointer parameter that reperesents an array
+// Describes the valid range of a pointer parameter that represents an array
 class Range<string Begin, string End> {
   string begin = Begin;
   string end = End;

offload/liboffload/API/README.md

@@ -62,7 +62,7 @@ which preserves ABI compatibility with C.
 Represents a C-style enum. Contains a list of `etor` values, which have a name
 and description.
 
-A `TaggedEtor` record type also exists which addtionally takes a type. This type
+A `TaggedEtor` record type also exists which additionally takes a type. This type
 is used when the enum is used as a parameter to a function with a type-tagged
 function parameter (e.g. `olGetDeviceInfo`).
 

offload/liboffload/src/OffloadLib.cpp

@@ -36,7 +36,7 @@ OffloadConfig &offloadConfig() {
   return Config;
 }
 
-// Pull in the declarations for the implementation funtions. The actual entry
+// Pull in the declarations for the implementation functions. The actual entry
 // points in this file wrap these.
 #include "OffloadImplFuncDecls.inc"