From 4df7fe5d35fbf2a207e79f1cd074af2eaa46dd14 Mon Sep 17 00:00:00 2001
From: Jakub Szuppe <jakub@streamhpc.com>
Date: Thu, 28 Sep 2017 11:40:50 +0200
Subject: [PATCH] Update for CUDA 9.0

---
 library/src/rng/philox4x32_10.hpp | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/library/src/rng/philox4x32_10.hpp b/library/src/rng/philox4x32_10.hpp
index 1cbb82408..5a7a73769 100644
--- a/library/src/rng/philox4x32_10.hpp
+++ b/library/src/rng/philox4x32_10.hpp
@@ -127,7 +127,11 @@ namespace detail {
 
     inline __device__ unsigned int warp_reduce_min(unsigned int val, int size) {
       for (int offset = size/2; offset > 0; offset /= 2) {
+        #if defined(__HIP_PLATFORM_NVCC__) && __CUDACC_VER_MAJOR__ >= 9
+        unsigned int temp = __shfl_xor_sync(0xffffffff, (int)val, offset);
+        #else
         unsigned int temp = __shfl_xor((int)val, offset);
+        #endif
         val = (temp < val) ? temp : val;
       }
       return val;