AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

Fix stack size

Bläddra i källkod
This commit is contained in:
BertanDogancay
2024-01-31 17:09:07 -08:00
förälder 95f87232c4
incheckning da85abab54
2 ändrade filer med 8 tillägg och 10 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
src/device/common.cu
+7 -3
Visa fil
@@ -17,13 +17,17 @@ struct RunWorkNop {
__device__ void run(ncclWork *w) {}
};
__global__ void ncclDevKernel_Generic(struct ncclDevComm* comm, uint64_t channelMask, struct ncclWork* workHead) {
__launch_bounds__(NCCL_MAX_NTHREADS, 1) __global__ void ncclDevKernel_Generic(struct ncclDevComm* comm, uint64_t channelMask, struct ncclWork* workHead) {
ncclKernelMain<-1, RunWorkNop, false>(comm, channelMask, workHead);
}
#ifdef ENABLE_COLLTRACE
__global__ void ncclDevKernelDebug_Generic(struct ncclDevComm* comm, uint64_t channelMask, struct ncclWork* workHead) {
__launch_bounds__(NCCL_MAX_NTHREADS, 1) __global__ void ncclDevKernelDebug_Generic(struct ncclDevComm* comm, uint64_t channelMask, struct ncclWork* workHead) {
ncclKernelMain<-1, RunWorkNop, true>(comm, channelMask, workHead);
}
#endif
__device__ void ncclDevFunc_Nop() {}
#ifdef USE_INDIRECT_FUNCTION_CALL
__device__ void ncclDevFunc_Nop();
#else
__device__ __attribute__((noinline)) void ncclDevFunc_Nop();
#endif
src/device/common.h
+1 -7
Visa fil
@@ -225,7 +225,7 @@ static __forceinline__ __device__ void ncclRedopPtrDeref(struct ncclWorkElem* we
}
template<int SpecializedFnId, typename SpecializedRunWork, bool COLLTRACE>
__device__ void ncclKernelMain(struct ncclDevComm* comm, uint64_t channelMask, struct ncclWork* workHead) {
__forceinline__ __device__ void ncclKernelMain(struct ncclDevComm* comm, uint64_t channelMask, struct ncclWork* workHead) {
const int tid = threadIdx.x;
int x = tid;
switch (tid/WARP_SIZE) {
@@ -365,12 +365,6 @@ __global__ void ncclDevKernel_Generic(struct ncclDevComm* comm, uint64_t channel
__global__ void ncclDevKernelDebug_Generic(struct ncclDevComm* comm, uint64_t channelMask, struct ncclWork* workHead);
#endif
#ifdef USE_INDIRECT_FUNCTION_CALL
__device__ void ncclDevFunc_Nop();
#else
__device__ __attribute__((noinline)) void ncclDevFunc_Nop();
#endif
#ifdef ENABLE_COLLTRACE
#define DEFINE_ncclDevKernel(suffix, coll, redop, ty, algo, proto, specializedFnId) \
__global__ void ncclDevKernel_##suffix(struct ncclDevComm* comm, uint64_t channelMask, struct ncclWork* workHead) { \