Skip to content

NCCL4Py v0.4.1 Release

Choose a tag to compare

@xiakun-lu xiakun-lu released this 11 Aug 23:00

Highlights

  • Added host APIs for NCCL teams, rank translation, and device resource configuration.
  • Expanded the experimental CuTe DSL device API with additional GIN operations, resource addressing, and direct support for host-created resources.
  • Added live NCCL parameter access and improved version reporting.

New Features

Teams and Device Resources

  • Added APIs for inspecting communicator teams, translating ranks, and configuring resources for device-side communication.

    APIs:

    • NCCLTeam
    • Communicator.team_world, Communicator.team_lsa, Communicator.team_rail
    • Communicator.team_rank_to_world(), Communicator.team_rank_to_lsa()
    • NCCLDevCommRequirements.teams, NCCLDevCommRequirements.resources
    • NCCLDevCommRequirements.gin_traffic_class
    • TeamRequirement
    • LsaBarrierRequirement, GinBarrierRequirement, LLA2ARequirement
    • MultimemHandle, LsaBarrierHandle, GinBarrierHandle, LLA2AHandle
    • DevCommResource.multimem_handle(), DevCommResource.resource_handles
    • RegisteredWindowHandle.get_multimem_device_pointer()

CuTe DSL Device API

  • Added device-side rank translation, resource addressing, and additional GIN operations.

    APIs:

    • DevComm.team_rank_to_world(), DevComm.team_rank_to_lsa()
    • DevComm.resource_buffer_local_pointer()
    • DevComm.resource_buffer_lsa_pointer()
    • DevComm.resource_buffer_peer_pointer()
    • DevComm.resource_buffer_multimem_pointer()
    • DevComm.resource_buffer_lsa_multimem_pointer()
    • Window.multimem_pointer(), Window.lsa_multimem_pointer()
    • Gin.put_value(), Gin.get(), Gin.signal(), Gin.flush()
    • Gin.read_signal(), Gin.reset_signal(), Gin.signal_shadow_pointer()
    • Gin.read_counter(), Gin.wait_counter(), Gin.reset_counter()
    • GIN_ALL_CONTEXTS
    • ThreadScope, GinResourceSharingMode
    • DevComm.gin(..., resource_sharing_mode=...)
    • opt_flags support for Gin.put(), Gin.put_value(), Gin.get(), and Gin.signal()
  • Host-created device communicators, registered windows, multimem handles, and LSA/GIN barrier handles can now be passed directly as arguments to @cute.jit functions.

Runtime Inspection

  • Added live NCCL parameter access and consolidated version reporting under nccl.core.

    APIs:

    • nccl.core.params
    • nccl.core.dump_params()
    • nccl.core.get_version()
    • nccl.core.show_versions()

Examples and Documentation

  • Added CuTe DSL examples covering teams, LSA all-reduce, multimem, GIN operations, barriers, and resource buffers.

Breaking Changes

  • nccl.ep is no longer included in nccl4py. The NCCL EP Python bindings are now distributed separately from the NVIDIA/nccl-extensions repository.
  • Version helpers are no longer exported from the nccl namespace. Import them from nccl.core.
  • Version reporting was consolidated:
    • nccl.core.get_version() now returns VersionInfo.
    • nccl.core.Version, get_lib_version(), and get_lib_path() were removed.
    • VersionInfo.nccl was renamed to VersionInfo.libnccl.
    • VersionInfo.nccl_ep was removed.
    • VersionInfo.nccl_bindings was added.

Fixes and Enhancements

  • Fixed the default release semantics of Gin.put().
  • Improved NCCL parameter error handling and string decoding.

Compatibility Notes

  • The top-level nccl Python package is now a PEP 420 namespace package.
  • The CuTe DSL device API remains experimental.
  • The selected libnccl_device.bc must exactly match the NCCL version used by the bindings; nccl4py does not perform this check automatically.
  • When one thread manages multiple local ranks, call Communicator.finalize() within nccl.core.group().

Known Issues

  • The following CuTe DSL APIs require NCCL 2.31:

    • Gin.get()
    • gin_session(), world_gin(), rail_gin(), hybrid_session(), world_hybrid()
  • CuTe DSL JIT compilation with CUDA 12 fails for kernels using the following APIs. CUDA 13 is not affected.

    Affected APIs:

    • Gin.put(), Gin.put_value(), Gin.get(), Gin.signal(), Gin.flush()
    • GinBarrierSession.sync(), BarrierSession.sync()