From 9442285526c082ba85520e722389fb8b87092c35 Mon Sep 17 00:00:00 2001 From: Nathan Luehr Date: Thu, 11 Feb 2016 12:59:31 -0800 Subject: [PATCH] Fixed buffer overflow in ReduceOrCopy Bug caused AllGathers and ReduceScatters of less than 8 bytes to fail in certain cases. Change-Id: I33e1beb50805bfdb457ae16a90e3f91c1b283b9b Reviewed-on: http://git-master/r/1011505 Reviewed-by: Przemek Tredak Tested-by: Przemek Tredak --- src/common_kernel.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/common_kernel.h b/src/common_kernel.h index e30bf5c..e785b71 100644 --- a/src/common_kernel.h +++ b/src/common_kernel.h @@ -1,5 +1,5 @@ /************************************************************************* - * Copyright (c) 2015, NVIDIA CORPORATION. All rights reserved. + * Copyright (c) 2015-2016, NVIDIA CORPORATION. All rights reserved. * * Redistribution and use in source and binary forms, with or without * modification, are permitted provided that the following conditions @@ -274,7 +274,7 @@ __device__ inline void ReduceOrCopy(const int tid, const int UNROLL2 = (UNROLL >= 2) ? (UNROLL / 2) : 1; const bool NOUNROLL2 = ((UNROLL / 2) == 0); - int Npreamble = AlignUp(dest0, alignof(PackType)) - dest0; + int Npreamble = (N