#include <stdio.h>


__global__ void ArrSum(int* arr, int* result, int n){

	__shared__ int s[256]; //  this is shared memory for ALL threads in a block	
	int thread_id = threadIdx.x;
	int i = blockIdx.x * blockDim.x + threadIdx.x;

	s[thread_id] = (i < n) ? arr[i] : 0;

	__syncthreads();
	

	for(int stride = blockDim.x/2; stride > 0; stride = stride/2){
	
		if(thread_id < stride) s[thread_id] += s[thread_id + stride];
		__syncthreads(); // wait for all threads to reach this line

	}
	
	// move the s0 back to original result var
	if(thread_id == 0) atomicAdd(result, s[0]);

}





int test1(){
    int n = 1024;
    int *h_a = new int[n];
    int *d_a, *d_res;
    int res = 0;

    for(int i = 0; i < n; i++) h_a[i] = 1;

    cudaMalloc(&d_a, n*sizeof(int));      // & on both
    cudaMalloc(&d_res, sizeof(int));

    cudaMemcpy(d_a, h_a, n*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_res, &res, sizeof(int), cudaMemcpyHostToDevice);

    ArrSum<<<4, 256>>>(d_a, d_res, n);    // no & here

    cudaMemcpy(&res, d_res, sizeof(int), cudaMemcpyDeviceToHost);  // sizeof(int) not n*
    printf("Sum = %d\n", res);
    return 0;
}



int main(){

test1();
return 0;
}
