Signaling Operations

NVSHMEM provides signaling operations that can be used to update a remote flag variable. When used in conjunction with wait/test routines at the remote PE, these routines can provide efficient point-to-point synchronization. The following example shows signal operations used to implement neighbor communication in a ring.

nvshmem_putmem(dest, src, size, (pe+1) % npes);
nvshmem_quiet();
nvshmemx_int_signal(flag, 1, (pe+1) % npes);
nvshmem_int_wait_until(flag, NVSHMEM_CMP_EQ, 1);

This section specifies the NVSHMEM support for put-with-signal, nonblocking put-with-signal, and signal-fetch routines. The put-with-signal routines provide a method for copying data from a contiguous local data object to a data object on a specified PE and subsequently updating a remote flag to signal completion. The signal-fetch routine provides support for fetching a signal update operation.

The NVSHMEM counted signaling extensions copy a payload and increment a remote counter by the number of bytes delivered. They also provide operations to reset, load, and wait on the counted signal data object. See Using Counted Writes with NVSHMEM for an overview and usage example.

Atomicity Guarantees for Signaling Operations

All signaling operations put-with-signal, nonblocking put-with-signal, and signal-fetch are performed on a signal data object, a remotely accessible symmetric object of type uint64_t. A signal operator in the put-with-signal routine is a NVSHMEM library constant that determines the type of update to be performed as a signal on the signal data object.

All signaling operations complete as if performed atomically with respect to the following:

• other signal operations that update the signal data object using the same datatype;

• signal-fetch routine that fetches the signal data object; and

• any point-to-point synchronization routine that accesses the signal data object using the same datatype.

Available Signal Operators

With the atomicity guarantees as described in Section Atomicity Guarantees for Signaling Operations, the following options can be used as a signal operator.

NVSHMEM_SIGNAL_SET

An update to signal data object is an atomic set operation. It writes an unsigned 64-bit value as a signal into the signal data object on a remote PE as an atomic operation.

NVSHMEM_SIGNAL_ADD

An update to signal data object is an atomic add operation. It adds an unsigned 64-bit value as a signal into the signal data object on a remote PE as an atomic operation.

NVSHMEM_PUT_SIGNAL

void nvshmem_TYPENAME_put_signal(

TYPE *dest,

const TYPE *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe

)

void nvshmemx_TYPENAME_put_signal_on_stream(

TYPE *dest,

const TYPE *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe,

cudaStream_t stream

)

__device__ void nvshmem_TYPENAME_put_signal(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_TYPENAME_put_signal_block(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_TYPENAME_put_signal_warp(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

where TYPE is one of the standard RMA types and has a corresponding TYPENAME specified by Table Standard RMA Types and Names.

void nvshmem_putSIZE_signal(

void *dest,

const void *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe

)

void nvshmemx_putSIZE_signal_on_stream(

void *dest,

const void *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe,

cudaStream_t stream

)

__device__ void nvshmem_putSIZE_signal(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_putSIZE_signal_block(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_putSIZE_signal_warp(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

where SIZE is one of 8, 16, 32, 64, 128.

void nvshmem_putmem_signal(

void *dest,

const void *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe

)

void nvshmemx_putmem_signal_on_stream(

void *dest,

const void *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe,

cudaStream_t stream

)

__device__ void nvshmem_putmem_signal(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_putmem_signal_block(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_putmem_signal_warp(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

Symmetric address of the data object to be updated on the remote PE. The type of dest should match that implied in the SYNOPSIS section.

Symmetric address or host/device address registered via nvshmemx_buffer_register of data object containing the data to be copied. The type of source should match that implied in the SYNOPSIS section.Additionally, it can also be backed by device shared memory when devices are connected via peer-to-peer transport.

Number of elements in the dest and source arrays. For nvshmem_putmem_signal elements are bytes.

Symmetric address of the signal data object to be updated on the remote PE as a signal.

Unsigned 64-bit value that is used for updating the remote sig_addr signal data object.

Signal operator that represents the type of update to be performed on the remote sig_addr signal data object.

PE number of the remote PE.

Description

The put-with-signal routines provide a method for copying data from a contiguous local data object to a data object on a specified PE and subsequently updating a remote flag to signal completion. The routines return after the data has been copied out of the source array on the local PE.

The sig_op signal operator determines the type of update to be performed on the remote sig_addr signal data object. The completion of signal update based on the sig_op signal operator using the signal flag on the remote PE indicates the delivery of its corresponding dest data words into the data object on the remote PE.

An update to the sig_addr signal data object through a put-with-signal routine completes as if performed atomically as described in Section Atomicity Guarantees for Signaling Operations. The various options as described in Section Available Signal Operators can be used as the sig_op signal operator.

Returns

None.

Notes

The dest and sig_addr data objects must both be remotely accessible.

sig_addr and dest may not be overlapping in memory.

The completion of signal update using the signal flag on the remote PE indicates only the delivery of its corresponding dest data words into the data object on the remote PE. Without a memory-ordering operation, there is no implied ordering between the signal update of a put-with-signal routine and another data transfer. For example, the completion of the signal update in a sequence consisting of a put routine followed by a put-with-signal routine does not imply delivery of the put routine’s data.

The following CUDA C++ example uses nvshmem_put_signal to implement a broadcast from PE 0 to itself and all other PEs in the job as a simple ring-based algorithm:

#include <cuda_runtime.h>
#include <nvshmem.h>
#include <nvshmemx.h>
#include <stdint.h>

#define MESSAGE_SIZE 2048

__global__ void ring_put_signal(uint64_t *data, uint64_t *signal) {
    int mype = nvshmem_my_pe();

    if (mype == 0) {
        for (size_t i = 0; i < MESSAGE_SIZE; i++) {
            data[i] = i;
        }
        *signal = 1;
    }

    nvshmem_signal_wait_until(signal, NVSHMEM_CMP_EQ, 1);
    if (mype != nvshmem_n_pes() - 1) {
        nvshmem_uint64_put_signal(data, data, MESSAGE_SIZE, signal, 1, NVSHMEM_SIGNAL_SET,
                                  mype + 1);
    }
}

int main(void) {
    nvshmem_init();
    cudaSetDevice(nvshmem_team_my_pe(NVSHMEMX_TEAM_NODE));

    uint64_t *data = (uint64_t *)nvshmem_malloc(MESSAGE_SIZE * sizeof(uint64_t));
    uint64_t *signal = (uint64_t *)nvshmem_calloc(1, sizeof(uint64_t));
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    void *args[] = {&data, &signal};
    nvshmemx_collective_launch((const void *)ring_put_signal, dim3(1), dim3(1), args, 0,
                               stream);
    cudaStreamSynchronize(stream);

    cudaStreamDestroy(stream);
    nvshmem_free(signal);
    nvshmem_free(data);
    nvshmem_finalize();
    return 0;
}

NVSHMEM_PUT_SIGNAL_NBI

void nvshmem_TYPENAME_put_signal_nbi(

TYPE *dest,

const TYPE *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe

)

void nvshmemx_TYPENAME_put_signal_nbi_on_stream(

TYPE *dest,

const TYPE *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe,

cudaStream_t stream

)

__device__ void nvshmem_TYPENAME_put_signal_nbi(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_TYPENAME_put_signal_nbi_block(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_TYPENAME_put_signal_nbi_warp(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

where TYPE is one of the standard RMA types and has a corresponding TYPENAME specified by Table Standard RMA Types and Names.

void nvshmem_putSIZE_signal_nbi(

void *dest,

const void *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe

)

void nvshmemx_putSIZE_signal_nbi_on_stream(

void *dest,

const void *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe,

cudaStream_t stream

)

__device__ void nvshmem_putSIZE_signal_nbi(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_putSIZE_signal_nbi_block(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_putSIZE_signal_nbi_warp(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

where SIZE is one of 8, 16, 32, 64, 128.

void nvshmem_putmem_signal_nbi(

void *dest,

const void *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe

)

void nvshmemx_putmem_signal_nbi_on_stream(

void *dest,

const void *source,

size_t nelems,

uint64_t *sig_addr,

uint64_t signal,

int sig_op,

int pe,

cudaStream_t stream

)

__device__ void nvshmem_putmem_signal_nbi(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_putmem_signal_nbi_block(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

__device__ void nvshmemx_putmem_signal_nbi_warp(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

Symmetric address of the data object to be updated on the remote PE. The type of dest should match that implied in the SYNOPSIS section.

Symmetric address or host/device address registered via nvshmemx_buffer_register of data object containing the data to be copied. The type of source should match that implied in the SYNOPSIS section.Additionally, it can also be backed by device shared memory when devices are connected via peer-to-peer transport.

Number of elements in the dest and source arrays. For nvshmem_putmem_signal_nbi and nvshmem_ctx_putmem_signal_nbi, elements are bytes.

Symmetric address of the signal data object to be updated on the remote PE as a signal.

Unsigned 64-bit value that is used for updating the remote sig_addr signal data object.

Signal operator that represents the type of update to be performed on the remote sig_addr signal data object.

PE number of the remote PE.

Description

The nonblocking put-with-signal routines provide a method for copying data from a contiguous local data object to a data object on a specified PE and subsequently updating a remote flag to signal completion.

The routines return after initiating the operation. The operation is considered complete after a subsequent call to nvshmem_quiet. At the completion of nvshmem_quiet, the data has been copied out of the source array on the local PE and delivered into the dest array on the destination PE.

The delivery of signal flag on the remote PE indicates only the delivery of its corresponding dest data words into the data object on the remote PE. Furthermore, two successive nonblocking put-with-signal routines, or a nonblocking put-with-signal routine with another data transfer may deliver data out of order unless a call to nvshmem_fence is introduced between the two calls.

The sig_op signal operator determines the type of update to be performed on the remote sig_addr signal data object.

An update to the sig_addr signal data object through a nonblocking put-with-signal routine completes as if performed atomically as described in Section Atomicity Guarantees for Signaling Operations. The various options as described in Section Available Signal Operators can be used as the sig_op signal operator.

Returns

None.

Notes

The dest and sig_addr data objects must both be remotely accessible.

sig_addr and dest may not be overlapping in memory.

NVSHMEM_SIGNAL_FETCH

__device__ uint64_t nvshmem_signal_fetch(const uint64_t *sig_addr)

sig_addr [IN]

Local address of the remotely accessible signal variable.

Description

nvshmem_signal_fetch performs a fetch operation and returns the contents of the sig_addr signal data object. Access to sig_addr signal object at the calling PE is expected to satisfy the atomicity guarantees as described in Section Atomicity Guarantees for Signaling Operations.

Returns

Returns the contents of the signal data object, sig_addr, at the calling PE.

NVSHMEMX_SIGNAL_COUNTED_RESET

void nvshmemx_signal_counted_reset(uint64_t *signal_addr)

__device__ void nvshmemx_signal_counted_reset(uint64_t *signal_addr)

signal_addr [OUT]

Symmetric address of the counted signal data object on the calling PE.

Description

The nvshmemx_signal_counted_reset routine stores zero in the counted signal data object at signal_addr on the calling PE. Applications can use this routine to initialize a counted signal or to begin a new protocol epoch after all operations from the preceding epoch have completed.

Returns

None.

Notes

signal_addr must point to an 8-byte uint64_t object in symmetric memory. Its offset in the symmetric heap must be 256-byte aligned. Otherwise, the behavior is undefined.

Device code may initialize the counter with an ordinary local store when no counted write or waiter can access it concurrently. Host code should use this routine instead of dereferencing signal_addr because the symmetric counter must reside in device memory.

This routine does not synchronize with producers, remote counted writes, or threads waiting on the counter. Resetting a counted signal while a counted write targeting it is in flight, or while a receiver is waiting on an epoch derived from its previous value, results in undefined behavior.

NVSHMEMX_SIGNAL_COUNTED_LOAD

__device__ uint64_t nvshmemx_signal_counted_load(const uint64_t *signal_addr)

signal_addr [IN]

Symmetric address of the counted signal data object on the calling PE.

Description

The nvshmemx_signal_counted_load routine returns the current value of the counted signal data object at signal_addr on the calling PE. Applications can use the returned value for a subsequent byte-count epoch.

Returns

Returns the current value of the counted signal data object.

Notes

signal_addr must point to an 8-byte uint64_t object in symmetric memory. Its offset in the symmetric heap must be 256-byte aligned. Otherwise, the behavior is undefined.

A plain dereference of signal_addr is not a substitute for this routine when observing an asynchronously updated counter. This routine uses a volatile load so that each call reads the counter.

This routine only loads the counter. Applications should use nvshmemx_signal_counted_wait_until when payload visibility is required.

NVSHMEMX_SIGNAL_COUNTED_WAIT_UNTIL

__device__ void nvshmemx_signal_counted_wait_until(const uint64_t *signal_addr, uint64_t expected)

signal_addr [IN]

Symmetric address of the counted signal data object on the calling PE.

expected [IN]

The byte-count epoch for which to wait.

Description

The nvshmemx_signal_counted_wait_until routine blocks until the counted signal at signal_addr has reached or passed expected. After the condition is satisfied, the routine guarantees visibility of the payload associated with the observed counted writes.

Counter comparison uses modulo-\(2^{64}\) ordering. A counter value is ready when the unsigned difference between the counter and expected is less than \(2^{63}\).

Returns

None.

Notes

signal_addr must point to an 8-byte uint64_t object in symmetric memory. Its offset in the symmetric heap must be 256-byte aligned. Otherwise, the behavior is undefined.

Directly polling signal_addr is not equivalent to this routine because it does not guarantee payload visibility.

Applications must keep the number of outstanding bytes between observed epochs below \(2^{63}\). Each producer must contribute its agreed byte count exactly once to the epoch.

NVSHMEMX_PUTMEM_SIGNAL_COUNTED_NBI_BLOCK

__device__ int nvshmemx_putmem_signal_counted_nbi_block(void *dest, const void *source, size_t bytes, uint64_t *signal_addr, int pe)

dest [OUT]

Symmetric address of the destination data object on PE pe.

source [IN]

Address of the local data object containing the bytes to be copied. The object may reside in global or shared memory.

bytes [IN]

Number of bytes to copy from source to dest.

signal_addr [OUT]

Symmetric address of the counted signal data object to be incremented on PE pe.

pe [IN]

Number of the destination PE.

Description

The nvshmemx_putmem_signal_counted_nbi_block routine is a collective device operation over the calling CTA. All threads in the CTA must call the routine with the same arguments and in the same control flow.

When the operation is accepted, the routine copies bytes bytes from source to dest on PE pe. After those destination bytes are delivered, it atomically adds bytes to signal_addr on the same PE. Multiple counted writes may contribute to one counted signal, including writes from different producers, provided that the application accounts for every contribution in the expected byte-count epoch.

The routine returns after local submission and completion work needed to make the source or internal staging buffer reusable. The receiver uses nvshmemx_signal_counted_wait_until to wait for destination visibility. Completion of a counted signal epoch orders only the payload bytes covered by the corresponding counted writes. It does not order unrelated communication.

Returns

• NVSHMEMX_SUCCESS if the operation is accepted.

• NVSHMEMX_ERROR_INVALID_VALUE if pe, an address, an address range, an alignment, the payload size, or the relationship between dest and signal_addr is invalid.

• NVSHMEMX_ERROR_NOT_SUPPORTED if CFT counted operations are unavailable for the build, device, destination PE, or registered shared-memory configuration.

• NVSHMEMX_ERROR_INTERNAL if an unexpected CFT operation or completion error occurs.

Notes

This routine requires CUDA Compute Fabric Transport (CFT) counted-operation support. CFT handles and logical endpoints must be enabled, TMA must be enabled, and dest and signal_addr must be reachable through the same counted-capable logical endpoint. The routine does not fall back to an ordinary put-with-signal operation or a network transport.

dest and signal_addr must be symmetric addresses. source and dest must be 16-byte aligned, and bytes must be a multiple of 16. signal_addr must point to an 8-byte uint64_t object whose offset in the symmetric heap is 256-byte aligned. The counted signal must not overlap the destination payload.

Every issuing CTA must register shared memory as described in Using TMA with NVSHMEM. A shared-memory source requires at least NVSHMEMX_SMEM_BARRIERS_ONLY. A global-memory source requires NVSHMEMX_SMEM_MINIMUM or NVSHMEMX_SMEM_RECOMMENDED for internal staging.

The counted signal is a monotonically increasing byte counter during a protocol epoch. An application must not reset it while a counted write is in flight or while a receiver is waiting on a value from the current epoch.

The following datatypes are supported by the NVSHMEM signal operation.

Signal Types and Names

TYPE	TYPENAME
short	short
int	int
long	long
long long	longlong
unsigned short	ushort
unsigned int	uint
unsigned long	ulong
unsigned long long	ulonglong
int32_t	int32
int64_t	int64
uint32_t	uint32
uint64_t	uint64
size_t	size
ptrdiff_t	ptrdiff

NVSHMEMX_SIGNAL

Deprecated. Refer to nvshmemx_signal_op or nvshmem_*_atomic_set.

__device__ inline void nvshmemx_TYPENAME_signal(TYPE *dest, const TYPE value, int pe)

where TYPE is one of the standard RMA types and has a corresponding TYPENAME specified by Table Signal Types and Names.

dest [OUT]

Symmetric address of the signal word to be updated.

value [IN]

The value to be placed in dest.

pe [IN]

PE number of the remote PE.

Description

The nvshmemx_signal operation atomically sets dest to value on the specified PE. This operation can be used together with wait and test routines for efficient point-to-point synchronization.

Returns

None.

NVSHMEMX_SIGNAL_OP

__device__ inline void nvshmemx_signal_op(uint64_t *sig_addr, uint64_t signal, int sig_op, int pe)

sig_addr [OUT]

Symmetric address of the signal word to be updated.

signal [IN]

The value used to update sig_addr.

sig_op [IN]

Operation used to update sig_addr with signal.

pe [IN]

PE number of the remote PE.

Description

The nvshmemx_signal_op operation atomically updates sig_addr with signal using operation sig_op on the specified PE. This operation can be used together with wait and test routines for efficient point-to-point synchronization.

Returns

None.