# explain me the vector reduction 

**URL:** <https://forums.developer.nvidia.com/t/explain-me-the-vector-reduction/27769>\
**Category:** CUDA Programming and Performance\
**Created:** [December 26, 2012, 11:03am UTC](https://forums.developer.nvidia.com/t/explain-me-the-vector-reduction/27769 "2012-12-26T11:03:36Z")\
**Posts on this page:** 2\
**Page:** 1

<div class="post-metadata">

**Author:** ![light86](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@light86](https://forums.developer.nvidia.com/u/light86)\
**Post date:** [December 26, 2012, 11:03am UTC](https://forums.developer.nvidia.com/t/explain-me-the-vector-reduction/27769/1 "2012-12-26T11:03:36Z")

</div>

Hi all I’m new in cuda word  
in this code

```auto
global__ void dot( float *a, float *b, float *c ) {
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int cacheIndex = threadIdx.x;
    float temp = 0;
    while (tid < N) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }
    // set the cache values
    cache[cacheIndex] = temp;
    // synchronize threads in this block
    __syncthreads();
    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0) {
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }
    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}

```

I don’t understand the aim of int i = blockDim.x/2;  
(the reduction is in the same block) PLease explain me

[/code]  
Thank you

---

<div class="post-metadata">

**Author:** ![pasoleatis](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@pasoleatis](https://forums.developer.nvidia.com/u/pasoleatis)\
**Post date:** [December 26, 2012, 9:04pm UTC](https://forums.developer.nvidia.com/t/explain-me-the-vector-reduction/27769/2 "2012-12-26T21:04:12Z")

</div>

Let assume we have tpb=blockDim.x threads in our block. The reduction algorithm reduces at each iteration the amount of data to half. The first iteration transforms the data from tpb to tpb/2 this is why you start with that. In the next iteration you reduce the data to tpb/4 and so on until you get to 1 element which will contain the sum of everything.  
You start with tpb/2 because this way you avoid the bank conflicts.  
Check this image:

![](https://global.discourse-cdn.com/nvidia/original/3X/5/4/548bcffc906c5fc881357f22aee76deaee437cfb.png)