 		"mov.s32 %0, r0; }"
 		: "=r"(result) : "r"(x), "r"(offset), "r"(mask), "r"(x));
 #endif
-	return result;
-}
-
-MGPU_DEVICE int shfl_max(int x, int offset, int width = WARP_SIZE) {
+#endif
+	return result;
+}
+
+MGPU_DEVICE int shfl_max(int x, 
+		int offset, int width = WARP_SIZE, unsigned threadmask=MEMBERMASK) {
 	int result = 0;
-#if __CUDA_ARCH__ >= 300
+#ifdef USE_SHFL_SYNC
+	int mask = (WARP_SIZE - width)<< 8;
+	asm(
+		"{.reg .s32 r0;"
+		".reg .pred p;"
+		"shfl.sync.up.b32 r0|p, %1, %2, %3, %4;"
+		"@p max.s32 r0, r0, %5;"
+		"mov.s32 %0, r0; }"
+		: "=r"(result) : "r"(x), "r"(offset), "r"(mask), "r"(threadmask), "r"(x));
+#else
+#if ( __CUDA_ARCH__ >= 300)
 	int mask = (WARP_SIZE - width)<< 8;
 	asm(
 		"{.reg .s32 r0;"


////////////////////////////////////////////////////////////////////////////////
// shfl_up
#pragma push_macro("__shfl_up")
#undef __shfl_up
__device__ __forceinline__ float shfl_up(float var, 
	unsigned int delta, int width = 32) {

#if __CUDA_ARCH__ >= 300
	var = __shfl_up(var, delta, width);
#endif	
	return var;
}

__device__ __forceinline__ double shfl_up(double var, 
	unsigned int delta, int width = 32) {

#if __CUDA_ARCH__ >= 300
	int2 p = mgpu::double_as_int2(var);
	p.x = __shfl_up(p.x, delta, width);
	p.y = __shfl_up(p.y, delta, width);
	var = mgpu::int2_as_double(p);
#endif
	
	return var;
}
#pragma pop_macro("__shfl_up")
////////////////////////////////////////////////////////////////////////////////
// shfl_add

MGPU_DEVICE int shfl_add(int x, int offset, int width = WARP_SIZE) {
	int result = 0;
#if __CUDA_ARCH__ >= 300
	int mask = (WARP_SIZE - width)<< 8;
	asm(
		"{.reg .s32 r0;"

		"@p max.s32 r0, r0, %4;"
		"mov.s32 %0, r0; }"
		: "=r"(result) : "r"(x), "r"(offset), "r"(mask), "r"(x));
#endif
#endif
	return result;
}
