From 4ea8700ce3a54cf6c00219c06eb21bb21ec86a4e Mon Sep 17 00:00:00 2001 From: Ashwin Shenai Date: Sat, 6 Mar 2021 12:54:56 +0530 Subject: [PATCH] Patch LookupTable See torch/cunn/issues/492 for reference --- lib/THCUNN/LookupTable.cu | 5 +++-- rocks/cunn-scm-1.rockspec | 3 ++- 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/lib/THCUNN/LookupTable.cu b/lib/THCUNN/LookupTable.cu index 116639b4..1d457012 100644 --- a/lib/THCUNN/LookupTable.cu +++ b/lib/THCUNN/LookupTable.cu @@ -13,6 +13,7 @@ #include "THCHalf.h" #include "THCHalfAutoNumerics.cuh" #include "THCTensorSort.cuh" +#define FULL_MASK 0xffffffff const int WARP_SIZE = 32; @@ -29,7 +30,7 @@ __device__ __forceinline__ bool warpHasCollision(int val) #pragma unroll for (int i = 1; i <= 16; i++) { - dup |= (__shfl(val, (laneId + i) % 32) == val); + dup |= (__shfl_sync(FULL_MASK, val, (laneId + i) % 32) == val); } #else @@ -46,7 +47,7 @@ __device__ __forceinline__ bool warpHasCollision(int val) #endif - return __any(dup) != 0; + return __any_sync(FULL_MASK, dup) != 0; } template diff --git a/rocks/cunn-scm-1.rockspec b/rocks/cunn-scm-1.rockspec index 3efe9435..bb8501d2 100644 --- a/rocks/cunn-scm-1.rockspec +++ b/rocks/cunn-scm-1.rockspec @@ -2,7 +2,8 @@ package = "cunn" version = "scm-1" source = { - url = "git://github.com/torch/cunn.git", + url = "git://github.com/ashwin2802/cunn.git", + branch = "cuda11" } description = {