Skip to content

Commit d865179

Browse files
DEVICE/API: Add nixlGpuWriteSignal API for local signal writing (#805)
1 parent 3daa987 commit d865179

File tree

1 file changed

+16
-0
lines changed

1 file changed

+16
-0
lines changed

src/api/gpu/ucx/nixl_device.cuh

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -253,4 +253,20 @@ nixlGpuReadSignal(const void *signal) {
253253
return ucp_device_counter_read<static_cast<ucs_device_level_t>(level)>(signal);
254254
}
255255

256+
/**
257+
* @brief Write value to the local signal.
258+
*
259+
* This function can be used to set a signal to a specific value.
260+
*
261+
* The signal must be initialized with the host function @ref prepGpuSignal.
262+
*
263+
* @param signal [in,out] Address of the signal.
264+
* @param value [in] Value to write to the signal.
265+
*/
266+
template<nixl_gpu_level_t level = nixl_gpu_level_t::THREAD>
267+
__device__ void
268+
nixlGpuWriteSignal(void *signal, uint64_t value) {
269+
ucp_device_counter_write<static_cast<ucs_device_level_t>(level)>(signal, value);
270+
}
271+
256272
#endif // _NIXL_DEVICE_CUH

0 commit comments

Comments
 (0)