[TEXPR][PASS] Fix thread all reduce to avoid write after read hazzard

[tqchen]

The bug was rare but can affect some of the code. Insert explicit sync before write to shared memory in thread_allreduce

__shared__ sbuf[65]
loop:
    // need add sync here, before the read may not yet finished.
    __syncthreads();
    sbuf[threadIdx.x] = x
    // all reduce code ...
    // read
    result = sbuf[0]

[tqchen]

Fixes #2767