11institutetext: University of New Mexico, Albuquerque NM, 87131
11email<EMAIL_ADDRESS>22institutetext: Tennessee Technological
University, Cookeville, TN 38505
22email<EMAIL_ADDRESS>
33institutetext: The University of Alabama, Tuscaloosa, AL 35401
33email<EMAIL_ADDRESS>

# Understanding GPU Triggering APIs
for MPI+X Communication

Patrick G. Bridges 11 0000-0003-4801-03903 Anthony Skjellum 22
0000-0001-5252-6600 Evan D. Suggs
22 0000-0002-8210-8992 Derek Schafer 11 0000-0001-8438-5144 Purushotham V.
Bangalore 33 0000-0002-1098-9997

###### Abstract

GPU-enhanced architectures are now dominant in HPC systems, but message-
passing communication involving GPUs with MPI has proven to be both complex
and expensive, motivating new approaches that lower such costs. We compare and
contrast stream/graph- and kernel-triggered MPI communication abstractions,
whose principal purpose is to enhance the performance of communication when
GPU kernels create or consume data for transfer through MPI operations.
Researchers and practitioners have proposed multiple potential APIs for stream
and/or kernel triggering that span various GPU architectures and approaches,
including MPI-4 partitioned point-to-point communication, stream
communicators, and explicit MPI stream/queue objects. Designs breaking
backward compatibility with MPI are duly noted. Some of these strengthen or
weaken the semantics of MPI operations. A key contribution of this paper is to
promote community convergence toward a stream- and/or kernel-triggering
abstraction by highlighting the common and differing goals and contributions
of existing abstractions. We describe the design space in which these
abstractions reside, their implicit or explicit use of stream and other non-
MPI abstractions, their relationship to partitioned and persistent operations,
and discuss their potential for added performance, how usable these
abstractions are, and where functional and/or semantic gaps exist. Finally, we
provide a taxonomy for stream- and kernel-triggered abstractions, including
disambiguation of similar semantic terms, and consider directions for future
standardization in MPI-5.

###### Keywords:

Stream-triggered message passing, Kernel-striggered message passing, MPI, GPU,
MPI+X, partitioned communication

##  1 Introduction

Most MPI implementations now include support for GPU data paths, enabling
standard MPI host-based APIs to move data to and from GPU memory. This feature
is generally termed _GPU-aware MPI_ and significant progress has been made
toward standardizing this implementation feature through work on _MPI memory
types_ [11]. Importantly, the MPI operation stages of _initiation_ ,
_starting_ , _completion_ , and _freeing_ [10] are all conducted by code
running on the host CPU in this approach; GPU code that needs to cause or
coordinate MPI operation state changes must do so indirectly by synchronizing
with CPU code that makes those updates on behalf of the GPU code.

Integrating MPI into GPU _control paths_ is also important because it could
eliminate the GPU/host communication latencies in current MPI implementations.
This is a more challenging problem since it requires addressing multiple
complex issues, including:

  * •

Which GPU control paths (stream execution and kernel execution) should support
interacting with MPI operations?

  * •

To what degree can current MPI abstractions be reused on GPU control paths?

  * •

What are the concurrency semantics of MPI operations on host and GPU control
paths?

The myriad proposed approaches for addressing these issues [1, 3, 5, 9, 10,
12, 13, 15, 16] have focused on _GPU-triggered communication APIs_ in which
host calls initialize and free MPI communication requests and the GPU control
path initiates and in some cases completes MPI communication requests. When
the GPU control path used for initiation and completion is the GPU stream
scheduler, here and elsewhere this is termed _stream-triggered communication_
and generally does not involve GPU-callable routines; instead the host API
interacts with the GPU stream control program to schedule implicit GPU
initiation and completion of communication operations. In contrast, using GPU
kernel thread execution for initiation and completion is generally termed
_kernel-triggered communication_ and requires the addition of explicit GPU-
callable APIs for initiating and completing communication.

Despite the many proposals, there is neither a consensus nor standard, either
formal or informal, on an appropriate MPI API for GPU-triggered communication.
To provide information and guidance for the MPI community to systematically
address this issue (such as in MPI-5, or ideally sooner), this paper collects,
organizes, and summarizes information about the various proposed MPI GPU-
triggered communication approaches and highlights their similarities and
differences. After providing background on representative GPU communication
triggering proposals (Section 2), this paper makes the following
contributions:

  * •

Presents a taxonomy for approaches to GPU-triggered communication that
highlights the major design decisions that an MPI API needs to address
(Section 3);

  * •

Summarizes the key features of and classifies eight different MPI stream- and
kernel-triggering API proposals using this taxonomy and compares and contrasts
key requirements and API design elements of these approaches (Section 4);

  * •

Highlights key gaps in the existing GPU triggering MPI APIs and in the MPI
standard that will need to be addressed by future proposals for a robust MPI
GPU triggering API (Section 5).

Finally, the paper offers conclusions and recommends a path forward in Section
6.

##  2 Representative GPU-Triggered Communication Abstractions

To illustrate the wide range of design choices and challenges in this area, we
begin by providing information on three proposals for GPU-triggered MPI
communication. Specifically, we discuss the stream triggered, two-sided
communication support in recent MPICH implementations [16], the stream-
triggered MPI one-sided communication support recently proposed by researchers
at Hewlett Packard Enterprise (HPE) [13], and extensions to the MPI 4.0
Partitioned Communication interface [10] to support GPU kernel triggering of
MPI communication [8]. We also present simple ping-pong communication for each
proposal to further elucidate their proposed approaches. Together, this
section highlights and provides essential background on the diverse design
choices we explored in designing the GPU triggered communication taxonomy
described below in Section 3. Note that while the examples in this section
utilize CUDA APIs, appropriate ROCM APIs could be utilized instead.

###  2.1 MPICH Triggering

The GPU triggering support in recent MPICH implementations [16] focuses
onstream triggering of MPI communication using two abstractions:

MPIX_Stream

–an MPI object to sequence MPI operations and to which external execution
contexts, for example GPU streams, can be bound. This provides clearly defined
concurrency semantics for MPI operations in multiple execution contexts
including multi-threaded host execution contexts and accelerator-based MPI+X
environments; and,

GPU Stream communicators

–MPI communicators to which an MPIX_-
Stream object is attached to a GPU stream execution context, and on which MPI
operations are enqueued to the attached GPU stream and return immediately in
the host execution context.

MPICH provides explicit enqueue variants of the supported GPU-triggered
operations that return immediately to the host; GPU synchronization routines
must be used to await for completion of these operations on the GPU. MPICH
implementers considered making `MPI_Send` and `MPI_Recv` on communicators with
GPU streams attached have enqueue semantics (similar to the previously-
proposed MPI-GDS API [15], but in the end rejected this because it would
change the semantics of well-known MPI APIs. Finally, MPICH stream
communicators also include additional features to support APIs similar to the
previous MPI Endpoints proposal [2].

⬇

1 MPIX_Stream mpistream;

2 MPI_Comm stream_comm;

3 MPI_Info info;

4

5 MPI_Info_create(&info);

6 MPI_Info_set(info, "type", "cudaStream_t");

7 MPIX_Info_set_hex(info, "value", &cudastream,

8 sizeof(cudastream));

9 MPIX_Stream_create(info, &mpistream);

10 MPI_Info_free(&info);

11 MPIX_Stream_comm_create(MPI_COMM_WORLD, mpistream,

12 &stream_comm);

13

14 for (int i = 0; i < niters; i++) {

15 if (my_rank == 0) {

16 MPIX_Send_enqueue(src_buf, 1, MPI_INT, 1, 123,

17 stream_comm);

18 MPIX_Recv_enqueue(src_buf, 1, MPI_INT, 1, 123,

19 stream_comm, MPI_STATUS_IGNORE);

20 } else if (my_rank == 1) {

21 MPIX_Recv_enqueue(dst_buf, 1, MPI_INT, 0, 123,

22 stream_comm, MPI_STATUS_IGNORE);

23 MPIX_Send_enqueue(dst_buf, 1, MPI_INT, 0, 123,

24 stream_comm);

25 }

26 }

27 cudaStreamSynchronize(cudastream);

Figure 1: Stream-triggered MPI Ping Pong using MPICH Stream Triggering [16]

Figure 1 shows an example ping-pong communication using this interface. We
highlight the following key features of this example:

  * •

To use GPU stream-triggered communications, the programmer must create a new
communicator;

  * •

The stream communicator concept, similar to the one previously proposed for
MVAPICH [15], can be naturally extended to include additional operations
beyond the point-to-point operations supported in the current implementation,
including collective communications; and,

  * •

MPICH’s stream communicators enqueued operations require the calling process
must explicitly synchronize with the underlying GPU stream using GPU-specific
operations to ensure completion of the corresponding MPI operation in the GPU
execution context.

###  2.2 HPE One-Sided

The HPE one-sided GPU triggering API [13] supports stream-triggered, one-sided
communication111HPE also provides stream-triggered, two-sided communication,
not considered here.. The stream abstractions included in HPE’s approach rely
on GPU-NIC asynchronous stream triggering that integrates the NIC command
queue and the GPU stream. In this approach, the CPU puts commands in the Cray
Cassini NIC’s command queue and then GPU stream operations triggers these
commands. This removes unnecessary synchronization on the CPU. HPE’s one-sided
API is designed to allow for GPU-NIC progress where the host CPU is fully
bypassed in progressing communication. HPE’s previous two-sided GPU triggering
APIs [12] required host progress to receive messages due to the need for tag
magtching and unexpected message handling.

⬇

1 /* Normal MPI communicator, window, group, and */

2 /* buffers assumed to exist */

3 for(int i = 0; i < niters; i++) {

4 if(rank == 0){

5 /* Send ping */

6 MPI_Win_start(group, MPI_MODE_STREAM, win);

7 MPI_Put(src,n,MPI_INT,1, disp, n, MPI_INT,win);

8 /* Puts triggered here */

9 MPIX_Win_complete_stream(win,stream);

10 /* Receive pong */

11 MPIX_Win_post_stream(group, win, stream);

12 MPIX_Win_wait_stream(win,stream);

13 } else { /* Receive ping */

14 MPIX_Win_post_stream(group, win, stream);

15 MPIX_Win_wait_stream(win, stream);

16 /* Send pong */

17 MPI_Win_start(group, MPI_MODE_STREAM, win);

18 MPI_Put(src, n, MPI_INT, 0, disp, n, MPI_INT,win);

19 /* Puts triggered here */

20 MPIX_Win_complete_stream(win, stream);

21 }

22 }

23 cudaStreamSynchronize(stream);

Figure 2: Stream-triggered MPI Ping Pong using HPE One-Sided Stream Triggering

Figure 2, shows an example of a one-sided ping-pong test using the HPE Cray
MPI interface where stream-triggered variants of window post, wait, and
complete operations are combined with window start and put operations to
achieve fully offloaded GPU communication. We highlight the following key
features of this example:

  * •

One-sided communication mixes new `MPIX_*` functions and standard MPI
functions, because the Window is associated with the stream.

  * •

`MPI_Put` and `MPI_Win_start` do not have stream-triggered `MPIX_`
alternatives and simply initialize future communications operations to be
triggered by the stream-triggered complete operation.

  * •

The `MPI_Win_complete_enqueue` initiates operations initialized by the host in
the command queue. This utilizes the GPU and NIC without weakening MPI RMA
synchronization semantics, and in reality strengthens MPI RMA semantics by
_always_ deferring one-sided operations until `MPI_Win_complete_` `enqueue` is
executed by the GPU stream control program.

###  2.3 Kernel-Triggered Partitioned Communication

⬇

1__global__ void ready(MPIX_Prequest preq) {

2 MPIX_Pready(threadIdx.x, preq);

3}

4

5__global__ void arrived(MPIX_Prequest preq) {

6 int flag = -1;

7 for(;;){

8 MPIX_Parrived(preq, threadIdx.x, &flag);

9 if(flag) break;

10 }

11}

12

13void pingpong(...) {

14 int sendbuf[1], recvbuf[1];

15 MPIX_Request srequest, rrequest;

16 MPIX_Prequest psrequest, prrequest;

17 int otherrank = (rank == 0) ? 1 : 0;

18 MPIX_Psend_init(sendbuf, 1, 1, MPI_INT, otherrank, 0,

19 MPI_COMM_WORLD, MPI_INFO_NULL,

20 &srequest);

21 MPIX_Precv_init(recvbuf, 1, 1, MPI_INT, otherrank, 0,

22 MPI_COMM_WORLD, MPI_INFO_NULL,

23 &rrequest);

24 MPIX_Prequest_create(srequest, &psrequest);

25 MPIX_Prequest_create(rrequest, &prrequest);

26 for(int i = 0; i < maxiters; i++) {

27 MPIX_Start(&rrequest); // ping if rank==1, else pong

28 MPIX_Start(&srequest); // pong if rank==1, else ping

29 if(rank == 0) {

30 ready<<<1, 1, 0, 0>>>(psrequest);

31 MPIX_Wait(&srequest, MPI_STATUS_IGNORE);

32 arrived<<<1, 1, 0, 0>>>(prrequest);

33 MPIX_Wait(&rrequest, MPI_STATUS_IGNORE);

34 } else {

35 arrived<<<1, 1, 0, 0>>>(prrequest);

36 MPIX_Wait(&rrequest, MPI_STATUS_IGNORE);

37 ready<<<1, 1, 0, 0>>>(psrequest);

38 MPIX_Wait(&srequest, MPI_STATUS_IGNORE);

39 }

40 }

41}

Figure 3: MPI Ping Pong using MPI-ACX Partitioned Stream Triggering

The MPI-ACX implementation [1] includes a prototype of the partitioned
communication detailed in MPI-4. Partitioned point-to-point communication was
specifically added to MPI-4 for a variety of reasons, including providing
matching semantics (matching at initialization) and lightweight interfaces for
data transfer initiation 222It also removes the restrictions of persistent
send and receive operations—MPI_Send_init and MPI_Recv_init—from MPI-1, which
are forbidden to communicate and do not match only once. These fail to offer a
channel semantic or obviate receive-queue matching. By contrast, single
partitioned mode provides a uni-directional channel concept.. A specific
nuance in this approach is that both `MPIX_` `Parrived` and `MPIX_Pready` can
be called from the host and the CUDA device. There are device-variants of the
host requests to support the device-invoked operations.

Notably, the absence of a strong progress guarantee in MPI limits the
guarantee that the partitioned communication initialization phase (i.e.,
MPI_Psend_init has already been matched proactively with MPI_Precv_init) when
`MPI_Start` is called; it can instead occur at the first `MPI_Wait` for both
sides of the partitioned channel. As such, proposals were offered for MPI-4.1
denoted MPI_Pbuf_prepare to guarantee initialization completion prior to first
use and ensure readiness of buffers upon reuse [6]. This has not yet been
adopted for the standard, but appears essential to avoid unexpected messages
that would introduce complex logic that is considered untenable in GPU
kernels.

Figure 3 details a ping-pong test based on the ring program supplied with MPI-
ACX; and the following are some key observations:

  * •

In `MPI-ACX`, the programmer uses `MPIX_Request` to create a `MPIX_Prequest`
for the device to use in operations. These objects are tied to each other with
`MPIX_Prequest``_create`.

  * •

The host-based `MPIX_Request` object is used for regular MPI functions such as
`MPIX_Start` and `MPIX_Wait`.

  * •

Different request types and funcation signatures are needed for GPU and CPU
interfaces due to differences between the state needed in each context.

##  3 A Taxonomy of MPI GPU-Triggered Communication

Based on a careful examination of both the MPI standard and prior proposals
for GPU-triggered communication, we created the following categories for
classifying key characteristics of MPI GPU-triggered communication proposals
in four different overarching areas:

_Area 1: GPU control path used:_

Does the API support the GPU stream control path or kernel control path for
initiating and completing MPI communication operations? In the case of APIs
where both stream and kernel triggering are supported, we split the API into
separate stream- and kernel-triggered sub-APIs and classify them separately
because of the frequent semantic and API differences between the such
approaches in other categories.

_Area 2: API Design Considerations_

Reuses Existing MPI APIs or abstractions:

Does the GPU triggering API reuse existing MPI communication operations or
abstractions to enable GPU triggering of communication operation initiation or
completion, or does it instead create new MPI operations or abstractions?

Changes Existing MPI API Semantics:

If the API reuses existing APIs or abstractions, does it change the semantics
of these API elements to behave differently than they currently do? For
proposals that modify the API, we classify whether the API _strengthens_ MPI
API semantics in a way that preserves backward compatibility or _weakens_ MPI
semantics in a way that may compromise backward compatibility.

Separate MPI Operation Initialization and Initiation:

Does the API use a single call to execute the _starting_ stages of MPI
communication operations or does it separate this into explicit API calls to
_initialize_ and _initiate_ MPI communication operations?

GPU MPI Operation Completion Support:

Does the API support GPU detection of the completion of GPU-initiated MPI
operations? We identify the classifications of _all_ , _some_ , and _none_ for
this category, corresponding to GPU control paths supporting waiting or
polling for the completion of (1) all GPU-initiated MPI operation, (2) some
GPU-initiated MPI operations, or (3) requiring host code to wait or poll for
the completion of MPI operations.

Collective communication support:

Does the API specify operations enabling a single GPU communication operation
to communicate with multiple processes? We classify collective support as
either _full_ , _group_ , or _none_ , cooresponding to (1) an API that
supports the full semantics of MPI two-sided collecgtives, (2) an API that
supports collective pairwise exchanges between subgroups of a communicator
such as one-sided group or two-sided neighbor collectives, or (3) no specified
support for collective communication.

_Area 3: Ordering and Concurrency Considerations_

MPI operation sequencing abstraction:

Does the API include an explicit MPI abstraction or object for ordering GPU
communication operations separate from GPU-provided abstractions?

Sequencing Abstraction semantics:

If the API includes an explicit sequencing abstraction, are communication
operations sequenced using that abstraction fully ordered or partially ordered
(e.g., using a fencing or graph abstraction)?

Specification of MPI Concurrency Standard Integration:

Does the API address concurrency-related issues arising from GPU triggering of
MPI communication through _explicit_ specification or _implicit_ specification
via the existing MPI _thread multiple_ mode, or are concurrency semantics
_unspecified_?

_Area 4: Implementation Considerations_

GPU/NIC Progress:

Does the API enable message initiation and completion by the GPU and network
device without the involvement of the host CPUS? This is generally
accomplished by supporting one-sided data movement that avoids the control
flow complexity of matching in the receive path.

Available Implementation or Evaluation:

Is an implementation or evaluatio of this API available either for download
and installation, from a vendor, or in a peer-reviewed publication so that it
can potentially be tested or compared with other APIs?

Multi-architecture Support:

Does the implementation, if available, target a single vendor architecture or
has it been demonstrated for multiple GPU and/or NIC accelerator
architectures? We categorize an implementation as _Yes_ when it supports
multiple GPUs and network architectures, as _GPU_ or _NIC_ when it has been
demonstrated on multiple GPU or NIC architectures, respectively, and as _No_
when it has been demonstrated on neither multiple NIC nor multiple GPU
architectures.

##  4 Classifying Existing MPI GPU Triggering Proposals

In this section, we classify these proposals and their availablity or
described implementations according to the taxonomy described in Section 3, as
shown in Table 1. We then use this taxonomical classification to compare and
contrast key features of these proposals and their corresponding
implementations.

###  4.1 API Classification

We focus on the eight proposed GPU-triggered MPI programming interface
extensions listed below:

MPI-GDS [15]

– a GPU-triggering interface extension for the MVAPICH MPI implementation
[14].

MCI-ACX Enqueued [1]

– stream-triggered two-sided communication operations as implemented in the
MPI-ACX GPU-triggering interface extension for NVIDIA network devices that
support MLX4 device functionality.

MPICH Triggering [16]

– a GPU triggering interface extension for the MPICH MPI implementation [7].

HPE Send/Recv [12]

– a stream-triggered MPI interface extension for HPE systems with the
Slingshot network interface [4].

Project Delorean [9]

– a proposal for graph-based sequencing of MPI operations that including GPU
triggering.

HPE One-sided [13]

– a GPU-triggered Post/Start/Complete/Wait interface extension for HPE systems
with the Slingshot network interface.

Partitioned Communication

– the MPI partitioned communication API [10] supplemented with the
`Pbuf_prepare` and partitioned collective operations recently proposed [8],
and partially implemented by MPI-ACX [1].

HPE Persistent [5]

– a kernel-triggered communication MPI interface extension for HPE systems
with the Slingshot network interface [4].

Note that multiple variants of some of these designs have been proposed in the
MPI Forum Working groups, BOFs, and other venues; we focus on the above listed
API descriptions and implementations as representative of the main existing
approaches at the time of this publication.

Table 1: Taxonomy of MPI GPU-Triggered Communication Proposals

| Area 1 | Area 2: API
---|---|---
|  | Features
Proposal | Control | Reuses | Changes | Separate | GPU | Collective
| Path | Existing | API | Initialize | Completion | Support
| Used | APIs | Semantics | and Initiate | Support |
MPI-GDS | Stream | Yes | Weaker | No | Full | Full1
MPI-ACX | Stream | Yes1 | No | Yes1 | Full | No
Enqueued |  |  |  |  |  |
MPICH | Stream | Yes | No2 | No | Full | No
Triggering |  |  |  |  |  |
HPE | Stream | No | No | Yes | Full | No
Send-Recv |  |  |  |  |  |
Delorean | Stream | No3 | No | Yes | Full | Full
HPE | Stream | Yes | Stronger4 | Yes | Full | Group
One-sided |  |  |  |  |  |
Partitioned | Kernel | Yes | No | Yes | Partial5 | Full1
Comm. |  |  |  |  |  |
HPE | Kernel | Yes | No | Yes | Full | No
Persistent |  |  |  |  |  |

| Area 3: Ordering and | Area 4: Implementation
---|---|---
| Concurrency | Considerations
Proposal | MPI | Sequencing | Concurrency | GPU/NIC | Impl. | Multiple
| Sequencing | Model | Model | Progress | Available | Arch.
| Object | Semantics | Specified |  |  | Support
MPI-GDS | No | N/A | Implicit | No | Yes | No
MPI-ACX | No | N/A | No | No | Yes | No
Enqueued |  |  |  |  |  |
MPICH | Yes | Full | Explicit | No | Yes | Yes
Triggering |  |  |  |  |  |
HPE | Yes | Partial | No | No | Yes | GPU
Send/Recv |  |  |  |  |  |
Delorean | Yes | Partial | Implicit | No | No | No
HPE | No | N/A | Implicit | Yes | Yes | GPU
One-sided |  |  |  |  |  |
Partitioned | No | N/A | Explicit | Yes | Yes | No
Comm. |  |  |  |  |  |
HPE | No | N/A | No | No | Yes | GPU
Persistent |  |  |  |  |  |

  * 1

Proposed in publication but not in the described or available implementations.

  * 2

The initial proposal considered changing the semantics of `MPI_Send` and
`MPI_Recv` on GPU stream communicators to be identical to their `_enqueue`
variants and return to the caller prior to local completion of the operation;
MPICH developers have since abandoned this approach.

  * 3

Proposes “deferred” operations as a generalization of MPI persistent
operations.

  * 4

Explicitly defers `MPI_Put`/`MPI_Get` execution to when
`MPI_Win_complete_stream` is called.

  * 5

There is no GPU call to determine if a partition that has been marked ready to
send has completed locally.

###  4.2 API Feature Comparisons

Based on the comparison shown in Table 1, there is no clear agreement on
whether stream or kernel triggering is preferable, either from an API,
programmability/usability, implementation, or performance perspective. Most
implementations seek to reuse portions of the API and maintain or strengthen
the semantics of the API operations they use, and also separate initialization
and initiation of API operations.

Almost every API allows the GPU to check for or wait for the completion of
every MPI operation that they initiate, allowing the GPU to avoid returning to
the host execution context to ensure communication completion. The main
exception is partitioned communication, where receives completion can be
tested on the GPU using `Parrived`, but there is no equivalent call for
testing for the completion of sends initiated from the kernel using `Pready`
(the only available option is to check if all send partitions have completed
using a single call to `Wait` or `Test` or one of their variants).

While several APIs propose collective support, only limited support for any
GPU-triggered collective communication is available in the form of point-to-
point group communication in the HPE one-sided triggering API. This is likely
due to the complexity of combining complex collective communication with the
constrained GPU and NIC control paths.

There is little agreement on whether to introduce a new MPI-level primitive to
represent GPU execution ordering, with some APIs introducing new primitives
and others reusing opaque GPU streams. Note that this is the case even from
the same vendor, where the HPE two-sided API includes an `MPI_Queue` object
that is not used by the HPE one-sided API.

Implementation availability also varies significantly. For instance, Delorean
is just a specification with no known implementation; MPI-ACX is fully
available and we were able to run the ping-ping example on representative
hardware. Few implementations support full GPU/NIC offload of communication,
with most APIs designed with the assumption that the host will be involved in
data movement.

##  5 Key Gaps

Based on the classifications described in the previous section, this section
highlights covers gaps and perceived limitations observed in prior work and
prototypes as well as the MPI standard.

###  5.1 Limitations in the MPI Standard

In the MPI Standard, we observe that there are two main areas that needs to be
addressed.

  1. 1.

Problems surrounding persistence:

The separation of initialization and initiation in GPU triggering APIs make
persistent operations attractive for this purpose, as demonstrated by
partitioned communication. Unfortunately, MPI-1 specified persistent send and
receive operations cannot communicate and their initialization does not result
in matching; indeed, persistent and non-persistent operations can match each
time they are initiated.

  2. 2.

Lack of MPI_Pbuf_prepare and similar operations for two-sided communication:

There are two separate issues that have been conflated in the prior
discussions of MPI_Pbuf_prepare and MPI_Pbuf_prepareall, yet to be
standardized [6]:

    1. (a)

Ensuring that MPI two-sided requests are matched (once and for all) so that
there is agreement between the sender and receiver side about the exposed
buffer on which to make operations (such as with MPI_Puts for early-bird
communication as MPI_Pready’s are issued on the send side).

    2. (b)

Ensuring or asserting that the remote buffer is _ready_ for one-sided data
movement; that is, there are no unexpected transfers.

The former evidently subsumes the latter; otherwise, matching/synchronization
would have to be done prior to every call to MPI_Start, which is likely to be
unacceptably high.

Further, because of the allowance for weak progress in MPI implementations,
there is no guarantee even in two sided operations that match at
initialization (e.g. partitioned communication) that the initial match has
occured between before the first complete communication. This means that
initial communication can happens before any agreement on the readiness of the
receive buffer; for instance, agreement and transmission could first happen
during the first wait calls, or the initial transfers could be done
differently than subsequent ones.

In addition, the desire for these conditions is not unique to partitioned
communication — the various stream-triggered communication proposals (MPICH,
HPE two-sided, MPI-ACX enqueuing, etc.) would want this functionality to push
their data movement through the one-sided path as well.

###  5.2 Limitations across Proposed APIs

In addition to gaps in the MPI standard that increase the difficulty of
supporting GPU triggering, there are also recurring limitations across the
various APIs that have been proposed. We highlight these gaps below to
encourage efforts to address them consistently in new proposals:

  * •

Clear concurrency model—If the host execution context and accelerator context
both initiate operations, the ordering and concurrency is ill-defined in many
APIS; any (left to the vagueness of MPI), or are silent. MPICH stream
triggering is the only one that provides a clear, specific concurrent model
for both host- and accelerator-initiated communication.

  * •

Matching—There is no explicit support for matching of send and receive buffers
so that this occurs only once, and receive queue matching overheads are
eliminated.

  * •

GPU completion support with partitioning—If something is initiated on a GPU,
there is no way to know if something is completed on the send side. This makes
it difficult or impossible to write programs that do not need to synchronizae
with the CPU to interleave communication and computation.

  * •

Collective communication support—This is broadly absent and needed to support
real applications that avoid unnecessary synchronization with the CPU. In
particular, _no_ current stream triggering approach supports full stream-
triggered collective communication.

  * •

GPU/NIC Progress—Lack of this feature is a gap for any performant
implementation. Out of the eight APIs considered only two (partitioned
communication augmented with `Pbuf_prepare` and HPE one-sided communication)
support GPU/NIC progress.

##  6 Conclusion

This is the first comprehensive survey of all eight proposals known to the
authors that address stream/graph- and kernel-triggered communication
abstractions for MPI, whose principal purpose is to improve the performance of
communication when GPU kernels create or consume data for transfer through MPI
operations. As stated at the outset, GPU-enhanced architectures are now
dominant in HPC systems, but message-passing communication involving GPUs with
MPI has proven to be both complex and expensive, motivating new approaches
that lower such costs.

Here, we compared and contrasted stream/graph- and kernel-triggered MPI
communication abstractions. Various proposals with distinct potential APIs for
stream- and/or kernel-triggering have been proffered as designs and/or
prototypes, which span various GPU architectures and approaches, including
MPI-4 partitioned point-to-point communication, stream communicators, and
explicit MPI stream/queue objects. Notably, some of these proposals either
strengthen or weaken the semantics of MPI operations. We enumerated eight
different proposals and provided a taxonomy of these proposals and their
operations against twelve categories of features and semantics.

We described the design space in which these abstractions reside, their
implicit or explicit use of stream and other non-MPI abstractions, their
relationship to partitioned and persistent operations, and discussed their
potential for added performance, how usable these abstractions are, and where
functional and/or semantic gaps exist. Designs breaking backward compatibility
with MPI were also duly noted (and cautioned). We also presented code examples
of how some of these operate on the ubiquitous ping-pong use case (some
utilize CUDA and others ROCM APIs).

The next steps to advance performance-portable programming with accelerator-
based MPI+X programming should start with robust, inclusive conversations and
design bake-offs among the various proposers, implementers, and users to
identify a coherent path forward within the MPI community, inclusive of
application developers. One or more proposed standard APIs should be competed
and down-selected for MPI-5 (or an earlier MPI-4.x release). Significantly,
the existing proposals/prototypes cited in this paper may not include all
hardware vendors or systems currently available and certainly there may be
other solutions forthcoming or that are not yet public (e.g., for FPGA-based
accelerators or other GPU accelerator products). Yet, fitting additional
and/or forthcoming proposals, prototypes, and designs into our taxonomy will
provide a coherent basis to compare and contrast these with the eight
presented here.

####  6.0.1 Acknowledgements

This work was performed with partial support from the National Science
Foundation under Grants Nos. OAC-2103510, CCF-2405142 and CCF-2412182, the
U.S. Department of Energy’s National Nuclear Security Administration (NNSA)
under the Predictive Science Academic Alliance Program (PSAAP-III), Award DE-
NA0003966, and Tennessee Technological University.

Any opinions, findings, and conclusions or recommendations expressed in this
material are those of the authors and do not necessarily reflect the views of
the National Science Foundation, or the U.S. Department of Energy’s National
Nuclear Security Administration.

####  6.0.2

The authors have no competing interests to declare that are relevant to the
content of this article.

## References

  * [1] Dinan, J.: MPI-ACX: MPI Accelerator Extensions Prototype (2023), https://github.com/NVIDIA/mpi-acx
  * [2] Dinan, J., Grant, R.E., Balaji, P., Goodell, D., Miller, D., Snir, M., Thakur, R.: Enabling communication concurrency through flexible MPI endpoints. The International Journal of High Performance Computing Applications 28(4), 390–405 (2014). https://doi.org/10.1177/1094342014548772, https://doi.org/10.1177/1094342014548772
  * [3] Elis, B., Pearce, O., Böhme, D., Burmark, J., Schulz, M.: Non-blocking GPU-CPU notifications to enable more GPU-CPU parallelism. In: Proceedings of the International Conference on High Performance Computing in Asia-Pacific Region, HPCAsia 2024, Nagoya, Japan, January 25-27, 2024. pp. 1–11. ACM (2024). https://doi.org/10.1145/3635035.3635036, https://doi.org/10.1145/3635035.3635036
  * [4] Enterprise, H.P.: HPE Slingshot Interconnect, https://www.hpe.com/in/en/compute/hpc/slingshot-interconnect.html
  * [5] Enterprise, H.P.: MPIX_start - Cray MPICH Documentation (2024), https://cpe.ext.hpe.com/docs/mpt/mpich/mpix_start.html
  * [6] Forum, M.: Synchronization on Partitioned Communication for Accelerator Optimization (2020), https://github.com/mpi-forum/mpi-issues/issues/302
  * [7] Gropp, W.: MPICH2: A New Start for MPI Implementations. In: Kranzlmüller, D., Volkert, J., Kacsuk, P., Dongarra, J. (eds.) Recent Advances in Parallel Virtual Machine and Message Passing Interface. pp. 7–7. Springer Berlin Heidelberg, Berlin, Heidelberg (2002)
  * [8] Holmes, D.J., Skjellum, A., Jaeger, J., Grant, R.E., Bangalore, P.V., Dosanjh, M.G., Bienz, A., Schafer, D.: Partitioned Collective Communication. In: 2021 Workshop on Exascale MPI (ExaMPI). pp. 9–17. IEEE, St. Louis, MO, USA (Nov 2021). https://doi.org/10.1109/ExaMPI54564.2021.00007, https://ieeexplore.ieee.org/document/9652828/
  * [9] Koziol, Q.: Composable Asynchronous Communication Graphs and Streams in MPI (2023), https://github.com/mpiwg-hybrid/hybrid-issues/files/11680187/Project_Delorean-2023-06-07a.pdf
  * [10] Message Passing Interface Forum: MPI: A Message-Passing Interface Standard Version 4.0 (Jun 2021), https://www.mpi-forum.org/docs/mpi-4.0/mpi40-report.pdf
  * [11] Message Passing Interface Forum: MPI: A Message-Passing Interface Standard Version 4.1 (Nov 2023), https://www.mpi-forum.org/docs/mpi-4.1/mpi41-report.pdf
  * [12] Namashivayam, N., Kandalla, K., White, T., Radcliffe, N., Kaplan, L., Pagel, M.: Exploring GPU Stream-Aware Message Passing using Triggered Operations (Aug 2022), http://arxiv.org/abs/2208.04817, arXiv:2208.04817 [cs]
  * [13] Namashivayam, N., Kandalla, K., White III, J.B., Kaplan, L., Pagel, M.: Exploring Fully Offloaded GPU Stream-Aware Message Passing (Jun 2023), http://arxiv.org/abs/2306.15773, arXiv:2306.15773 [cs]
  * [14] Panda, D.K., Subramoni, H., Chu, C.H., Bayatpour, M.: The MVAPICH project: Transforming research into high-performance MPI library for HPC community. Journal of Computational Science 52, 101208 (2021). https://doi.org/https://doi.org/10.1016/j.jocs.2020.101208, https://www.sciencedirect.com/science/article/pii/S1877750320305093, case Studies in Translational Computer Science
  * [15] Venkatesh, A., Hamidouche, K., Potluri, S., Rosetti, D., Chu, C.H., Panda, D.K.: MPI-GDS: High Performance MPI Designs with GPUDirect-aSync for CPU-GPU Control Flow Decoupling. In: 2017 46th International Conference on Parallel Processing (ICPP). pp. 151–160. IEEE, Bristol, United Kingdom (Aug 2017). https://doi.org/10.1109/ICPP.2017.24, http://ieeexplore.ieee.org/document/8025289/
  * [16] Zhou, H., Raffenetti, K., Guo, Y., Thakur, R.: MPIX Stream: An Explicit Solution to Hybrid MPI+X Programming. In: Proceedings of the 29th European MPI Users’ Group Meeting. pp. 1–10. ACM, Chattanooga TN USA (Sep 2022). https://doi.org/10.1145/3555819.3555820, https://dl.acm.org/doi/10.1145/3555819.3555820


# Small oscillations of a heavy symmetric top

V. Tanrıverdi
<EMAIL_ADDRESS>
Address: Bahadın Kasabası 66710 Sorgun-Yozgat TURKEY

###### Abstract

This study is on small oscillations of a heavy symmetric top. A different
method than previous works is applied, and differently from previous works,
the explicit formulas for the amplitudes for oscillations are given. This
method can be helpful in teaching.

##  1 Introduction

Studying the motion of a heavy symmetric top by considering small oscillations
is not a new topic, it dates back to $19^{th}$ century. Different scientists
had studied this topic from different perspectives. Routh studied small
oscillations of a heavy symmetric top by considering regular precession and
oscillations of both inclination angle and precession angular velocity [2].
Crabtree also used a similar method and obtained the same result [3]. Lamb
also considered similar arguments to Routh and obtained an equivalent result
[4]. Scarborough considered the motion with cusps and studied small
oscillations [5]. There are also more recent works considering small
oscillations of a heavy symmetric top. Goldstein et al., like Scarborough,
consider tops making the motion with cusps and look for the fast top limit
[6]. Arnold also studies small oscillations of a heavy symmetric top
considering motion with cusps but from a different aspect [7]. Greiner
considers the effective potential and obtains angular frequency for small
oscillations [8].

In general, a solution for the motion of a heavy symmetric top can be found
either by using elliptic integrals or numerical integrations of angular
accelerations [6, 9]. These can be found abstract by students, especially when
they first meet the concept of rigid body rotations. And, as a result of this,
they may find bizarre changes in angular velocities. On the other hand,
considering the motion in terms of small oscillations provides an approximate
analytical solution, which can be used to show different aspects of the
motion. This can provide a nice path for teaching.

In this work, differently from previous works, we will use angular
acceleration of the inclination angle to obtain equations for small
oscillations. This derivation is a bit easier to follow and it can increase
the understandings of students. Besides this, we will calculate amplitudes of
oscillations, and we will also find changes in precession and spin angles and
angular velocities in terms of small oscillations. Scarborough studied motion
with cups in a similar way and used this to plot the shapes for locus [5], and
we will also plot them in a similar way. We will consider two cases: The first
one is easier and can be used while teaching to undergraduate students, and
the second one can be used for upper-level undergraduate and graduate
students.

In section 2, we will give basic equations for a heavy symmetric top which
will be used for small oscillations. In section 3, we will obtain equations
for small oscillations. In section 4, we will conclude. In the appendix, we
will give other relations for angular frequency.

##  2 Heavy symmetric top

For a heavy symmetric top, Lagrangian can be written as [10]

$L=\frac{I_{x}}{2}(\dot{\theta}^{2}+\dot{\phi}^{2}\sin^{2}\theta)+\frac{I_{z}}{2}(\dot{\psi}+\dot{\phi}\cos\theta)^{2}-Mgl\cos\theta,$
(1)

where $\theta$, $\phi$ and $\psi$ are Euler angles, and the time derivatives
of them correspond angular velocities, $M$ is the mass, $l$ is the distance
between the center of mass and the fixed point of the top, $I_{x}$ and $I_{z}$
are moments of inertia, and $g$ is the gravitational acceleration.

By using Euler-Lagrange equations, the angular acceleration $\ddot{\theta}$
can be obtained as [9]

$\ddot{\theta}=\frac{\sin\theta}{I_{x}}\left[Mgl+I_{x}\dot{\phi}^{2}\cos\theta-
I_{z}\dot{\phi}^{2}\cos\theta-I_{z}\dot{\phi}\dot{\psi}\right],$ (2)

and two conserved angular momenta can be obtained as

$\displaystyle L_{z}$ $\displaystyle=$ $\displaystyle
I_{z}(\dot{\psi}+\dot{\phi}\cos\theta),$ $\displaystyle L_{z^{\prime}}$
$\displaystyle=$ $\displaystyle
I_{x}\dot{\phi}\sin^{2}\theta+I_{z}(\dot{\psi}+\dot{\phi}\cos\theta)\cos\theta.$
(3)

From these, one can define two constants: $a=L_{z}/I_{x}$ and
$b=L_{z^{\prime}}/I_{x}$ [6]. By using these constants, one can obtain
$\dot{\phi}$ and $\dot{\psi}$ as

$\displaystyle\dot{\phi}(\theta)$ $\displaystyle=$
$\displaystyle\frac{b-a\cos\theta}{\sin^{2}\theta},$ (4)
$\displaystyle\dot{\psi}(\theta)$ $\displaystyle=$
$\displaystyle\frac{I_{x}}{I_{z}}a-\frac{b-a\cos\theta}{\sin^{2}\theta}\cos\theta.$
(5)

These two equations show that finding the change of $\theta$ means finding
changes of $\dot{\phi}$ and $\dot{\psi}$.

One can define a constant from energy as

$E^{\prime}=\frac{I_{x}}{2}\dot{\theta}^{2}+\frac{I_{x}}{2}\frac{(b-a\cos\theta)^{2}}{\sin^{2}\theta}+Mgl\cos\theta,$
(6)

and by using this constant, one can define an effective potential [6]

$U_{eff}(\theta)=\frac{I_{x}}{2}\frac{(b-a\cos\theta)^{2}}{\sin^{2}\theta}+Mgl\cos\theta.$
(7)

This effective potential goes to infinity at the domain boundaries of $\theta$
($0$ and $\pi$), and it has a minimum between these when $|b|\neq|a|$.

##  3 Small oscillations

For a physical motion, $E^{\prime}$ is either equal to the minimum of the
effective potential or greater than that. When $E^{\prime}$ is slightly
greater than $U_{eff_{min}}$, changes in $\theta$ become small and one can use
small oscillation approximation.

We will take $I_{x}=I_{y}=22.8\times 10^{-5}\,kg\,m^{2}$, $I_{z}=5.72\times
10^{-5}\,kg\,m^{2}$ and $Mgl=0.068J$ for the below examples. We will compare
the results with the ones obtained from numerical integrations of angular
accelerations, the details of which can be found in previous work [9].

###  3.1 Small oscillations for motion with cusps

If one chooses initial conditions of a heavy symmetric top as
$\dot{\theta}_{0}=0$ and $\dot{\phi}_{0}=0$ ($\theta_{0}\neq 0$ and $\pi$, and
$\dot{\psi}_{0}\neq 0$), then $b$ becomes equal to $a\cos\theta_{0}$ and
$E^{\prime}=Mglb/a$. In these conditions, motion with cusps is observed.

One can rewrite equation (2) in terms of $a$ and $b$ as

$\ddot{\theta}=\frac{Mgl}{I_{x}}\sin\theta+\frac{(b-a\cos\theta)^{2}\cos\theta}{\sin^{3}\theta}-\frac{a(b-a\cos\theta)}{\sin\theta}.$
(8)

By considering small oscillations and using $\theta=\theta_{0}+\delta$, one
can approximate $\sin\theta\approx\sin\theta_{0}+\delta\cos\theta_{0}$ and
$\cos\theta\approx\cos\theta_{0}-\delta\sin\theta_{0}$. For small
oscillations, one can expand the denominators of equation 8 and ignore higher-
order terms in $\delta$. One can write $b=a\cos\theta_{0}$ for motions with
cusps, and then one can obtain

$\ddot{\delta}+\left(a^{2}-\frac{Mgl}{I_{x}}\cos\theta_{0}\right)\delta\approx\frac{Mgl}{I_{x}}\sin\theta_{0}.$
(9)

Here, one can define two constants of motion, which can be obtained from
initial conditions as: $w_{1}^{2}=a^{2}-(Mgl/I_{x})\cos\theta_{0}$ and
$C_{1}=(Mgl/I_{x})\sin\theta_{0}$. The first one corresponds to angular
frequency. The left-hand side of this equation is the same as the harmonic
oscillator and by using this property, one can write a solution as
$\delta(t)=A\cos(w_{1}t)+B\sin(w_{1}t)+C_{1}/w_{1}^{2}$. If initially
$\dot{\delta}(t=0)=0$ and $\delta(t=0)=0$ or $\dot{\theta}(t=0)=0$ and
$\theta(t=0)=\theta_{0}$, then one can get the amplitude of the oscilation as
$A=-C_{1}/w_{1}^{2}$ and $B=0$, and the solution becomes

$\delta(t)\approx\frac{C_{1}}{w_{1}^{2}}(1-\cos(w_{1}t)).$ (10)

Then, the solution for $\theta$ and angular velocities can be obtained as

$\displaystyle\theta(t)$ $\displaystyle\approx$
$\displaystyle\theta_{0}+\frac{C_{1}}{w_{1}^{2}}(1-\cos(w_{1}t)),$ (11)
$\displaystyle\dot{\theta}(t)$ $\displaystyle\approx$
$\displaystyle\frac{C_{1}}{w_{1}}\sin(w_{1}t),$ (12)
$\displaystyle\dot{\phi}(t)$ $\displaystyle\approx$
$\displaystyle\frac{a}{\sin\theta_{0}}\frac{C_{1}}{w_{1}^{2}}(1-\cos(w_{1}t)),$
(13) $\displaystyle\dot{\psi}(t)$ $\displaystyle\approx$
$\displaystyle\dot{\psi}_{0}-\frac{a\cos\theta_{0}}{\sin\theta_{0}}\frac{C_{1}}{w_{1}^{2}}(1-\cos(w_{1}t)).$
(14)

While obtaining solutions for $\dot{\phi}(t)$ and $\dot{\psi}(t)$, we have
used equations (4) and (5) and applied small oscillations procedure. One can
also integrate $\dot{\phi}(t)$ and $\dot{\psi}(t)$ with respect to time and by
considering $\phi(t=0)=0$ and $\psi(t=0)=0$, and find

$\displaystyle\phi(t)$ $\displaystyle\approx$
$\displaystyle\frac{a}{\sin\theta_{0}}\frac{C_{1}}{w_{1}^{2}}\left(t-\frac{\sin(w_{1}t)}{w_{1}}\right),$
(15) $\displaystyle\psi(t)$ $\displaystyle\approx$
$\displaystyle\dot{\psi}_{0}t-\frac{a\cos\theta_{0}}{\sin\theta_{0}}\frac{C_{1}}{w_{1}^{2}}\left(t-\frac{\sin(w_{1}t)}{w_{1}}\right).$
(16)

When $|a|>>|Mgl|/I_{x}$, which is valid for very fast tops, one can
approximate angular frequency as $a$ which we will designate as $w_{2}$. This
is an equivalent result to the ones given in classical mechanics books [4, 5,
6, 7].

Now, we will consider an example and compare results for $w_{1}$ and $w_{2}$
with the solution obtained from numerical integrations of angular
accelerations. In this example, the initial conditions are chosen as
$\dot{\phi}_{0}=0$, $\dot{\theta}_{0}=0$, $\theta_{0}=0.524\,rad$ and
$\dot{\psi}_{0}=700rad\,s^{-1}$. Results of numerical integrations and small
oscillations for $w_{1}$ and $w_{2}$ can be seen in figure 1. The shapes for
the locus for this case can be seen in figure 2(a), which are obtained both
from small oscillation approximation with $w_{1}$ and numerical integration of
angular accelerations.

(a) $\theta$ (b) $\dot{\theta}$ (c) $\dot{\phi}$ (d) $\dot{\psi}$

Figure 1:  Time evolution of $\theta$ (a), $\dot{\theta}$ (b), $\dot{\phi}$
(c) and $\dot{\psi}$ (d) for motion with cusps. Continuous (black) curves show
the results of numerical integration of angular accelerations, dotted-dashed
(red) curves represent results for the small oscillations with $w_{2}$, dashed
(blue) curves show results for the small oscillations with $w_{1}$. Initial
values are $\theta_{0}=0.524\,rad$, $\dot{\theta}_{0}=0$, $\dot{\phi}_{0}=0$
and $\dot{\psi}_{0}=700\,rad\,s^{-1}$, and $T=0.0364\,s$.

From figure 1, it can be seen that the small oscillation approximation gives
close results to the numerical solution obtained from the integration of
angular accelerations. One can say that considering small oscillations gives
slightly different angular frequency with respect to the integration of
angular accelerations, i.e. $w_{1}=175rad\,s^{-1}>w=173rad\,s^{-1}$ , and
slightly different amplitudes which are expected results due to ignoring
higher-order terms. The percentage difference in angular frequency is about
$1.3\%$, and the percentage difference in the amplitude of oscillations of
$\theta$ is $1.6\%$.

From figure 1, it can be seen that ignoring $(Mgl/I_{x})\cos\theta_{0}$ causes
small differences. As a result of this ignoring, the angular frequency becomes
equal to $a$, i.e. $w_{2}=176rad\,s^{-1}$, and the percentage difference is
slightly greater, i.e. $1.7\%$. The percentage difference for the amplitude of
oscillations of $\theta$ is the same since we have used the same amplitude.

As $|\dot{\psi}_{0}|$ and dependently $|a|$ increase, the percentage
differences become smaller. And, there are two reasons for this: One of them
is that as $|a|$ increases the contribution from $Mgl\cos\theta$ becomes
smaller which results in the increase of the resemblance of effective
potential to the potential of the harmonic oscillator as long as $|b|$ is not
close to $|a|$. The other one is related to difference between
$E^{\prime}=Mglb/a$ and $U_{eff_{min}}$, which becomes smaller as $|a|$
increases [11].

Figure 2:  Shapes for the locus. Continuous (black) curves show results of
numerical integration of angular accelerations, and dashed (blue) curves show
results of small oscillation approximation, i.e. obtained by using equations
(11) and (15) or (20) and (24). (a) Motion with cusps, initial values are
given in figure 1. (b) Wavy precession, initial values are given in figure 3.
(c) Looping motion, initial values are given in figure 4. (c) Motion when
$b>a$, initial values are given in figure 5.

###  3.2 Small oscillations in different types of motion

Now, we will study the small oscillations of a heavy symmetric top in general,
i.e. $b=a\cos\theta_{0}$ condition does not need to be satisfied. By
considering small oscillations similar to the previous case, one can write
$\theta=\theta_{0}+\eta$. Then, from equation (8), one can obtain

$\displaystyle\ddot{\eta}$ $\displaystyle\approx$
$\displaystyle\frac{Mgl}{I_{x}}\sin\theta_{0}+\frac{(a^{2}+b^{2})\cos\theta_{0}-ab(1+\cos^{2}\theta_{0})}{\sin^{3}\theta_{0}}$
$\displaystyle+\eta\left[\frac{Mgl}{I_{x}}\cos\theta_{0}+\frac{ab\cos\theta_{0}(5+\cos^{2}\theta_{0})-(a^{2}+b^{2})(1+2\cos^{2}\theta_{0})}{\sin^{4}\theta_{0}}\right].$

Then, one can define

$\displaystyle w_{g}^{2}$ $\displaystyle=$
$\displaystyle-\frac{Mgl}{I_{x}}\cos\theta_{0}-\frac{ab\cos\theta_{0}(5+\cos^{2}\theta_{0})-(a^{2}+b^{2})(1+2\cos^{2}\theta_{0})}{\sin^{4}\theta_{0}},$
(18) $\displaystyle C_{g}$ $\displaystyle=$
$\displaystyle\frac{Mgl}{I_{x}}\sin\theta_{0}+\frac{(a^{2}+b^{2})\cos\theta_{0}-ab(1+\cos^{2}\theta_{0})}{\sin^{3}\theta_{0}}.$
(19)

Then, one can write $\ddot{\eta}+w_{g}^{2}\eta\approx C_{g}$ whose solution
can be written as $\eta(t)=A\cos(w_{g}t)+B\sin(w_{g}t)+C_{g}/w_{g}^{2}$.
Similar to previous case, one can consider the case $\theta(t=0)=\theta_{0}$
and $\dot{\theta}(t=0)=0$ and obtain for small oscillations

$\displaystyle\theta(t)$ $\displaystyle\approx$
$\displaystyle\theta_{0}+\frac{C_{g}}{w_{g}^{2}}(1-\cos(w_{g}t)),$ (20)
$\displaystyle\dot{\theta}(t)$ $\displaystyle\approx$
$\displaystyle\frac{C_{g}}{w_{g}}\sin(w_{g}t),$ (21)
$\displaystyle\dot{\phi}(t)$ $\displaystyle\approx$
$\displaystyle\dot{\phi}_{0}+\frac{a-2\dot{\phi}_{0}\cos\theta_{0}}{\sin\theta_{0}}\frac{C_{g}}{w_{g}^{2}}(1-\cos(w_{g}t)),$
(22) $\displaystyle\dot{\psi}(t)$ $\displaystyle\approx$
$\displaystyle\dot{\psi}_{0}+\frac{2\dot{\phi}_{0}-b}{\sin\theta_{0}}\frac{C_{g}}{w_{g}^{2}}(1-\cos(w_{g}t)).$
(23)

Both angular frequency $w_{g}$ and $C_{g}$ can be obtained from initial values
so do amplitudes for all oscillations. Similar to previous case, one can
integrate $\dot{\phi}(t)$ and $\dot{\psi}(t)$ with respect to time and find

$\displaystyle\phi(t)$ $\displaystyle\approx$
$\displaystyle\dot{\phi}_{0}t+\frac{a-2\dot{\phi}_{0}\cos\theta_{0}}{\sin\theta_{0}}\frac{C_{g}}{w_{g}^{2}}\left(t-\frac{\sin(w_{g}t)}{w_{g}}\right),$
(24) $\displaystyle\psi(t)$ $\displaystyle\approx$
$\displaystyle\dot{\psi}_{0}t+\frac{2\dot{\phi}_{0}-b}{\sin\theta_{0}}\frac{C_{g}}{w_{g}^{2}}\left(t-\frac{\sin(w_{1}t)}{w_{1}}\right),$
(25)

where $\phi(t=0)=0$ and $\psi(t=0)=0$, again.

(a) $\theta$ (b) $\dot{\theta}$ (c) $\dot{\phi}$ (d) $\dot{\psi}$

Figure 3:  Time evolution of $\theta$ (a), $\dot{\theta}$ (b), $\dot{\phi}$
(c) and $\dot{\psi}$ (d) for wavy precession. Continuous (black) curves show
the results of numerical integration of angular accelerations, dashed (blue)
curves show results of equations 20, 21, 22 and 23. Initial values are
$\theta_{0}=0.524\,rad$, $\dot{\theta}_{0}=0$,
$\dot{\phi}_{0}=10\,rad\,s^{-1}$ and $\dot{\psi}_{0}=150\,rad\,s^{-1}$, and
$T=0.264\,s$.

Now, we will consider three different examples: Wavy precession, looping
motion and a motion when $b>a$. Shapes for the locus, obtained from both
numerical integration of angular accelerations and small oscillations for
these examples, can be seen in figure 2.

The first example is wavy precession which is the most possible motion for
small oscillations when $|a|>|b|$, since $U_{eff_{min}}<E^{\prime}<Mglb/a$ in
this case [12]. Initial values of this example can be found in the
explanations of figure 3. From that figure, it can be seen that results of
equations (20), (21), (22) and (23) are close to the results of numerical
integrations of angular accelerations. For this example,
$w_{g}=23.3rad\,s^{-1}$ and the angular frequency obtained from numerical
integration is $w=23.8rad\,s^{-1}$, and the percentage difference is $2.1\%$,
which is $0.2\%$ for the amplitude of $\theta$.

(a) $\theta$ (b) $\dot{\theta}$ (c) $\dot{\phi}$ (d) $\dot{\psi}$

Figure 4:  Time evolution of $\theta$ (a), $\dot{\theta}$ (b), $\dot{\phi}$
(c) and $\dot{\psi}$ (d) for looping motion. Legends are the same as figure 3.
Initial values are $\theta_{0}=0.524\,rad$, $\dot{\theta}_{0}=0$,
$\dot{\phi}_{0}=4.5\,rad\,s^{-1}$ and $\dot{\psi}_{0}=700\,rad\,s^{-1}$, and
$T=0.0362\,s$.

The second example of this part is looping motion which is seen when
$E^{\prime}>Mglb/a$ ($|a|>|b|$) [12]. Then, to obtain an adequate example for
looping motion with small oscillations, one may need to choose initial
conditions resulting in greater $|a|$ values. This is the reason for choosing
a greater $\dot{\psi}_{0}$ value, which can be found in the explanations of
figure 4 together with other initial values. From that figure, it can be seen
that one can obtain an approximate solution for looping motion by using small
oscillation approximation. For this example, $w_{g}=170rad\,s^{-1}$ and
$w=174rad\,s^{-1}$, and the percentage difference is $2\%$, which is $1.9\%$
for the amplitude of $\theta$.

(a) $\theta$ (b) $\dot{\theta}$ (c) $\dot{\phi}$ (d) $\dot{\psi}$

Figure 5:  Time evolution of $\theta$ (a), $\dot{\theta}$ (b), $\dot{\phi}$
(c) and $\dot{\psi}$ (d) for the motion when $b>a$. Legends are the same as
figure 3. Initial values are $\theta_{0}=0.524\,rad$, $\dot{\theta}_{0}=0$,
$\dot{\phi}_{0}=50\,rad\,s^{-1}$ and $\dot{\psi}_{0}=150\,rad\,s^{-1}$, and
$T=0.142\,s$.

The third example of this part is a motion when $b>a$. The shapes for the
locus for these types of motions are always similar [12]. Results of numerical
integrations of angular accelerations and small oscillation approximation can
be seen in figure 5 where one can also find initial values. One can see from
that figure that the results of small oscillation approximation are close to
numerical integrations’s. For this example, $w_{g}=45.2rad\,s^{-1}$ and
$w=42.3rad\,s^{-1}$, and the percentage difference is $7\%$, which is $3.1\%$
for the amplitude of $\theta$.

##  4 Conclusion

In this work, we studied small oscillations of a heavy symmetric top. We
followed a different way from the previous works and obtained an equivalent
result for angular frequency to that of the previous works, which can be found
in the appendix. In some cases, other forms of angular frequency can be more
convenient and can be used. In this work, we used the form given above to keep
the derivation simple. We should note that solutions are obtained when
$\dot{\theta}_{0}=0$, and for other cases, the solutions should be modified.

We have obtained equations describing the changes in angles and angular
velocities for small oscillations. We have seen that angular velocities change
as $\theta$ changes. These changes in angular velocities occur due to the
conservation of angular momenta and energy as they should be. These
conservation laws can not be explicitly seen from these equations. But, we
know that these equations are obtained from the equations of motion by using
Lagrangian, which is time-independent and energy is conserved, and, in fact,
two of the equations of motion are noting but conservation of angular momenta
for this case. On the other hand, due to considering small oscillations and
ignoring higher-order terms, these equations approximately satisfy
conservation laws and they are valid up to some point, for total concordance,
naturally, one needs to consider full equations.

We have considered different types of motion and chosen moderate examples, for
which percentage differences are less than $10\%$. However, one can obtain
examples with less than $1\%$ percentage differences when the top is fastly
spinning, i.e. $|a|$ is much greater than used values. On the other hand, as
$|a|$ decreases or the difference between $E^{\prime}$ and $U_{eff_{min}}$
increases the difference becomes bigger, and for larger values, small
oscillation approximation can not be used as expected. Shapes for the locus
can also be approximately obtained from small oscillations by the formulas
received for $\theta$ and $\phi$, i.e. equations (11) and (15) or (20) and
(24), which can be helpful in teaching. The studied examples show that both
the angular frequency and amplitudes are convenient for small oscillations of
a heavy symmetric top, and studying different types of motion as examples
shows that they can be used in different types of motion. We should note that
this small oscillation approximation also works for the negative $Mgl$ case.

##  5 Appendix

Routh approaches the small oscillations from a different perspective and takes
into account small oscillations both in $\theta$ and $\dot{\phi}$ by
considering small oscillations around the regular precession [2] occurring at
the minimum of the effective potential. By assuming the form of the solution
as $\eta_{R}=F\sin(w_{R}t+f)$, Routh finds the square of the angular frequency
for small oscillations of $\theta$ which can be written in terms of parameters
of this work as

$w_{R}^{2}=\dot{\phi}_{rp}^{2}-\frac{2mgl}{I_{x}}\cos\theta_{rp}+\left(\frac{mgl}{I_{x}\dot{\phi}_{rp}}\right)^{2},$
(26)

where $\theta_{rp}$ is the angle giving regular precession which is the angle
making the effective potential minimum, and $\dot{\phi}_{rp}$ is the
precession angular velocity corresponding to that regular precession. Crabtree
obtains the same result with Routh [3].

Lamb follows a similar path to Routh’s approach and finds angular frequency
which can be written as [4]

$w_{L}^{2}=\dot{\phi}_{rp}^{2}\sin^{2}\theta_{rp}+(a-2\dot{\phi}_{rp}\cos\theta_{rp})^{2}.$
(27)

Greiner obtains angular frequency from the second derivative of effective
potential with respect to $\theta$, which can be written as [8]

$w_{G}^{2}=\frac{I_{x}ab-Mgl(4-3\sin^{2}\theta_{rp})}{I_{x}\cos\theta_{rp}}.$
(28)

One can show that these three results are the same as $w_{g}^{2}$ by using
equations (2), (4) and (5) by considering $\ddot{\theta}=0$ for regular
precession.

## References

  * [1]
  * [2] Routh E J 1955 Advanced Dynamics of a System of Rigid Bodies (New York: Dover) p.145
  * [3] Crabtree H 1909 An elementary treatment of the theory of spinning tops and gyroscopic motion (London: Longmans, Green and Co.) pp.123-124
  * [4] Lamb H 1920 Higher Mechanics (Cambridge: Cambridge University Press) pp.138-139
  * [5] Scarborough J B 1958 The Gyroscope Theory and Applications (London: Interscience Publishers) pp.75-80
  * [6] Goldstein H, Poole C and Safko J 2002 Classical Mechanics 3rd Ed (New York: Addison-Wesley) pp.208-217
  * [7] Arnold V I 1989 Mathematical Methods of Classical Mechanics 2nd Ed (New York: Springer-Verlag) pp.155-158
  * [8] Greiner W 2003 Classical Mechanics, Systems of particles and Hamiltonian dynamics (New York: Springer) p.266
  * [9] Tanrıverdi V 2019 Can a gyroscope reverse its spin direction? Eur. J. Phys. 40 065004 https://doi.org/10.1088/1361-6404/ab335d
  * [10] Thornton S T and Marion J B 2004 Classical dynamics of particles and systems 5th Ed (Belmont: Thomson Brooks/Cole)
  * [11] Tanrıverdi V 2021 Effective potential of a spinning heavy symmetric top when magnitudes of conserved angular momenta are not equal https://arxiv.org/abs/2101.06168
  * [12] Tanrıverdi V 2021 ”Motion of the heavy symmetric top when magnitudes of conserved angular momenta are different”, arXiv:2011.09348 https://arxiv.org/abs/2011.09348