# 32 thread block doesn't need \_syncthreads()?

**URL:** <https://forums.developer.nvidia.com/t/32-thread-block-doesnt-need-syncthreads/1490>\
**Category:** CUDA Programming and Performance\
**Created:** [September 17, 2007, 3:52pm UTC](https://forums.developer.nvidia.com/t/32-thread-block-doesnt-need-syncthreads/1490 "2007-09-17T15:52:36Z")\
**Posts on this page:** 1\
**Showing post:** 18

<div class="post-metadata">

**Author:** ![czxu](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@czxu](https://forums.developer.nvidia.com/u/czxu)\
**Post date:** [January 21, 2024, 5:26pm UTC](https://forums.developer.nvidia.com/t/32-thread-block-doesnt-need-syncthreads/1490/18 "2024-01-21T17:26:22Z")

</div>

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

```auto
__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:

```auto
    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”?

---

_[View the full topic](https://forums.developer.nvidia.com/t/32-thread-block-doesnt-need-syncthreads/1490)._