Following very simple code prints value of 1 to output (tensor(1., device=‘cuda:1’)); instead of 0, i’d expect to see
Using default stream (instead of custom) ‘fixes’ the issue, but defeats the purpose of original code - copy from one GPU to several others at the same time (async). Would be happy for any other solutions which archive this goal?
can only guess that something odd happens with memory when “x1 + 1” starts executing (see nvprof) and at the same time device-to-device copy starts happening
pytorch 1.0, CUDA 10
12.9429s  4.6388ms         (878907 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  ZN2at6native18elementwise_kernelILi512ELi1EZNS0_16gpu_unary_kernelIZNS0_17gpu_binary_kernelIZNS0_15add_kernel_implIfEEvRNS_14TensorIteratorEN3c106ScalarEEUlffE_EEvS6_RKT_EUlfE0_EEvS6_SC_EUliE_EEviT1 [886]
12.9468s  74.185ms                    -               -         -         -         -  3.3528GB  45.195GB/s      Device      Device  Tesla V100-SXM2         1        28  Tesla V100-SXM2         1  Tesla V100-SXM2         2  [CUDA memcpy PtoP]

import torch
n = int(3e4)
d0 = torch.device(‘cuda:0’)
x0 = torch.zeros(n, n, device=d0)
d1 = torch.device(‘cuda:1’)
x1 = torch.zeros(n, n, device=d1)
x1 + 1
x1 + 1
s = torch.cuda.Stream()
with torch.cuda.stream(s):
x2 = x0.to(d1)
torch.cuda.synchronize()  # sync on dev1 doesn’t help, using events don’t help either
print(x2[0, 0])
nvprof
Start  Duration            Grid Size      Block Size     Regs*    SSMem*    DSMem*      Size  Throughput  SrcMemType  DstMemType           Device   Context    Stream          Src Dev   Src Ctx          Dst Dev   Dst Ctx  Name
6.69668s  2.0160us                    -               -         -         -         -  3.3528GB   2e+06GB/s      Device           -  Tesla V100-SXM2         1         7                -         -                -         -  [CUDA memset]
13.0215s  1.9520us                    -               -         -         -         -  3.3528GB   2e+06GB/s      Device           -  Tesla V100-SXM2         2        18                -         -                -         -  [CUDA memset]
13.0261s  4.7312ms         (878907 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  ZN2at6native18elementwise_kernelILi512ELi1EZNS0_16gpu_unary_kernelIZNS0_17gpu_binary_kernelIZNS0_15add_kernel_implIfEEvRNS_14TensorIteratorEN3c106ScalarEEUlffE_EEvS6_RKT_EUlfE0_EEvS6_SC_EUliE_EEviT1 [886]
13.0292s  74.185ms                    -               -         -         -         -  3.3528GB  45.194GB/s      Device      Device  Tesla V100-SXM2         1        28  Tesla V100-SXM2         1  Tesla V100-SXM2         2  [CUDA memcpy PtoP]
13.0308s  5.0902ms         (878907 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  ZN2at6native18elementwise_kernelILi512ELi1EZNS0_16gpu_unary_kernelIZNS0_17gpu_binary_kernelIZNS0_15add_kernel_implIfEEvRNS_14TensorIteratorEN3c106ScalarEEUlffE_EEvS6_RKT_EUlfE0_EEvS6_SC_EUliE_EEviT1 [889]
13.0359s  5.0761ms         (878907 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  ZN2at6native18elementwise_kernelILi512ELi1EZNS0_16gpu_unary_kernelIZNS0_17gpu_binary_kernelIZNS0_15add_kernel_implIfEEvRNS_14TensorIteratorEN3c106ScalarEEUlffE_EEvS6_RKT_EUlfE0_EEvS6_SC_EUliE_EEviT1 [902]
13.0410s  5.0793ms         (878907 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  ZN2at6native18elementwise_kernelILi512ELi1EZNS0_16gpu_unary_kernelIZNS0_17gpu_binary_kernelIZNS0_15add_kernel_implIfEEvRNS_14TensorIteratorEN3c106ScalarEEUlffE_EEvS6_RKT_EUlfE0_EEvS6_SC_EUliE_EEviT1 [905]
13.1045s  2.6560us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply3<TensorEQOp<float, unsigned char>, unsigned char, float, float, unsigned int, int=1, int=1, int=1>(OffsetInfo<unsigned char, float, float>, OffsetInfo<TensorEQOp<float, unsigned char>, float, unsigned int>, OffsetInfo<unsigned char, float, int=1>, float, float) [1014]
13.1046s  1.5040us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply2<Tensor_abs_Float_Op, float, float, unsigned int, int=1, int=1>(OffsetInfo<float, unsigned int, int=1>, OffsetInfo<float, unsigned int, int=1>, unsigned int, Tensor_abs_Float_Op) [1028]
13.1047s  1.7280us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply2<TensorNEValueOp<float, unsigned char>, unsigned char, float, unsigned int, int=1, int=1>(OffsetInfo<unsigned char, unsigned char, float>, OffsetInfo<TensorNEValueOp<float, unsigned char>, unsigned char, unsigned int>, unsigned char, float) [1040]
13.1048s  1.6630us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply3<TensorBitAndOp, unsigned char, unsigned char, unsigned char, unsigned int, int=1, int=1, int=1>(OffsetInfo<TensorBitAndOp, unsigned char, unsigned int>, OffsetInfo<unsigned char, unsigned char, int=1>, OffsetInfo<unsigned char, unsigned char, int=1>, unsigned char, unsigned char) [1051]
13.1048s  1.1840us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply2<TensorNEValueOp<float, unsigned char>, unsigned char, float, unsigned int, int=1, int=1>(OffsetInfo<unsigned char, unsigned char, float>, OffsetInfo<TensorNEValueOp<float, unsigned char>, unsigned char, unsigned int>, unsigned char, float) [1062]
13.1048s  1.2160us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply3<TensorBitAndOp, unsigned char, unsigned char, unsigned char, unsigned int, int=1, int=1, int=1>(OffsetInfo<TensorBitAndOp, unsigned char, unsigned int>, OffsetInfo<unsigned char, unsigned char, int=1>, OffsetInfo<unsigned char, unsigned char, int=1>, unsigned char, unsigned char) [1073]
13.1049s  3.6480us              (1 1 1)      (1024 1 1)        30        0B  8.0000KB         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelReduceAll<unsigned char, unsigned int, long, thrust::identity, ReduceAdd, int=1>(TensorInfo<unsigned char, unsigned int>, unsigned int, long, long, thrust::identity, long*) [1086]
13.1049s  3.2000us                    -               -         -         -         -        8B  2.3842MB/s      Device    Pageable  Tesla V100-SXM2         2        18                -         -                -         -  [CUDA memcpy DtoH]
13.1051s  1.9520us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  ZN2at4cuda75_GLOBAL__N__51_tmpxft_000055a3_00000000_11_Copy_compute_75_cpp1_ii_dd3fb9a321kernelPointwiseApply2IZN75_GLOBAL__N__51_tmpxft_000055a3_00000000_11_Copy_compute_75_cpp1_ii_dd3fb9a36CopyOpIlhE5applyERNS_6TensorERKS6_EUlRlRKhE_lhjLi1ELi1ELi1EEEvNS0_6detail10TensorInfoIT0_T2_EENSF_IT1_SH_EESH_T [1101]
13.1051s  1.3440us              (1 1 1)       (128 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void thrust::cuda_cub::core::_kernel_agent<thrust::cuda_cub::__scan::InitAgent<thrust::cuda_cub::cub::ScanTileState<long, bool=1>, int>, thrust::cuda_cub::cub::ScanTileState<long, bool=1>, int>(bool=1, thrust::cuda_cub::cub::ScanTileState<long, bool=1>) [1121]
13.1051s  7.0080us              (1 1 1)        (96 1 1)        42        0B  2.0781KB         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void thrust::cuda_cub::core::_kernel_agent<thrust::cuda_cub::__scan::ScanAgent<thrust::device_ptr, thrust::device_ptr, thrust::plus, int, long, thrust::detail::integral_constant<bool, bool=0>>, thrust::device_ptr, thrust::device_ptr, thrust::plus, int, thrust::cuda_cub::cub::ScanTileState<long, bool=1>, thrust::cuda_cub::__scan::AddInitToExclusiveScan<long, thrust::plus>>(thrust::device_ptr, thrust::device_ptr, long, thrust::plus, int, long) [1126]
13.1052s  1.7600us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply3<TensorMaskedSelectOp<float, unsigned char, long>, unsigned char, long, float, unsigned int, int=1, int=1, int=1>(OffsetInfo<unsigned char, unsigned char, long>, OffsetInfo<long, unsigned char, float>, OffsetInfo<TensorMaskedSelectOp<float, unsigned char, long>, unsigned char, unsigned int>, unsigned char, float) [1132]
13.1052s  1.2480us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply2<Tensor_abs_Float_Op, float, float, unsigned int, int=1, int=1>(OffsetInfo<float, unsigned int, int=1>, OffsetInfo<float, unsigned int, int=1>, unsigned int, Tensor_abs_Float_Op) [1146]
13.1053s  1.8240us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  ZN2at4cuda75_GLOBAL__N__51_tmpxft_000055a3_00000000_11_Copy_compute_75_cpp1_ii_dd3fb9a321kernelPointwiseApply2IZN75_GLOBAL__N__51_tmpxft_000055a3_00000000_11_Copy_compute_75_cpp1_ii_dd3fb9a36CopyOpIdfE5applyERNS_6TensorERKS6_EUlRdRKfE_dfjLi1ELi1ELi1EEEvNS0_6detail10TensorInfoIT0_T2_EENSF_IT1_SH_EESH_T [1159]
13.1054s  5.6320us              (1 1 1)      (1024 1 1)        16        0B  8.0000KB         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelReduceAll<double, unsigned int, double, thrust::identity, ReduceMin, int=1>(TensorInfo<double, unsigned int>, unsigned int, double, double, thrust::identity, double*) [1172]
13.1054s  1.8880us                    -               -         -         -         -        8B  4.0410MB/s      Device    Pageable  Tesla V100-SXM2         2        18                -         -                -         -  [CUDA memcpy DtoH]
13.1055s  1.3440us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply1<TensorFillOp, double, unsigned int, int=1>(OffsetInfo<TensorFillOp, double, unsigned int>, double, double) [1184]
13.1055s  5.4080us              (1 1 1)      (1024 1 1)        16        0B  8.0000KB         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelReduceAll<double, unsigned int, double, thrust::identity, ReduceMax, int=1>(TensorInfo<double, unsigned int>, unsigned int, double, double, thrust::identity, double*) [1196]
13.1055s  1.5040us                    -               -         -         -         -        8B  5.0727MB/s      Device    Pageable  Tesla V100-SXM2         2        18                -         -                -         -  [CUDA memcpy DtoH]
13.1056s  1.1840us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply1<TensorFillOp, double, unsigned int, int=1>(OffsetInfo<TensorFillOp, double, unsigned int>, double, double) [1208]
13.1057s  1.5680us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply2<Tensor_ceil_Float_Op, float, float, unsigned int, int=1, int=1>(OffsetInfo<float, unsigned int, int=1>, OffsetInfo<float, unsigned int, int=1>, unsigned int, Tensor_ceil_Float_Op) [1226]
13.1057s  1.5680us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply3<TensorNEOp<float, unsigned char>, unsigned char, float, float, unsigned int, int=1, int=1, int=1>(OffsetInfo<unsigned char, float, float>, OffsetInfo<TensorNEOp<float, unsigned char>, float, unsigned int>, OffsetInfo<unsigned char, float, int=1>, float, float) [1238]
13.1058s  1.6000us                    -               -         -         -         -        1B  610.35KB/s      Device    Pageable  Tesla V100-SXM2         2        18                -         -                -         -  [CUDA memcpy DtoH]
13.1059s  2.3680us              (1 1 1)       (128 1 1)        26        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  ZN2at6native18elementwise_kernelILi128ELi4EZNS0_17gpu_binary_kernelIZNS0_15div_kernel_implIdEEvRNS_14TensorIteratorEEUlddE_EEvS5_RKT_EUliE0_EEviT1 [1259]
13.1060s  1.5360us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply2<TensorGTValueOp<double, unsigned char>, unsigned char, double, unsigned int, int=1, int=1>(OffsetInfo<unsigned char, unsigned char, double>, OffsetInfo<TensorGTValueOp<double, unsigned char>, unsigned char, unsigned int>, unsigned char, double) [1270]
13.1060s  1.4720us                    -               -         -         -         -        1B  663.42KB/s      Device    Pageable  Tesla V100-SXM2         2        18                -         -                -         -  [CUDA memcpy DtoH]
13.1060s  1.2480us              (1 1 1)       (512 1 1)        16        0B        0B         -           -           -           -  Tesla V100-SXM2         2        18                -         -                -         -  void kernelPointwiseApply2<TensorGTValueOp<double, unsigned char>, unsigned char, double, unsigned int, int=1, int=1>(OffsetInfo<unsigned char, unsigned char, double>, OffsetInfo<TensorGTValueOp<double, unsigned char>, unsigned char, unsigned int>, unsigned char, double) [1289]
13.1060s  1.4720us                    -               -         -         -         -        1B  663.42KB/s      Device    Pageable  Tesla V100-SXM2         2        18                -         -                -         -  [CUDA memcpy DtoH]
13.1061s  1.4720us                    -               -         -         -         -        4B  2.5915MB/s      Device    Pageable  Tesla V100-SXM2         2        18                -         -                -         -  [CUDA memcpy DtoH]
13.1061s  1.9520us                    -               -         -         -         -        4B  1.9543MB/s      Device    Pageable  Tesla V100-SXM2         2        18                -         -                -         -  [CUDA memcpy DtoH]