 #include <iostream>
 #include <numeric>
 
-#define LEN 65536
-#define SIZE LEN*sizeof(float)
+#define LEN 1024*16
+#define SIZE LEN*sizeof(int)
 
 __global__ void reduce0(int *g_idata, int *g_odata)
 {

    if(tid == 0) g_odata[blockIdx.x] = sdata[0];
}

int main()
{
    int *A = new int[LEN];
    int *B = new int[LEN];

    cudaMemcpy(Bd, B, SIZE, cudaMemcpyHostToDevice);
    dim3 dimGrid(LEN/1024,1,1);
    reduce0<<<dimGrid, dim3(1024,1,1)>>>(Ad, Bd);
//    reduce0<<<dim3(1,1,1), dim3(dimGrid.x,1,1)>>>(Bd, Bd);
    cudaDeviceSynchronize();
    cudaMemcpy(B, Bd, SIZE, cudaMemcpyDeviceToHost);    
    std::cout<<B[0]<<std::endl;
    delete A;
    delete B;
    cudaFree(Ad);
    cudaFree(Bd);
}


void doReduce1()
{
    int *A = new int[LEN];
    int *B = new int[LEN];

    int *Ad, *Bd;

    for(int i=0;i<LEN;i++) {
        A[i] = 1;
        B[i] = 0;
    }

    cudaMalloc(&Ad, SIZE);
    cudaMalloc(&Bd, SIZE);

    cudaMemcpy(Ad, A, SIZE, cudaMemcpyHostToDevice);
    cudaMemcpy(Bd, B, SIZE, cudaMemcpyHostToDevice);
    dim3 dimGrid(LEN/1024,1,1);
    reduce1<<<dimGrid, dim3(1024,1,1)>>>(Ad, Bd);
//    reduce1<<<dim3(1,1,1), dim3(dimGrid.x,1,1)>>>(Bd, Bd);
    cudaDeviceSynchronize();
    cudaMemcpy(B, Bd, SIZE, cudaMemcpyDeviceToHost);    
    std::cout<<B[0]<<std::endl;
    delete A;
    delete B;
    cudaFree(Ad);
    cudaFree(Bd);
}


void doReduce2()
{
    int *A = new int[LEN];
    int *B = new int[LEN];

    int *Ad, *Bd;

    for(int i=0;i<LEN;i++) {
        A[i] = 1;
        B[i] = 0;
    }

    cudaMalloc(&Ad, SIZE);
    cudaMalloc(&Bd, SIZE);

    cudaMemcpy(Ad, A, SIZE, cudaMemcpyHostToDevice);
    cudaMemcpy(Bd, B, SIZE, cudaMemcpyHostToDevice);
    dim3 dimGrid(LEN/1024,1,1);
    reduce2<<<dimGrid, dim3(1024,1,1)>>>(Ad, Bd);
//    reduce2<<<dim3(1,1,1), dim3(dimGrid.x,1,1)>>>(Bd, Bd);
    cudaDeviceSynchronize();
    cudaMemcpy(B, Bd, SIZE, cudaMemcpyDeviceToHost);    
    std::cout<<B[0]<<std::endl;
    delete A;
    delete B;
    cudaFree(Ad);
    cudaFree(Bd);
}

void doReduce3()
{
    int *A = new int[LEN];
    int *B = new int[LEN];

    int *Ad, *Bd;

    for(int i=0;i<LEN;i++) {
        A[i] = 1;
        B[i] = 0;
    }

    cudaMalloc(&Ad, SIZE);
    cudaMalloc(&Bd, SIZE);

    cudaMemcpy(Ad, A, SIZE, cudaMemcpyHostToDevice);
    cudaMemcpy(Bd, B, SIZE, cudaMemcpyHostToDevice);
    dim3 dimGrid(LEN/1024,1,1);
    reduce3<<<dimGrid, dim3(1024,1,1)>>>(Ad, Bd);
//    reduce3<<<dim3(1,1,1), dim3(dimGrid.x,1,1)>>>(Bd, Bd);
    cudaDeviceSynchronize();
    cudaMemcpy(B, Bd, SIZE, cudaMemcpyDeviceToHost);    
    std::cout<<B[0]<<std::endl;
    delete A;
    delete B;
    cudaFree(Ad);
    cudaFree(Bd);
}


void doReduce6()
{
    int *A = new int[LEN];
    int *B = new int[LEN];

    int *Ad, *Bd;

    for(int i=0;i<LEN;i++) {
        A[i] = 1;
        B[i] = 0;
    }

    cudaMalloc(&Ad, SIZE);
    cudaMalloc(&Bd, SIZE);

    cudaMemcpy(Ad, A, SIZE, cudaMemcpyHostToDevice);
    cudaMemcpy(Bd, B, SIZE, cudaMemcpyHostToDevice);
    dim3 dimGrid(LEN/1024,1,1);
    dim3 dimBlock(1024,1,1);
    
    reduce6<int, 1024, true> <<<dimGrid, dim3(1024,1,1), dimBlock.x*2*sizeof(int)>>>(Ad, Bd, LEN);
    cudaDeviceSynchronize();
    cudaMemcpy(B, Bd, SIZE, cudaMemcpyDeviceToHost);    
    std::cout<<B[0]<<std::endl;
    delete A;
    delete B;
    cudaFree(Ad);
    cudaFree(Bd);
}





int main(){
    doReduce0();
    doReduce1();
    doReduce2();
    doReduce3();
    doReduce6();
}
