Skip to content

UCT/CUDA_IPC: Skip zero-length elements in SGL zcopy - #12219

Open
dmitrygx wants to merge 12 commits into
openucx:masterfrom
dmitrygx:topic/ucx/cuda-ipc-sgl-zero-length
Open

dmitrygx wants to merge 12 commits into
openucx:masterfrom
dmitrygx:topic/ucx/cuda-ipc-sgl-zero-length

Conversation

@dmitrygx

@dmitrygx dmitrygx commented Nov 3, 2026 •

Copy link
Copy Markdown
Member

UCT/CUDA_IPC: Skip zero-length elements in SGL zcopy

What?

A cuda_ipc put_sgl_zcopy/get_sgl_zcopy with a zero-length element among non-empty ones fails with UCS_ERR_IO_ERROR. This PR leaves zero-length elements out of the copy batch, so such an SGL transfers the other elements and completes successfully, as on rc_mlx5.

Why?

uct_cuda_ipc_post_cuda_sgl_async_copy() skips only an SGL whose total length is zero. Otherwise it passes all the elements to cuMemcpyBatchAsync(), which rejects a copy of zero length:

cuda_ipc_ep.c:379  UCX  ERROR cuMemcpyBatchAsync_v2(dsts, srcs, (size_t *)lengths, count, &attr, &attrs_idx, 1, *stream) failed: invalid argument

The SGL zcopy API does not forbid zero-length elements, and rc_mlx5 allows them for put.

How?

  • UCT/CUDA_IPC: build the batch only from the elements with data: the mapping loop skips empty elements and stores the local address of the mapped remote memory, the buffer address and the length of each element with data (local_addrs, buffer_addrs and copy_lengths) in arrays allocated together with the SGL mapping, so mapping->count is the number of copies in the batch. The CUDA context and the stream are taken from the first element with data instead of element 0, which may be empty.
  • GTEST/UCT: add test_uct_sgl.zero_length_element, zero_length_first_last_elements and zero_length_all_elements: a 3-element SGL put/get with an empty middle element, with empty first and last elements, and with all elements empty. They check that the other elements are transferred, that the targets of the empty ones keep their data, and that the completion is not called when the operation completes in place. They are skipped where zero-length zcopy is not allowed (min_zcopy > 0, e.g. get on rc_mlx5).

Tests:

  • Test commit only, on master: cuda_ipc/test_uct_sgl.zero_length_element fails for put and get (Input/output error); rc_mlx5 put passes.
  • With the fix: cuda_ipc/test_uct_sgl.* and rc_mlx5/test_uct_sgl.* passed UCP all/test_ucp_rma_sgl.* passed.
  • cuda_ipc/* UCT tests, with the first version of the fix: all pass except test_uct_loopback_cuda.envelope ("Destination is unreachable" in the perf endpoint setup), which fails the same way without the fix on this machine.

Found while testing #12205 on cuda_ipc.

dmitrygx and others added 2 commits November 3, 2026 09:55
An element of an SGL put or get may be of zero length. Such an element
transfers no data, while the other elements of the SGL are transferred.
Check it on the transports that allow zero-length zcopy operations.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
cuMemcpyBatchAsync fails with an invalid argument error for a copy of
zero length, so an SGL put or get with a zero-length element failed with
UCS_ERR_IO_ERROR. Leave such elements out of the batch, and take the
CUDA context and the stream from the first element with data.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@dmitrygx
dmitrygx requested a review from a team as a code owner November 3, 2026 10:27
@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Starting review — findings will be posted here when done.

Comment thread test/gtest/uct/v2/test_uct_sgl.cc Outdated
void test_sgl_zero_length_element()
{
static const size_t count = 3;
static const size_t empty_index = 1;

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

pls cover a zero length in the first element too (e.g. loop the index over 0..count-1) - the new first selection in cuda_ipc only matters when lengths[0] == 0, and with empty_index = 1 it is never exercised.

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Added zero_length_first_last_elements with {0, 2} in eefee38: the first element with data is element 1, and the last element is empty too. I did not loop over the index, to keep one named case per scenario; with {1} and {0, 1, 2}, each position is empty in some test. All elements in the test have valid buffers and rkeys on the same GPU, so the test runs the first > 0 path and checks the data, but it would also pass if the context and the stream were taken from element 0.

🤖 Generated with Claude Code

Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Outdated
mapping = ucs_malloc(sizeof(*mapping) +
count * (sizeof(*mapping->entries) +
sizeof(CUdeviceptr)),
(2 * sizeof(CUdeviceptr)) + sizeof(size_t)),

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

maybe keep the common case cheap: when no length is zero we could pass buffers/lengths directly as before, and skip allocating and filling the two extra arrays per element.

@dmitrygx dmitrygx Nov 3, 2026 •

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in b661636: the sum loop also counts the elements with data. If none is empty, buffers and lengths are passed to cuMemcpyBatchAsync as on master, and the allocation has the same size as on master. Only if some element is empty, the allocation also has room for the compacted local addresses and lengths, which a separate loop fills.

Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Resolved
An SGL put or get with all elements of zero length transfers no data.
Check that it keeps the target buffers unchanged, and that the completion
is not called when the operation completes in place. Pass a completion to
the test of a single zero-length element as well.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Starting review — findings will be posted here when done.

Comment thread test/gtest/uct/v2/test_uct_sgl.cc Resolved
Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Resolved
Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Outdated
total_length = 0;
first = count;
for (i = 0; i < count; i++) {
if ((first == count) && (lengths[i] != 0)) {

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

minor: maybe find the first non-empty element right after the total_length == 0 check (for (first = 0; lengths[first] == 0; ++first) {})? The sum loop stays branchless and the first = count sentinel is not needed.

@dmitrygx dmitrygx Nov 3, 2026 •

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in b661636: the first element with data is found right after the total_length == 0 check, the sum loop has no branch, and the first = count sentinel is gone. The search is a while loop rather than a for with an empty body.

dmitrygx and others added 3 commits November 3, 2026 14:33
An SGL put or get may start or end with an element of zero length. With
an empty first element, cuda_ipc takes the CUDA context and the stream
from the first element with data, a path that the other tests do not
reach.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
If no element of an SGL put or get is of zero length, pass the local
buffers and the lengths of the caller to cuMemcpyBatchAsync as before,
without allocating and filling compacted copies of them. Find the first
element with data after the check for zero total length, and trace the
number of copies in the batch instead of the number of elements.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Each length of an SGL put or get must not be less than the minimal zcopy
length of the interface. A length can be 0 only if that minimum is 0,
and such an element transfers no data.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Starting review — findings will be posted here when done.

Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Outdated
elem_size += sizeof(*local_addrs) + sizeof(*copy_lengths);
}

mapping = ucs_malloc(sizeof(*mapping) + (count * elem_size),

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

minor: only num_nonempty elements are ever stored in these arrays, so maybe allocate num_nonempty * elem_size and offset mapped_addrs/local_addrs/copy_lengths by num_nonempty instead of count? The fast path is unchanged since there num_nonempty == count.

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in bf15f67: the mapping is allocated for num_nonempty elements, and mapped_addrs, local_addrs and copy_lengths are offset by num_nonempty.

Comment thread test/gtest/uct/v2/test_uct_sgl.cc Outdated
UCS_TEST_SKIP_R("not enough sgl elements");
}

if (elem_size(0) > 0) {

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

minor: elem_size(0) > 0 as a min_zcopy check is hard to read; can we add a small min_zcopy() helper and use it here and in elem_size()?

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in b118b78: added a min_zcopy() helper, used here and in elem_size().

@svc-nvidia-pr-review

Copy link
Copy Markdown

Residual coverage gaps (not posted as inline comments): the new tests only exercise zero-length elements with valid buffers/memhs/rkeys — the doc does not state whether those may be invalid for a zero-length element, and cuda_ipc (skips them entirely) and rc_mlx5 (still reads remote_addrs[i]/rkeys[i]) differ in what they touch. Also, zero-length SGL elements are currently unreachable from UCP (ucp_proto_put_sgl_offload_send_func breaks on length == 0), so this path is covered by gtest only, and only on hosts with CUDA >= 13 / cuda_ipc or IB put.

Comment thread src/uct/api/v2/uct_v2.h Resolved
The rule tied each SGL length to min_zcopy, which uct_iface_attr defines
for the total length of the iov of a zcopy operation. Skipping the empty
elements in cuda_ipc does not need such a new contract.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Starting review — findings will be posted here when done.

Only the elements with data are stored in the arrays of an SGL mapping,
so allocate and lay out the arrays for that number of elements instead
of for all the elements.
Skip the zero-length tests by the minimal zcopy length itself instead of
by the size of an empty element, and use the helper in elem_size() too.
Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Outdated
elem_size += sizeof(*local_addrs) + sizeof(*copy_lengths);
}

mapping = ucs_malloc(sizeof(*mapping) + (count * elem_size),

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

minor: can we use num_nonempty instead of count here and in the offsets below? the tail of every sub-array is unused when some elements are empty.

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in bf15f67: the mapping is allocated for num_nonempty elements, and the arrays are offset by num_nonempty.

Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Outdated
/* Same elements and order as in the mapping loop below */
local_addrs = mapped_addrs + count;
copy_lengths = (size_t *)(local_addrs + count);
for (i = 0, j = 0; i < count; i++) {

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

maybe fill local_addrs/copy_lengths inside the mapping loop below instead of here, so the two loops don't have to be kept in the same order? this path already does a ucs_malloc plus a memhandle cache lookup per element, so one extra store per element is not significant (docs/OptimizationStyle.md).

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in c3082be: local_addrs and copy_lengths are filled in the mapping loop together with mapped_addrs, so there is one loop over the elements.

@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Starting review — findings will be posted here when done.

Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Outdated
/* If some elements are empty, the local addresses and lengths of the
* non-empty ones are compacted into the same allocation, otherwise the
* caller arrays are passed to cuMemcpyBatchAsync as is */
elem_size = sizeof(*mapping->entries) + sizeof(*mapped_addrs);

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

maybe we could always fill local_addrs/copy_lengths inside the mapping loop at index mapping->count? That would drop the elem_size branch, the extra pre-pass and j, and the "same elements and order as in the mapping loop below" coupling, at the cost of 16 bytes and two stores per element, which look negligible next to the per-element memhandle mapping.

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in c3082be: local_addrs and copy_lengths are always allocated with the mapping and filled in the mapping loop at index mapping->count, so the elem_size branch, the separate loop and j are gone. The caller arrays are no longer passed to cuMemcpyBatchAsync as is, which costs two stores and 16 bytes of the mapping per element.

@svc-nvidia-pr-review

Copy link
Copy Markdown

The lengths doc in uct_v2.h still says only "Array of transfer lengths in bytes"; the new behavior (zero-length elements allowed, and only when cap.{put,get}.min_zcopy == 0 — which is exactly what the test's min_zcopy() > 0 skip encodes) stays undocumented.

rc_mlx5 GET rejects zero-length elements via UCT_CHECK_LENGTH(lengths[i], max_inl_cqe+1, ...), so zero_length_* only runs for PUT there; on cuda_ipc it needs CUDA >= 13. Zero-length SGL elements are still unreachable from UCP (ucp_proto_put_sgl_offload_send_func breaks on length == 0), so gtest is the only coverage.

Pre-existing and out of scope: when cap.get.max_zcopy == 0 (cuda_ipc without NVLink) elem_size() returns 0 while UCT_IFACE_FLAG_V2_GET_SGL_ZCOPY is still advertised, so all test_uct_sgl GET cases degenerate to zero-length elements instead of skipping.

Store the local address and the length of each element with data in the
mapping loop, next to its remote address, instead of compacting them in
a separate loop only if some element is empty. This keeps one loop over
the elements, at the cost of two stores and 16 bytes of the mapping per
element, which is small next to the memory handle lookup per element.
@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Starting review — findings will be posted here when done.

@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Review complete — no issues found.

tvegas1
tvegas1 previously approved these changes Nov 4, 2026
@dmitrygx
dmitrygx enabled auto-merge November 4, 2026 11:31
@svc-ucx

svc-ucx commented Nov 4, 2026

Copy link
Copy Markdown

🤖 CI Triage Agent — UCX PR (AddressSanitizer BlueField on worker 3) · commit c3082be2

TL;DR: The gtest dc_mlx5/test_rc_iface_flush_remote.put_fence_no_flush_remote/0 hung for ~3m20s waiting for a fenced uct_ep_put_zcopy completion that never arrived, then crashed in teardown because uct_rc_txqp_purge_outstanding() invoked a user completion that lives on the test's (already unwound) stack — a lost DC fenced-PUT completion, not anything in this PR.

Full analysis

Summary: dc_mlx5/test_rc_iface_flush_remote.put_fence_no_flush_remote/0 failed (comp.count still 1 at test/gtest/uct/ib/test_rc.cc:776) and then aborted on Assertion 'comp->count > 0' failed ... count=-360728992 in uct_invoke_completion(), killing the whole gtest run (make: *** [test] Aborted (core dumped)).

Root cause: Two layered problems, both visible in the log:

  1. Lost completion (the real bug). The test posts uct_ep_fence() + uct_ep_put_zcopy() on a dc_mlx5 EP with IB_PCI_RELAXED_ORDERING?=try; the put returned UCS_INPROGRESS and then nothing happened for 3m20s (last test output 13:19:21, failure at 13:22:41 — the gap is the wait_for_value poll loop, so the operation hung, no CQE was ever reaped). Teardown confirms it: destroying txqp 0xffff7c046090 with uncompleted operation 0xffff75f3f800 handler uct_rc_ep_send_op_completion_handler (rc_ep.c:550). So the fenced PUT WQE on the DCI never produced a signaled completion.
  2. Test-harness amplification. comp is a stack variable in the test body (test_rc.cc:765). After the EXPECT failures, uct_test::cleanup() → uct_dc_mlx5_ep_t_cleanup (dc_mlx5_ep.c:1459) → uct_rc_txqp_purge_outstanding() → uct_invoke_completion(op->user_comp, ...) (rc_ep.c:566) dereferences that now-dead stack object, hence count=-360728992 and the fatal assert at src/uct/base/uct_iface.h:1012. That turns a soft test failure into a core dump that aborts the entire suite.

This is unrelated to PR #12219 (topic/ucx/cuda-ipc-sgl-zero-length touches cuda_ipc, not the IB/DC path) — treat it as a pre-existing/infra failure on the BlueField worker.

Implicated commit: unknown for the lost completion; most plausible recent candidate is 22cff74b "UCT/RC: Reserve TX credits for multi-credit sends (#11841)" (Michal Shalev, 2026-11-02 — two days before this build), which changed TX credit/CQE reservation (uct_rc_ep_reserve_txqp/uct_rc_ep_reserve_cqe, src/uct/ib/rc/base/rc_ep.h:393-435) and also edited this very test file. The DC fence path itself was last touched by 73a6a5c9 (relaxed-only fence ordering) and 827c995b (put fence for DDP).

File: test/gtest/uct/ib/test_rc.cc:765-776 (stack uct_completion_t); src/uct/ib/rc/base/rc_ep.c:566; src/uct/base/uct_iface.h:1012

Suggested fix:

  1. Make the test robust so this failure mode doesn't core-dump the suite: move comp out of the stack frame (class member / heap, or flush(CANCEL) + wait on the EP before TearDown) so uct_rc_txqp_purge_outstanding() can't write to a dead object. Add EXPECT_EQ with a bounded timeout and ASSERT so the test exits cleanly.
  2. For the hang itself: re-run the job to confirm reproducibility, then instrument the dc_mlx5 fenced put_zcopy path on BlueField with UCX_LOG_LEVEL=trace_data and check whether the fence/NOP WQE consumed a CQE slot without a signaled completion — i.e. verify cq_reserve/txqp_reserve accounting from UCT/RC: Reserve TX credits for multi-credit sends #11841 against the multi-WQE fence post, and that the doorbell is rung after the fence WQE.
  3. If the DC fenced-PUT stall reproduces, consider reverting/guarding UCT/RC: Reserve TX credits for multi-credit sends #11841's reservation logic for the fence path until fixed.

Related: #11841 (and its v1.23.x backport #12222); fence-path changes #11721, #10658; purge-outstanding series #11866/#11874/#12133.

Resolve the conflict in cuda_ipc_ep.c with the renames on master: the
array of the local addresses of the remote memory is now local_addrs, so
the array of the buffer addresses of the elements with data is renamed
to buffer_addrs. The comment and the sizeof() terms of the mapping
allocation, which merged without a conflict, are updated the same way.
@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Starting review — findings will be posted here when done.

@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Review complete — no issues found.

tvegas1
tvegas1 previously approved these changes Nov 4, 2026
@dmitrygx
dmitrygx disabled auto-merge November 4, 2026 15:00
@dmitrygx
dmitrygx enabled auto-merge (squash) November 4, 2026 15:00
@svc-ucx

svc-ucx commented Nov 5, 2026

Copy link
Copy Markdown

🤖 CI Triage Agent — UCX PR (Cuda Test Cuda Docker centos7_cuda_11_4) · commit 3ff81b8b

TL;DR: The Memory hooks test step failed because cudaMalloc() in test/apps/test_cuda_hook_dynamic returned CUDA error 807 (cudaErrorMpsServerNotReady) on the GPU agent, so 0/2 memory events were recorded — this is a GPU/MPS-daemon environment failure on the ucx_gpu_test host, unrelated to PR #12219; re-run on a healthy agent and harden the CI script to detect an unusable CUDA runtime as an infra error.

Full analysis

Summary: Azure job "Cuda Test Cuda Docker centos7_cuda_11_4" (build 142791) failed in the Memory hooks test step: ./test/apps/test_cuda_hook_dynamic printed cudaMalloc() failed: 807 / got 0/2 memory events and exited 255 (Bash exited with code '255').

Root cause: The build itself completed cleanly (all CC/CXXLD steps succeeded, no gap/hang in timestamps — compile ran 01:48→01:56 continuously, failure at 01:56:52 immediately after + ./test/apps/test_cuda_hook_dynamic). The test only does ucp_init(), installs UCM memtype event handlers, then calls cudaMalloc(4096). The allocation was rejected by the CUDA driver with status 807, which in the CUDA runtime/driver enums is cudaErrorMpsServerNotReady / CUDA_ERROR_MPS_SERVER_NOT_READY — i.e. the CUDA client could not get a usable context/allocation from the driver (MPS server on the GPU host not accepting clients). The zero memory events are a consequence, not a cause: no allocation happened, so no UCM_EVENT_MEM_TYPE_ALLOC callback fired. Nothing in the hook path (UCX_MEM_CUDA_HOOK_MODE=reloc) can turn a working cudaMalloc into an 80x driver-level error. check_gpu 'gpu' passed (so nvidia-smi -L saw a GPU) and check_cuda_driver passed (/proc/driver/nvidia/version exists), which is exactly the blind spot: both checks pass while the CUDA runtime is still unusable. PR #12219 ("UCT/CUDA_IPC: Skip zero-length elements in SGL zcopy") touches only cuda_ipc SGL zcopy and is not exercised by this standalone hook test.

Implicated commit: unknown — not a code regression; the failure is environmental (GPU agent / MPS state). Test app test/apps/test_cuda_hook.c last changed in 1c6ddd6 (rakhmets, 2025‑05‑30), long before this PR.

File: buildlib/pr/cuda/test_malloc_hook.sh:22 (check_cuda_driver, insufficient health check) and test/apps/test_cuda_hook.c:74 (failing cudaMalloc)

Suggested fix:

  1. Re-run the Cuda stage; if other matrix legs (*_cuda_11_x, *_cuda_12_x) failed the same way, the agent is bad — have the host's MPS control daemon/driver state fixed (restart nvidia-cuda-mps-control, or clear CUDA_MPS_PIPE_DIRECTORY/CUDA_MPS_* from the agent environment so containers talk to the driver directly).
  2. Harden buildlib/pr/cuda/test_malloc_hook.sh: extend check_cuda_driver() to actually validate the runtime before running the hook tests (e.g. nvidia-smi -q | grep -i mps, dump cat /proc/driver/nvidia/version and env | grep CUDA_MPS, and run a trivial cudaMalloc probe). On failure, call azure_log_error "CUDA runtime unusable on $(hostname -s)" so it is reported as an infrastructure error instead of a PR test failure.
  3. Optionally improve diagnostics in test/apps/test_cuda_hook.c:76 to print cudaGetErrorString(res) alongside the numeric code, so future logs say "MPS server not ready" instead of just 807.

Related: PR #12219 (#12219) — failure is unrelated to its change; no existing issue found for CUDA error 807 in CI.

Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Resolved
Count the empty elements before the first element with data in the loop
that sums the lengths, instead of searching for that element after the
check for zero total length, so that the search does not rely on the
total length to stay in bounds.
@svc-nvidia-pr-review

Copy link
Copy Markdown

🤖 Starting review — findings will be posted here when done.

Comment thread src/uct/cuda/cuda_ipc/cuda_ipc_ep.c Resolved
@dmitrygx

dmitrygx commented Nov 5, 2026

Copy link
Copy Markdown
Member Author

/azp run

@azure-pipelines

Copy link
Copy Markdown
Azure Pipelines successfully started running 4 pipeline(s).

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants