32 thread block doesn't need _syncthreads()?

Accelerated Computing CUDA CUDA Programming and Performance
18

Interesting topic, I tried on Ampere GPU with a reduced sum application, using 32 threads per block, and it seems that syncthreads() are needed

__global__ void global_sum_reduce_kernel(float * arr, float * sum_result) {
    // 32 threads per block, 64 data per block
    unsigned int arr_size = 64;
    const unsigned int data_id = blockIdx.x * arr_size + threadIdx.x;

    __shared__ __align__(2 * 1024) char smem[32 * 4];
    float * arr_s = reinterpret_cast<float *>(smem);
    // load global
    arr_s[threadIdx.x] = arr[data_id] + arr[data_id + arr_size / 2];
    arr_size /= 2;
    __syncthreads();
    while (arr_size > 1) {
        if (threadIdx.x < arr_size / 2) arr_s[threadIdx.x] += arr_s[threadIdx.x + arr_size / 2];
        arr_size /= 2;
        __syncthreads();
    }
    // the result is in arr_s[0];
    if (threadIdx.x == 0) atomicAdd(sum_result, arr_s[0]);
    return;
}

and in Host side:

    const unsigned int thread_num = 32;
    const unsigned int block_num = arr_size/(thread_num*2);
    global_sum_reduce_kernel<<<block_num, thread_num>>>(arr_device, sum_result_device);

The two syncthreads() are both needed or the sum result will be incorrect.

Is this an example for what Mark said about “shared memory cached on register”?