NCCL 2.27.6-1
Improve support for DirectNIC (CX8)
* Add support for XDR speed detection.
* When DirectNIC is enabled, report only the RDMA interfaces.
Extend the P2C (PXN over C2C) support to send/receive operations.
Support compilation with GCC 14 (Issues #1743, #1751).
Fix the unloading of network plugins that also provide tuner capability.
Fix the change of the current device across the calls to ncclCommDestroy()
and ncclCommAbort().
A note for users on MNNVL systems: please ensure an adequate stack size for
NCCL threads. While the default Linux stack size limit of 8192 KB is known
to be sufficient, we've seen crashes if the limit is changed to
"unlimited", as it causes the glibc library to unexpectedly *decrease* the
stack size of NCCL's background threads to just 2048 KB. Use "ulimit -s"
in bash to print the current limit; if needed, reset it to 8192 KB using
"ulimit -s 8192" (one also needs to ensure that the new setting is
propagated to other nodes when launching a multi-node NCCL job).
[ROCm/rccl commit: 7c12c627c6]
Этот коммит содержится в:
Kamil Iskra
@@ -9,6 +9,7 @@
|
||||
#include <stdint.h>
|
||||
#include <sys/types.h>
|
||||
#include <unistd.h>
|
||||
#include <string.h>
|
||||
|
||||
#if __GNUC__ >= 3
|
||||
# define __attribute_const __attribute__((const))
|
||||
@@ -39,7 +40,7 @@ union ibv_gid {
|
||||
#define vext_field_avail(type, fld, sz) (offsetof(type, fld) < (sz))
|
||||
|
||||
/*XXX:__VERBS_ABI_IS_EXTENDED produces warning "integer operation result is out of range" with g++ 4.8.2*/
|
||||
//static void *__VERBS_ABI_IS_EXTENDED = ((uint8_t *)NULL) - 1;
|
||||
static void *__VERBS_ABI_IS_EXTENDED = ((uint8_t *)NULL) - 1;
|
||||
|
||||
enum ibv_node_type {
|
||||
IBV_NODE_UNKNOWN = -1,
|
||||
@@ -208,7 +209,9 @@ struct ibv_port_attr {
|
||||
uint8_t active_speed;
|
||||
uint8_t phys_state;
|
||||
uint8_t link_layer;
|
||||
uint8_t reserved;
|
||||
uint8_t flags;
|
||||
uint16_t port_cap_flags2;
|
||||
uint32_t active_speed_ex;
|
||||
};
|
||||
|
||||
enum ibv_event_type {
|
||||
@@ -993,37 +996,50 @@ enum verbs_context_mask {
|
||||
|
||||
struct verbs_context {
|
||||
/* "grows up" - new fields go here */
|
||||
int (*_reserved_2) (void);
|
||||
int (*destroy_flow) (struct ibv_flow *flow);
|
||||
int (*_reserved_1) (void);
|
||||
struct ibv_flow * (*create_flow) (struct ibv_qp *qp,
|
||||
struct ibv_flow_attr *flow_attr);
|
||||
int (*query_port)(struct ibv_context *context, uint8_t port_num,
|
||||
struct ibv_port_attr *port_attr,
|
||||
size_t port_attr_len);
|
||||
int (*_reserved[25]) (void);
|
||||
struct verbs_ex_private *priv;
|
||||
int (*query_device_ex)(struct ibv_context *context,
|
||||
const struct ibv_query_device_ex_input *input,
|
||||
struct ibv_device_attr_ex *attr,
|
||||
size_t attr_size);
|
||||
int (*ibv_destroy_flow) (struct ibv_flow *flow);
|
||||
void (*ABI_placeholder2) (void); /* DO NOT COPY THIS GARBAGE */
|
||||
struct ibv_flow * (*ibv_create_flow) (struct ibv_qp *qp,
|
||||
struct ibv_flow_attr *flow_attr);
|
||||
void (*ABI_placeholder1) (void); /* DO NOT COPY THIS GARBAGE */
|
||||
struct ibv_qp * (*open_qp)(struct ibv_context *context,
|
||||
struct ibv_qp_open_attr *attr);
|
||||
struct ibv_qp * (*create_qp_ex)(struct ibv_context *context,
|
||||
struct ibv_qp_init_attr_ex *qp_init_attr_ex);
|
||||
int (*get_srq_num)(struct ibv_srq *srq, uint32_t *srq_num);
|
||||
struct ibv_srq * (*create_srq_ex)(struct ibv_context *context,
|
||||
struct ibv_srq_init_attr_ex *srq_init_attr_ex);
|
||||
struct ibv_xrcd * (*open_xrcd)(struct ibv_context *context,
|
||||
struct ibv_xrcd_init_attr *xrcd_init_attr);
|
||||
int (*close_xrcd)(struct ibv_xrcd *xrcd);
|
||||
uint64_t has_comp_mask;
|
||||
size_t sz; /* Must be immediately before struct ibv_context */
|
||||
struct ibv_context context;/* Must be last field in the struct */
|
||||
struct ibv_srq * (*create_srq_ex)(struct ibv_context *context,
|
||||
struct ibv_srq_init_attr_ex *srq_init_attr_ex);
|
||||
struct ibv_xrcd * (*open_xrcd)(struct ibv_context *context,
|
||||
struct ibv_xrcd_init_attr *xrcd_init_attr);
|
||||
int (*close_xrcd)(struct ibv_xrcd *xrcd);
|
||||
uint64_t _ABI_placeholder3;
|
||||
size_t sz; /* Must be immediately before struct ibv_context */
|
||||
struct ibv_context context; /* Must be last field in the struct */
|
||||
};
|
||||
|
||||
/*XXX:__VERBS_ABI_IS_EXTENDED produces warning "integer operation result is out of range" with g++ 4.8.2*/
|
||||
/*static inline struct verbs_context *verbs_get_ctx(struct ibv_context *ctx)
|
||||
static inline struct verbs_context *verbs_get_ctx(struct ibv_context *ctx)
|
||||
{
|
||||
return (!ctx || (ctx->abi_compat != __VERBS_ABI_IS_EXTENDED)) ?
|
||||
NULL : container_of(ctx, struct verbs_context, context);
|
||||
if (ctx->abi_compat != __VERBS_ABI_IS_EXTENDED)
|
||||
return NULL;
|
||||
|
||||
/* open code container_of to not pollute the global namespace */
|
||||
return (struct verbs_context *)(((uintptr_t)ctx) -
|
||||
offsetof(struct verbs_context,
|
||||
context));
|
||||
}
|
||||
|
||||
#define verbs_get_ctx_op(ctx, op) ({ \
|
||||
struct verbs_context *_vctx = verbs_get_ctx(ctx); \
|
||||
(!_vctx || (_vctx->sz < sizeof(*_vctx) - offsetof(struct verbs_context, op)) || \
|
||||
!_vctx->op) ? NULL : _vctx; })*/
|
||||
struct verbs_context *__vctx = verbs_get_ctx(ctx); \
|
||||
(!__vctx || (__vctx->sz < sizeof(*__vctx) - offsetof(struct verbs_context, op)) || \
|
||||
!__vctx->op) ? NULL : __vctx; })
|
||||
|
||||
#define verbs_set_ctx_op(_vctx, op, ptr) ({ \
|
||||
struct verbs_context *vctx = _vctx; \
|
||||
@@ -1055,4 +1071,20 @@ struct ibv_ece {
|
||||
uint32_t comp_mask;
|
||||
};
|
||||
|
||||
/**
|
||||
* ibv_query_port_ex - Get (extended) port properties
|
||||
*/
|
||||
static inline int ibv_query_port_ex(struct ibv_context *context,
|
||||
uint8_t port_num,
|
||||
struct ibv_port_attr *port_attr)
|
||||
{
|
||||
struct verbs_context *vctx = verbs_get_ctx_op(context, query_port);
|
||||
|
||||
if (vctx) {
|
||||
return vctx->query_port(context, port_num, port_attr, sizeof(*port_attr));
|
||||
}
|
||||
|
||||
return -1;
|
||||
}
|
||||
|
||||
#endif // NCCL_IBV_CORE_H_
|
||||
|
||||
@@ -9,10 +9,16 @@
|
||||
|
||||
#include "nccl.h"
|
||||
|
||||
enum ncclPluginType {
|
||||
ncclPluginTypeNet,
|
||||
ncclPluginTypeTuner,
|
||||
ncclPluginTypeProfiler,
|
||||
};
|
||||
|
||||
void* ncclOpenNetPluginLib(const char* name);
|
||||
void* ncclOpenTunerPluginLib(const char* name);
|
||||
void* ncclOpenProfilerPluginLib(const char* name);
|
||||
void* ncclGetNetPluginLib(void);
|
||||
ncclResult_t ncclClosePluginLib(void* handle);
|
||||
void* ncclGetNetPluginLib(enum ncclPluginType type);
|
||||
ncclResult_t ncclClosePluginLib(void* handle, enum ncclPluginType type);
|
||||
|
||||
#endif
|
||||
|
||||
Ссылка в новой задаче
Block a user