2018-09-24 16:06:59 -07:00
|
|
|
/*************************************************************************
|
2019-03-14 19:39:20 -07:00
|
|
|
* Copyright (c) 2015-2019, NVIDIA CORPORATION. All rights reserved.
|
2018-09-24 16:06:59 -07:00
|
|
|
*
|
|
|
|
|
* See LICENSE.txt for license information
|
|
|
|
|
************************************************************************/
|
|
|
|
|
|
|
|
|
|
#ifndef NCCL_ENQUEUE_H_
|
|
|
|
|
#define NCCL_ENQUEUE_H_
|
|
|
|
|
|
|
|
|
|
#include "core.h"
|
|
|
|
|
#include "group.h"
|
|
|
|
|
|
2019-03-14 19:39:20 -07:00
|
|
|
// Channels / LL tuning
|
|
|
|
|
#define NCCL_LL_CHANNEL_THRESHOLD 8 // Per thread size before we start increasing nrings
|
|
|
|
|
#define NCCL_THREAD_THRESHOLD 64 // Per thread size before we switch to non-LL
|
|
|
|
|
#define NCCL_THREAD_THRESHOLD_PREVOLTA 32 // Per thread size before we switch to non-LL for pre-Volta archs
|
|
|
|
|
#define NCCL_LL_MIN_NTHREADS 64
|
|
|
|
|
|
2018-12-13 15:56:12 -08:00
|
|
|
ncclResult_t ncclEnqueueCheck(struct ncclInfo* info);
|
2018-09-24 16:06:59 -07:00
|
|
|
ncclResult_t ncclCpuBarrierIn(ncclComm_t comm, int* isLast);
|
|
|
|
|
ncclResult_t ncclCpuBarrierLast(ncclComm_t comm);
|
|
|
|
|
ncclResult_t ncclCpuBarrierOut(ncclComm_t comm);
|
|
|
|
|
ncclResult_t ncclBarrierEnqueue(ncclComm_t comm);
|
|
|
|
|
ncclResult_t ncclBarrierEnqueueWait(ncclComm_t comm);
|
|
|
|
|
ncclResult_t ncclEnqueueEvents(ncclComm_t comm);
|
|
|
|
|
|
|
|
|
|
#endif // End include guard
|
