From cc374451be8430cdf2ca8bca2a94ab4022f22ada Mon Sep 17 00:00:00 2001
From: Vladislav Kassym <vladislav.kassym@gmail.com>
Date: Fri, 1 Nov 2019 17:01:27 +0300
Subject: [PATCH] Fix __shfl_down cuda error

---
 README.md                    | 2 +-
 rnnt_include/detail/reduce.h | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/README.md b/README.md
index 2fef9fd..f358eb7 100644
--- a/README.md
+++ b/README.md
@@ -45,7 +45,7 @@ class RNNTLoss(Loss):
 From the repo test with:
 
 ```bash
-python test/test.py 10 300 100 50 --mx
+python test/test.py --mx
 ```
 
 ## Reference
diff --git a/rnnt_include/detail/reduce.h b/rnnt_include/detail/reduce.h
index 6405e21..eb4dbcb 100644
--- a/rnnt_include/detail/reduce.h
+++ b/rnnt_include/detail/reduce.h
@@ -30,7 +30,7 @@ struct CTAReduce {
 
         T shuff;
         for (int offset = warp_size / 2; offset > 0; offset /= 2) {
-            shuff = __shfl_down(x, offset);
+            shuff = __shfl_down_sync(0xFFFFFFFF, x, offset);
             if (tid + offset < count && tid < offset)
                 x = g(x, shuff);
         }