From c545724abf460803c991e0279425b523a63212d2 Mon Sep 17 00:00:00 2001 From: Adrian Ostrowski Date: Fri, 25 Sep 2026 16:17:47 +0200 Subject: [PATCH] [xpupti] Collect peer-to-peer memory copies PTI reports a copy between two devices only through PTI_VIEW_DEVICE_GPU_MEM_COPY_P2P, a view xpupti never enabled, so such copies were missing from the trace altogether. Enable it along with the ordinary copy view for GPU_MEMCPY and log its records as gpu_memcpy on the source device, with the fromDevice/toDevice metadata that CUDA peer-to-peer copies carry. Co-Authored-By: Claude Opus 5.5 (1M context) --- libkineto/include/MetadataFieldCatalog.h | 2 + .../src/plugin/xpupti/XpuptiActivityApi.cpp | 47 +++++++++++------- .../plugin/xpupti/XpuptiActivityHandlers.cpp | 22 ++++++++- .../xpupti/XpuptiActivityProfilerSession.h | 2 + .../xpupti/XpuptiActivityHandlersTest.cpp | 41 ++++++++++++++++ libkineto/test/xpupti/XpuptiProfilerTest.cpp | 49 +++++++++++++++++++ libkineto/test/xpupti/XpuptiTestUtilities.cpp | 28 ----------- libkineto/test/xpupti/XpuptiTestUtilities.h | 31 ++++++++++++ libkineto/test/xpupti/compute/CMakeLists.txt | 4 +- .../xpupti/compute/XpuptiPeerCopyCompute.cpp | 32 ++++++++++++ .../xpupti/compute/XpuptiPeerCopyCompute.h | 16 ++++++ 11 files changed, 224 insertions(+), 50 deletions(-) create mode 100644 libkineto/test/xpupti/compute/XpuptiPeerCopyCompute.cpp create mode 100644 libkineto/test/xpupti/compute/XpuptiPeerCopyCompute.h diff --git a/libkineto/include/MetadataFieldCatalog.h b/libkineto/include/MetadataFieldCatalog.h index 1c85bcafb..604ae05c5 100644 --- a/libkineto/include/MetadataFieldCatalog.h +++ b/libkineto/include/MetadataFieldCatalog.h @@ -187,6 +187,7 @@ inline constexpr MetadataField kCorrelation{"correlation"}; inline constexpr MetadataField kDevice{"device"}; inline constexpr MetadataField kEngineIndex{"engine_index"}; inline constexpr MetadataField kEngineOrdinal{"engine_ordinal"}; +inline constexpr MetadataField kFromDevice{"fromDevice"}; inline constexpr MetadataField kKernelId{"kernel_id"}; inline constexpr MetadataField kMemoryBandwidthGbps{ "memory bandwidth (GB/s)"}; @@ -203,4 +204,5 @@ inline constexpr MetadataField kSyclInvocationId{ "sycl_invocation_id"}; inline constexpr MetadataField kSyclNodeId{"sycl_node_id"}; inline constexpr MetadataField kSyclQueue{"sycl queue"}; +inline constexpr MetadataField kToDevice{"toDevice"}; } // namespace libkineto::XpuMetadataFields diff --git a/libkineto/src/plugin/xpupti/XpuptiActivityApi.cpp b/libkineto/src/plugin/xpupti/XpuptiActivityApi.cpp index 540384c45..d68ef6b53 100644 --- a/libkineto/src/plugin/xpupti/XpuptiActivityApi.cpp +++ b/libkineto/src/plugin/xpupti/XpuptiActivityApi.cpp @@ -15,8 +15,8 @@ #include #include #include -#include #include +#include #include #include @@ -203,32 +203,39 @@ void warnIfIttNotifyLibInvalid() noexcept { } } -// The PTI view an activity type is collected through, if it has one of its own. -// Shared by enable and disable so the two cannot drift apart. -constexpr std::optional findPtiViewKind(ActivityType activity) { +template +constexpr pti_view_kind kViewKinds[] = {kViews...}; + +// The PTI views an activity type is collected through, if it has any of its +// own. Shared by enable and disable so the two cannot drift apart. +constexpr std::span findPtiViewKinds( + ActivityType activity) { switch (activity) { case ActivityType::GPU_MEMCPY: - return PTI_VIEW_DEVICE_GPU_MEM_COPY; + // PTI reports a peer-to-peer copy only through its own view. + return kViewKinds< + PTI_VIEW_DEVICE_GPU_MEM_COPY, + PTI_VIEW_DEVICE_GPU_MEM_COPY_P2P>; case ActivityType::GPU_MEMSET: - return PTI_VIEW_DEVICE_GPU_MEM_FILL; + return kViewKinds; case ActivityType::CONCURRENT_KERNEL: - return PTI_VIEW_DEVICE_GPU_KERNEL; + return kViewKinds; case ActivityType::EXTERNAL_CORRELATION: - return PTI_VIEW_EXTERNAL_CORRELATION; + return kViewKinds; case ActivityType::XPU_RUNTIME: - return PTI_VIEW_RUNTIME_API; + return kViewKinds; case ActivityType::XPU_DRIVER: - return PTI_VIEW_DRIVER_API; + return kViewKinds; case ActivityType::OVERHEAD: - return PTI_VIEW_COLLECTION_OVERHEAD; + return kViewKinds; case ActivityType::XPU_SYNC: - return PTI_VIEW_DEVICE_SYNCHRONIZATION; + return kViewKinds; case ActivityType::COLLECTIVE_COMM: - return PTI_VIEW_COMMUNICATION; + return kViewKinds; default: // XPU_SCOPE_PROFILER among others: enabled by // XpuptiScopeProfilerSession's constructor, not by a view. - return std::nullopt; + return {}; } } } // namespace @@ -240,11 +247,13 @@ void XpuptiActivityApi::enableXpuptiActivities( externalCorrelationEnabled_ = false; selected_activities.forEach([this](ActivityType activity) { - const auto view = findPtiViewKind(activity); - if (not view) { + const auto views = findPtiViewKinds(activity); + if (views.empty()) { return; } - XPUPTI_CALL(ptiViewEnable(*view)); + for (const auto view : views) { + XPUPTI_CALL(ptiViewEnable(view)); + } // What enabling a view takes on top of the view itself. switch (activity) { @@ -270,8 +279,8 @@ void XpuptiActivityApi::enableXpuptiActivities( void XpuptiActivityApi::disablePtiActivities( const ActivityTypeMask& selected_activities) { selected_activities.forEach([](ActivityType activity) { - if (const auto view = findPtiViewKind(activity)) { - XPUPTI_CALL(ptiViewDisable(*view)); + for (const auto view : findPtiViewKinds(activity)) { + XPUPTI_CALL(ptiViewDisable(view)); } }); externalCorrelationEnabled_ = false; diff --git a/libkineto/src/plugin/xpupti/XpuptiActivityHandlers.cpp b/libkineto/src/plugin/xpupti/XpuptiActivityHandlers.cpp index 19c75a7fb..4f8b929df 100644 --- a/libkineto/src/plugin/xpupti/XpuptiActivityHandlers.cpp +++ b/libkineto/src/plugin/xpupti/XpuptiActivityHandlers.cpp @@ -207,7 +207,9 @@ void XpuptiActivityProfilerSession::handleRuntimeKernelMemcpyMemsetActivities( std::is_same_v; constexpr bool handleKernelActivities = std::is_same_v; - constexpr bool handleMemcpyActivities = + constexpr bool handleMemcpyP2PActivities = + std::is_same_v; + constexpr bool handleMemcpyActivities = handleMemcpyP2PActivities || std::is_same_v; constexpr bool handleMemsetActivities = std::is_same_v; @@ -257,7 +259,17 @@ void XpuptiActivityProfilerSession::handleRuntimeKernelMemcpyMemsetActivities( trace_activity->device = activity->_process_id; trace_activity->resource = activity->_thread_id; } else { - trace_activity->device = getDeviceIdxFromUUID(activity->_device_uuid); + if constexpr (handleMemcpyP2PActivities) { + // The copy is submitted to, and runs on an engine of, the source device. + trace_activity->device = getDeviceIdxFromUUID(activity->_src_uuid); + trace_activity->addMetadata( + XpuFields::kFromDevice, static_cast(trace_activity->device)); + trace_activity->addMetadata( + XpuFields::kToDevice, + static_cast(getDeviceIdxFromUUID(activity->_dst_uuid))); + } else { + trace_activity->device = getDeviceIdxFromUUID(activity->_device_uuid); + } trace_activity->resource = activity->_sycl_queue_id; if constexpr (handleKernelActivities) { @@ -612,6 +624,12 @@ void XpuptiActivityProfilerSession::handlePtiActivity( reinterpret_cast(record), logger); break; + case PTI_VIEW_DEVICE_GPU_MEM_COPY_P2P: + handleRuntimeKernelMemcpyMemsetActivities( + ActivityType::GPU_MEMCPY, + reinterpret_cast(record), + logger); + break; case PTI_VIEW_DEVICE_GPU_MEM_FILL: handleRuntimeKernelMemcpyMemsetActivities( ActivityType::GPU_MEMSET, diff --git a/libkineto/src/plugin/xpupti/XpuptiActivityProfilerSession.h b/libkineto/src/plugin/xpupti/XpuptiActivityProfilerSession.h index 23973cdfa..e1687fc36 100644 --- a/libkineto/src/plugin/xpupti/XpuptiActivityProfilerSession.h +++ b/libkineto/src/plugin/xpupti/XpuptiActivityProfilerSession.h @@ -44,11 +44,13 @@ using DeviceUUIDsT = std::array; inline constexpr bool kReportsEngineIds = true; using pti_view_record_kernel_t = pti_view_record_kernel_v2; using pti_view_record_memcpy_t = pti_view_record_memory_copy_v2; +using pti_view_record_memcpy_p2p_t = pti_view_record_memory_copy_p2p_v2; using pti_view_record_memfill_t = pti_view_record_memory_fill_v2; #else inline constexpr bool kReportsEngineIds = false; using pti_view_record_kernel_t = pti_view_record_kernel; using pti_view_record_memcpy_t = pti_view_record_memory_copy; +using pti_view_record_memcpy_p2p_t = pti_view_record_memory_copy_p2p; using pti_view_record_memfill_t = pti_view_record_memory_fill; #endif diff --git a/libkineto/test/xpupti/XpuptiActivityHandlersTest.cpp b/libkineto/test/xpupti/XpuptiActivityHandlersTest.cpp index 8654a6fa5..4cd521325 100644 --- a/libkineto/test/xpupti/XpuptiActivityHandlersTest.cpp +++ b/libkineto/test/xpupti/XpuptiActivityHandlersTest.cpp @@ -37,6 +37,9 @@ static_assert( static_assert( offsetof(pti_view_record_memory_fill, _view_kind) == 0, "base record must be first member"); +static_assert( + offsetof(pti_view_record_memory_copy_p2p, _view_kind) == 0, + "base record must be first member"); // Mock XpuptiActivityApi that delivers hand-crafted PTI records // through the virtual processActivities without needing PTI runtime. @@ -652,6 +655,44 @@ TEST_F(XpuptiActivityHandlersTest, MemcpyAndMemsetRecordsAreFlowDestinations) { EXPECT_FALSE(memset_activity.flowStart()); } +TEST_F(XpuptiActivityHandlersTest, PeerToPeerMemcpyIsGpuMemcpy) { + constexpr uint32_t kCorrelationId = 53; + + KN::pti_view_record_memcpy_p2p_t p2p_record{}; + p2p_record._view_kind._view_kind = PTI_VIEW_DEVICE_GPU_MEM_COPY_P2P; + p2p_record._name = "zeCommandListAppendMemoryCopy(D2D - P2P)"; + p2p_record._start_timestamp = 200; + p2p_record._end_timestamp = 240; + p2p_record._thread_id = 7; + p2p_record._correlation_id = kCorrelationId; + p2p_record._sycl_queue_id = 3; + p2p_record._mem_op_id = 5; + p2p_record._bytes = 4096; + p2p_record._memcpy_type = PTI_VIEW_MEMCPY_TYPE_D2D; + p2p_record._mem_src = PTI_VIEW_MEMORY_TYPE_DEVICE; + p2p_record._mem_dst = PTI_VIEW_MEMORY_TYPE_DEVICE; + + mockApi_.records.push_back(&p2p_record._view_kind); + + auto traceBuffer = processAndGetTrace(); + ASSERT_EQ(traceBuffer->activities.size(), 1); + + auto& activity = *traceBuffer->activities[0]; + EXPECT_EQ(activity.type(), ActivityType::GPU_MEMCPY); + EXPECT_EQ(activity.name(), "Memcpy D2D (Device -> Device)"); + EXPECT_EQ(activity.resourceId(), 3); + EXPECT_EQ(activity.flowId(), kCorrelationId); + EXPECT_FALSE(activity.flowStart()); + EXPECT_EQ( + activity.getMetadataValue(XpuMetadataFields::kBytes), + std::optional{4096}); + EXPECT_EQ( + activity.getMetadataValue(XpuMetadataFields::kFromDevice), + std::optional{activity.deviceId()}); + EXPECT_TRUE( + activity.getMetadataValue(XpuMetadataFields::kToDevice).has_value()); +} + // --- Mixed dispatch test --- TEST_F(XpuptiActivityHandlersTest, MixedCommunicationAndSynchronization) { diff --git a/libkineto/test/xpupti/XpuptiProfilerTest.cpp b/libkineto/test/xpupti/XpuptiProfilerTest.cpp index 1868edabb..65ee2ea37 100644 --- a/libkineto/test/xpupti/XpuptiProfilerTest.cpp +++ b/libkineto/test/xpupti/XpuptiProfilerTest.cpp @@ -9,12 +9,17 @@ #include "XpuptiTestUtilities.h" #include "include/GenericTraceActivity.h" +#include "include/MetadataFieldCatalog.h" #include "include/libkineto.h" #include "src/plugin/xpupti/XpuptiActivityProfiler.h" +#include "test/xpupti/compute/XpuptiPeerCopyCompute.h" +#include #include #include +#include +#include #include #include @@ -313,3 +318,47 @@ TEST(XpuptiProfilerTest, GpuUserAnnotationLinkedActivityRetry) { runGpuUserAnnotationCase(LinkedActivityMode::MissFirstThenLinked); expectTwoUserAnnotations(*pBuffer); } + +TEST(XpuptiProfilerTest, PeerToPeerMemcpy) { + KN::XPUActivityProfiler profiler; + auto pSession = + profiler.configure({KN::ActivityType::GPU_MEMCPY}, KN::Config{}); + + const auto now = [] { + return std::chrono::duration_cast( + std::chrono::system_clock::now().time_since_epoch()) + .count(); + }; + const int64_t startTime = now(); + pSession->start(); + const bool copied = CopyPeerToPeerOnXpu(1024); + pSession->stop(); + const int64_t endTime = now(); + if (!copied) { + GTEST_SKIP() << "needs two GPUs with peer access"; + } + + TestActivityLogger logger; + pSession->processTrace( + logger, + [](int32_t) -> const KN::ITraceActivity* { return nullptr; }, + startTime, + endTime); + EXPECT_TRUE(pSession->errors().empty()) + << fmt::format("{}", fmt::join(pSession->errors(), ",")); + + auto pBuffer = pSession->getTraceBuffer(); + ASSERT_EQ(pBuffer->activities.size(), 1); + const auto& activity = *pBuffer->activities[0]; + EXPECT_EQ(activity.type(), KN::ActivityType::GPU_MEMCPY); + EXPECT_EQ( + activity.getMetadataValue(libkineto::XpuMetadataFields::kBytes), + std::optional{1024}); + const auto fromDevice = + activity.getMetadataValue(libkineto::XpuMetadataFields::kFromDevice); + const auto toDevice = + activity.getMetadataValue(libkineto::XpuMetadataFields::kToDevice); + EXPECT_EQ(fromDevice, std::optional{activity.deviceId()}); + ASSERT_TRUE(toDevice.has_value()); + EXPECT_NE(*fromDevice, *toDevice); +} diff --git a/libkineto/test/xpupti/XpuptiTestUtilities.cpp b/libkineto/test/xpupti/XpuptiTestUtilities.cpp index bc16f77c3..14874ccc0 100644 --- a/libkineto/test/xpupti/XpuptiTestUtilities.cpp +++ b/libkineto/test/xpupti/XpuptiTestUtilities.cpp @@ -202,34 +202,6 @@ static unsigned acceptOverheadActivities( return count; } -class TestActivityLogger : public KN::ActivityLogger { - void handleDeviceInfo( - [[maybe_unused]] const KN::DeviceInfo& info, - [[maybe_unused]] int64_t time) override {} - void handleResourceInfo( - [[maybe_unused]] const KN::ResourceInfo& info, - [[maybe_unused]] int64_t time) override {} - void handleOverheadInfo( - [[maybe_unused]] const KN::ActivityLogger::OverheadInfo& info, - [[maybe_unused]] int64_t time) override {} - void handleTraceSpan([[maybe_unused]] const KN::TraceSpan& span) override {} - void handleActivity( - [[maybe_unused]] const KN::ITraceActivity& activity) override {} - void handleGenericActivity( - [[maybe_unused]] const KN::GenericTraceActivity& activity) override {} - void handleTraceStart( - [[maybe_unused]] const std::unordered_map& - metadata, - [[maybe_unused]] const std::string& device_properties) override {} - void finalizeMemoryTrace( - [[maybe_unused]] const std::string&, - [[maybe_unused]] const KN::Config&) override {} - void finalizeTrace( - [[maybe_unused]] const KN::Config& config, - [[maybe_unused]] std::unique_ptr buffers, - [[maybe_unused]] int64_t endTime) override {} -}; - std::pair< std::unique_ptr, std::unique_ptr> diff --git a/libkineto/test/xpupti/XpuptiTestUtilities.h b/libkineto/test/xpupti/XpuptiTestUtilities.h index fd2240db8..dde7bb592 100644 --- a/libkineto/test/xpupti/XpuptiTestUtilities.h +++ b/libkineto/test/xpupti/XpuptiTestUtilities.h @@ -7,6 +7,7 @@ */ #include "include/output_base.h" +#include "src/ActivityBuffers.h" #include @@ -14,6 +15,36 @@ namespace KN = KINETO_NAMESPACE; bool IsEnvVerbose(); +// Discards everything: tests inspect the session's trace buffer instead. +class TestActivityLogger : public KN::ActivityLogger { + public: + void handleDeviceInfo( + [[maybe_unused]] const KN::DeviceInfo& info, + [[maybe_unused]] int64_t time) override {} + void handleResourceInfo( + [[maybe_unused]] const KN::ResourceInfo& info, + [[maybe_unused]] int64_t time) override {} + void handleOverheadInfo( + [[maybe_unused]] const KN::ActivityLogger::OverheadInfo& info, + [[maybe_unused]] int64_t time) override {} + void handleTraceSpan([[maybe_unused]] const KN::TraceSpan& span) override {} + void handleActivity( + [[maybe_unused]] const KN::ITraceActivity& activity) override {} + void handleGenericActivity( + [[maybe_unused]] const KN::GenericTraceActivity& activity) override {} + void handleTraceStart( + [[maybe_unused]] const std::unordered_map& + metadata, + [[maybe_unused]] const std::string& device_properties) override {} + void finalizeMemoryTrace( + [[maybe_unused]] const std::string&, + [[maybe_unused]] const KN::Config&) override {} + void finalizeTrace( + [[maybe_unused]] const KN::Config& config, + [[maybe_unused]] std::unique_ptr buffers, + [[maybe_unused]] int64_t endTime) override {} +}; + std::pair< std::unique_ptr, std::unique_ptr> diff --git a/libkineto/test/xpupti/compute/CMakeLists.txt b/libkineto/test/xpupti/compute/CMakeLists.txt index 567bc7f45..b47dc4849 100644 --- a/libkineto/test/xpupti/compute/CMakeLists.txt +++ b/libkineto/test/xpupti/compute/CMakeLists.txt @@ -20,7 +20,9 @@ endif() # it without seeing any SYCL. ComputeOnXpu is exported via the API macro in the # header (dllexport on Windows / default visibility on ELF) - no need for # WINDOWS_EXPORT_ALL_SYMBOLS or an exe<->dll symbol cycle. -add_library(xpupti_compute SHARED XpuptiScopeProfilerCompute.cpp) +add_library(xpupti_compute SHARED + XpuptiScopeProfilerCompute.cpp + XpuptiPeerCopyCompute.cpp) target_compile_options(xpupti_compute PRIVATE -fsycl) target_link_options(xpupti_compute PRIVATE -fsycl) diff --git a/libkineto/test/xpupti/compute/XpuptiPeerCopyCompute.cpp b/libkineto/test/xpupti/compute/XpuptiPeerCopyCompute.cpp new file mode 100644 index 000000000..4f644cdd4 --- /dev/null +++ b/libkineto/test/xpupti/compute/XpuptiPeerCopyCompute.cpp @@ -0,0 +1,32 @@ +/* + * Copyright (C) Intel Corporation + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + */ + +#define XPUPTI_COMPUTE_BUILDING +#include "XpuptiPeerCopyCompute.h" + +#include + +XPUPTI_COMPUTE_API bool CopyPeerToPeerOnXpu(std::size_t bytes) { + auto devices = sycl::device::get_devices(sycl::info::device_type::gpu); + if (devices.size() < 2) { + return false; + } + auto& src = devices[0]; + auto& dst = devices[1]; + if (!src.ext_oneapi_can_access_peer(dst)) { + return false; + } + + const sycl::context context({src, dst}); + sycl::queue queue(context, src, sycl::property::queue::in_order()); + auto* srcPtr = sycl::malloc_device(bytes, src, context); + auto* dstPtr = sycl::malloc_device(bytes, dst, context); + queue.memcpy(dstPtr, srcPtr, bytes).wait_and_throw(); + sycl::free(srcPtr, context); + sycl::free(dstPtr, context); + return true; +} diff --git a/libkineto/test/xpupti/compute/XpuptiPeerCopyCompute.h b/libkineto/test/xpupti/compute/XpuptiPeerCopyCompute.h new file mode 100644 index 000000000..5f2eaf7b1 --- /dev/null +++ b/libkineto/test/xpupti/compute/XpuptiPeerCopyCompute.h @@ -0,0 +1,16 @@ +/* + * Copyright (C) Intel Corporation + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + */ + +#pragma once + +#include "XpuptiScopeProfilerCompute.h" + +#include + +// Copies `bytes` from the first GPU to the second over a peer-to-peer link. +// Returns false, without copying, when there is no such pair of devices. +XPUPTI_COMPUTE_API bool CopyPeerToPeerOnXpu(std::size_t bytes);