Support CudaIpc connection within a single process

[chhwang]

• Allow CudaIpc connection between GPUs in a single process
• Added an example of connection in a single process
• Minor interface updates

[Copilot]

Pull Request Overview

Extends CudaIpc transport to support connections between GPUs within a single process, whereas previously it only worked across different processes. This enables GPU-to-GPU communication within the same application.

Key Changes

• Added process ID hash tracking to endpoints for intra-process GPU connections
• Enabled peer access for same-process GPU connections
• Updated constants and APIs for better maintainability

Reviewed Changes

Copilot reviewed 13 out of 13 changed files in this pull request and generated 2 comments.

Show a summary per file

File	Description
src/registered_memory.cc	Removed namespace prefix from SysError exception
src/include/endpoint.hpp	Added pidHash field to endpoint implementation
src/endpoint.cc	Added pidHash serialization and getter methods
src/context.cc	Added peer access enablement for same-process GPU connections
src/connection.cc	Moved connection constructor to implementation file
python/mscclpp/core_py.cpp	Removed deprecated Nvls transport enum value
include/mscclpp/port_channel_device.hpp	Replaced macros with constexpr constants
include/mscclpp/gpu.hpp	Added HIP compatibility for cudaDeviceEnablePeerAccess
include/mscclpp/core.hpp	Removed Nvls transport and updated constants
examples/tutorials/01-basic-concepts/	Added complete GPU ping-pong example with build files
apps/nccl/src/nccl.cu	Updated to use new constant naming

Comments suppressed due to low confidence (1)

examples/tutorials/01-basic-concepts/gpu_ping_pong.cu:10

• [nitpick] Function names 'gpuKernel0' and 'gpuKernel1' are not descriptive. Consider more meaningful names like 'pingKernel' and 'pongKernel' or 'initiatorKernel' and 'responderKernel'.

__global__ void gpuKernel0(mscclpp::BaseMemoryChannelDeviceHandle *devHandle, int iter) {