@@ -515,32 +515,32 @@ __device__ inline cuuint32_t __nvvm_get_smem_pointer(void *__ptr) {
515515#if !defined(__CUDA_ARCH__ ) || __CUDA_ARCH__ >= 800
516516__device__ inline unsigned __reduce_add_sync (unsigned __mask ,
517517 unsigned __value ) {
518- return __nvvm_redux_sync_add (__mask , __value );
518+ return __nvvm_redux_sync_add (__value , __mask );
519519}
520520__device__ inline unsigned __reduce_min_sync (unsigned __mask ,
521521 unsigned __value ) {
522- return __nvvm_redux_sync_umin (__mask , __value );
522+ return __nvvm_redux_sync_umin (__value , __mask );
523523}
524524__device__ inline unsigned __reduce_max_sync (unsigned __mask ,
525525 unsigned __value ) {
526- return __nvvm_redux_sync_umax (__mask , __value );
526+ return __nvvm_redux_sync_umax (__value , __mask );
527527}
528528__device__ inline int __reduce_min_sync (unsigned __mask , int __value ) {
529- return __nvvm_redux_sync_min (__mask , __value );
529+ return __nvvm_redux_sync_min (__value , __mask );
530530}
531531__device__ inline int __reduce_max_sync (unsigned __mask , int __value ) {
532- return __nvvm_redux_sync_max (__mask , __value );
532+ return __nvvm_redux_sync_max (__value , __mask );
533533}
534534__device__ inline unsigned __reduce_or_sync (unsigned __mask , unsigned __value ) {
535- return __nvvm_redux_sync_or (__mask , __value );
535+ return __nvvm_redux_sync_or (__value , __mask );
536536}
537537__device__ inline unsigned __reduce_and_sync (unsigned __mask ,
538538 unsigned __value ) {
539- return __nvvm_redux_sync_and (__mask , __value );
539+ return __nvvm_redux_sync_and (__value , __mask );
540540}
541541__device__ inline unsigned __reduce_xor_sync (unsigned __mask ,
542542 unsigned __value ) {
543- return __nvvm_redux_sync_xor (__mask , __value );
543+ return __nvvm_redux_sync_xor (__value , __mask );
544544}
545545
546546__device__ inline void __nv_memcpy_async_shared_global_4 (void * __dst ,
0 commit comments