EngineX-Hygon/sglang
5
0
Fork 0
Code Issues Pull Requests Actions 7 Projects Releases Wiki Activity

Fix sgl-kernel cu118 compile issue (#2750)

Browse Source
This commit is contained in:
Ke Bao
2025-01-06 21:59:31 +08:00
committed by GitHub
parent 2f0d386496
commit 439f65809f
1 changed files with 4 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
sgl-kernel/src/sgl-kernel/csrc/trt_reduce_internal.cu
View File

@@ -302,8 +302,10 @@ static __global__ void __launch_bounds__(512, 1) twoShotAllReduceKernel(AllReduc
buffers[ii] = reinterpret_cast<T*>(params.peer_comm_buffer_ptrs[rank]);
}
#if (defined(__CUDACC_VER_MAJOR__) && (__CUDACC_VER_MAJOR__ >= 12))
#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
cudaGridDependencySynchronize();
#endif
#endif
block_barrier(params.peer_barrier_ptrs_in, params.barrier_flag, params.local_rank, RANKS_PER_NODE, tidx, bidx,
@@ -350,10 +352,11 @@ static __global__ void __launch_bounds__(512, 1) twoShotAllReduceKernel(AllReduc
*reinterpret_cast<int4*>(&local_output_buffer[offset_rank]) = *reinterpret_cast<int4*>(&buffers[ii][offset_rank]);
}
}
#if (defined(__CUDACC_VER_MAJOR__) && (__CUDACC_VER_MAJOR__ >= 12))
#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
cudaTriggerProgrammaticLaunchCompletion();
#endif
#endif
}
////////////////////////////////////////////////////////////////////////////////////////////////////