int array[8] =  [3, 1, 7, 0, 4, 1, 6, 3]

int sum = 0;
for (int i = 0; i < N; i++)sum += array[i]

int main(int argc,char** argv)
{initDevice(0);//initializationint size = 1 << 24;printf("    with array size %d  ", size);//execution configurationint blocksize = 1024;if (argc > 1){blocksize = atoi(argv[1]);   //从命令行输入设置block大小}dim3 block(blocksize, 1);dim3 grid((size - 1) / block.x + 1, 1);printf("grid %d block %d n", grid.x, block.x);//allocate host memorysize_t bytes = size * sizeof(int);int *idata_host = (int*)malloc(bytes);int *odata_host = (int*)malloc(grid.x * sizeof(int));int * tmp = (int*)malloc(bytes);//initialize the arrayinitialData_int(idata_host, size);memcpy(tmp, idata_host, bytes);double timeStart, timeElaps;int gpu_sum = 0;// device memoryint * idata_dev = NULL;int * odata_dev = NULL;CHECK(cudaMalloc((void**)&idata_dev, bytes));CHECK(cudaMalloc((void**)&odata_dev, grid.x * sizeof(int)));//cpu reduction 对照组int cpu_sum = 0;timeStart = cpuSecond();//cpu_sum = recursiveReduce(tmp, size);for (int i = 0; i < size; i++)cpu_sum += tmp[i];timeElaps = 1000*(cpuSecond() - timeStart);printf("cpu sum:%d n", cpu_sum);printf("cpu reduction elapsed %lf ms cpu_sum: %dn", timeElaps, cpu_sum);//kernel reduceNeighboredCHECK(cudaMemcpy(idata_dev, idata_host, bytes, cudaMemcpyHostToDevice));CHECK(cudaDeviceSynchronize());timeStart = cpuSecond();reduceNeighbored <<<grid, block >>>(idata_dev, odata_dev, size);cudaDeviceSynchronize();cudaMemcpy(odata_host, odata_dev, grid.x * sizeof(int), cudaMemcpyDeviceToHost);gpu_sum = 0;for (int i = 0; i < grid.x; i++)gpu_sum += odata_host[i];   timeElaps = 1000*(cpuSecond() - timeStart);printf("gpu sum:%d n", gpu_sum);printf("gpu reduceNeighbored elapsed %lf ms     <<<grid %d block %d>>>n",timeElaps, grid.x, block.x);// free host memoryfree(idata_host);free(odata_host);CHECK(cudaFree(idata_dev));CHECK(cudaFree(odata_dev));//reset devicecudaDeviceReset();//check the resultsif (gpu_sum == cpu_sum){printf("Test success!n");}return EXIT_SUCCESS;
}

__global__ void reduceNeighbored(int * g_idata,int * g_odata,unsigned int n)
{//set thread IDunsigned int tid = threadIdx.x;//boundary checkif (tid >= n) return;//convert global data pointer to the int *idata = g_idata + blockIdx.x*blockDim.x;//in-place reduction in global memoryfor (int stride = 1; stride < blockDim.x; stride *= 2){if ((tid % (2 * stride)) == 0){idata[tid] += idata[tid + stride];}//synchronize within block__syncthreads();}//write result for this block to global memif (tid == 0)g_odata[blockIdx.x] = idata[0];}