Удаление __syncthreads() в уменьшении уровня деформации CUDA

Следующий код суммирует каждые 32элементов массива с самым первым элементом каждого 32element group:

int i = threadIdx.x;
int warpid = i&31;
if(warpid < 16){
    s_buf[i] += s_buf[i+16];__syncthreads();
    s_buf[i] += s_buf[i+8];__syncthreads();
    s_buf[i] += s_buf[i+4];__syncthreads();
    s_buf[i] += s_buf[i+2];__syncthreads();
    s_buf[i] += s_buf[i+1];__syncthreads();
}

Я думал, что смогу исключить все __syncthreads()в коде, так как все операции выполняются в одном и том же варпе. Но если я их устраню, я получу мусорные результаты. Это не сильно повлияет на производительность, но я хочу знать, зачем мне здесь нужна __syncthreads().