testGPU_FusionReduction3() failing on some configurations

[tlemo]

On my Ubuntu 18.04 / Cuda 10.2 / GTX 1070, I see the following test failure:

[ RUN      ] JitTest.GPU_FusionReduction3_CUDA
unknown file: Failure
C++ exception with description "Expected t5.allclose(cg_output) to be true, but got false.  (Could this error message be improved?  If so, please report an enhancement request to PyTorch.)
Exception raised from testGPU_FusionReduction3 at ../test/cpp/jit/test_gpu.cpp:2524 (most recent call first):
frame #0: <unknown function> + 0xe99a7 (0x7f19417529a7 in /home/lemo/work/pytorch/build/lib/libc10.so)
frame #1: std::function<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > ()>::operator()() const + 0x4c (0x55cf8b604b78 in bin/test_jit)
frame #2: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0x40 (0x7f1941751c3c in /home/lemo/work/pytorch/build/lib/libc10.so)
frame #3: torch::jit::testGPU_FusionReduction3() + 0xb57 (0x55cf8b5e7062 in bin/test_jit)
frame #4: torch::jit::JitTest_GPU_FusionReduction3_CUDA_Test::TestBody() + 0x11 (0x55cf8b42d10f in bin/test_jit)
frame #5: void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) + 0x65 (0x55cf8b63f175 in bin/test_jit)
frame #6: void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) + 0x5a (0x55cf8b639945 in bin/test_jit)
frame #7: testing::Test::Run() + 0xd2 (0x55cf8b618b0c in bin/test_jit)
frame #8: testing::TestInfo::Run() + 0xf3 (0x55cf8b619463 in bin/test_jit)
frame #9: testing::TestCase::Run() + 0x104 (0x55cf8b619ae4 in bin/test_jit)
frame #10: testing::internal::UnitTestImpl::RunAllTests() + 0x2a6 (0x55cf8b6248ce in bin/test_jit)
frame #11: bool testing::internal::HandleSehExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) + 0x65 (0x55cf8b64023c in bin/test_jit)
frame #12: bool testing::internal::HandleExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) + 0x5a (0x55cf8b63a6a5 in bin/test_jit)
frame #13: testing::UnitTest::Run() + 0xc0 (0x55cf8b623352 in bin/test_jit)
frame #14: <unknown function> + 0x1be738 (0x55cf8b42c738 in bin/test_jit)
frame #15: main + 0x18a (0x55cf8b42c655 in bin/test_jit)
frame #16: __libc_start_main + 0xe7 (0x7f1940efbb97 in /lib/x86_64-linux-gnu/libc.so.6)
frame #17: _start + 0x2a (0x55cf8b42c34a in bin/test_jit)
" thrown in the test body.
[  FAILED  ] JitTest.GPU_FusionReduction3_CUDA (303 ms)

[tlemo]

For the record, the identified root cause (thanks Naoya) is:

After blockReduce , all threads write to T5[blockIdx.x * T5.stride[0]]. I believe in this case only threadIdx.x==0 should do it as it is the only thread that has the valid reduction result.

__global__ void kernel(Tensor<float, 2> T0, Tensor<float, 2> T1, Tensor<float, 1> T4, Tensor<float, 1> T5){
  float T6[1];
  if ( ( ( ( 0 * 16 ) + threadIdx.x ) < T0.size[1] ) ) { 
    T6[ 0 ]
       = float(0);
  }
  for(size_t i28 = 0; i28 < ( ceilDiv(T0.size[1], 16) ); ++i28 ) {
    float T2[1];
    if ( ( ( ( i28 * 16 ) + threadIdx.x ) < T0.size[1] ) ) { 
      T2[ 0 ]
         = T0[ ( blockIdx.x * T0.stride[0] ) + ( ( ( i28 * 16 ) + threadIdx.x ) * T0.stride[1] ) ]
         + T1[ ( blockIdx.x * T1.stride[0] ) + ( ( ( i28 * 16 ) + threadIdx.x ) * T1.stride[1] ) ];
    }
    if ( ( ( ( i28 * 16 ) + threadIdx.x ) < T0.size[1] ) ) { 
      T6[ 0 ]
         = T6[ 0 ]
         + T2[ 0 ];
    }
  }
  float T3[1];
  T3[ 0 ]
     = float(0);
  blockReduce< true, false, false > ( T3[ 0 ], T6[ 0 ], reduction_add_float);
  T5[ ( blockIdx.x * T5.stride[0] ) ] // <----------------------- 
     = T3[ 0 ]
     * T4[ ( blockIdx.x * T4.stride[0] ) ];
}

[csarofeen]

@tlemo I believe this should be fixed now in our devel branch, can you verify and close?

[tlemo]

The new code generation looks correct, closing this issue.