From b0944dc72ad8bc1250dcc9a4e9c3eea6723c0206 Mon Sep 17 00:00:00 2001
From: Megvii Engine Team <megengine@megvii.com>
Date: Mon, 1 Mar 2021 23:42:12 +0800
Subject: [PATCH] fix(mgb): fix tx1 compile

GitOrigin-RevId: 0eeb62ff7389085d9a0d7f3488b01916423133d3
---
 src/opr/impl/standalone/nms_kern.cu | 6 +++++-
 1 file changed, 5 insertions(+), 1 deletion(-)

diff --git a/src/opr/impl/standalone/nms_kern.cu b/src/opr/impl/standalone/nms_kern.cu
index 1c5e70fea..fd543ff34 100644
--- a/src/opr/impl/standalone/nms_kern.cu
+++ b/src/opr/impl/standalone/nms_kern.cu
@@ -5,6 +5,10 @@
 
 namespace {
 
+#if __CUDACC_VER_MAJOR__ >= 9
+#define __shfl_down(x, y) __shfl_down_sync(0xffffffffu, x, y)
+#endif
+
 // each thread computs one bit
 const int THREADS_PER_BLOCK = 64;
 
@@ -95,7 +99,7 @@ __device__ __forceinline__ uint32_t warp_reduce_min_brdcst(uint32_t val) {
     static_assert(WARP_SIZE == 32, "warp size != 32");
 #pragma unroll
     for (uint32_t offset = WARP_SIZE / 2; offset; offset /= 2)
-        val = min(val, __shfl_down_sync(0xFFFFFFFF, val, offset));
+        val = min(val, __shfl_down(val, offset));
 
     if (!threadIdx.x)
         ans = val;
-- 
GitLab