scan-array CUDA

Question

I'm trying to scan a simple array using CUDA but it seems there is something wrong with the code below..I am trying to find what i am doing wrong but i can't.Can anyone please help me?

#include <stdio.h>
#include <stdlib.h>


__global__ void prescan(int *g_odata, int *g_idata, int n){

    extern __shared__ int temp[];// allocated on invocation
    int thid = threadIdx.x;
    int offset = 1;
    temp[2*thid] = g_idata[2*thid]; // load input into shared memory
    temp[2*thid+1] = g_idata[2*thid+1];


    for (int d = n>>1; d > 0; d >>= 1){ // build sum in place up the tree
        __syncthreads();
        if (thid < d){

            int ai = offset*(2*thid+1)-1;
            int bi = offset*(2*thid+2)-1;
            temp[bi] += temp[ai];
        }
        offset *= 2;
    }

    if (thid == 0) { temp[n - 1] = 0; } // clear the last element

    for (int d = 1; d < n; d *= 2){ // traverse down tree & build scan
        offset >>= 1;
        __syncthreads();
        if (thid < d){

            int ai = offset*(2*thid+1)-1;
            int bi = offset*(2*thid+2)-1;
            int t = temp[ai];
            temp[ai] = temp[bi];
            temp[bi] += t;
        }
    }

    __syncthreads();

    g_odata[2*thid] = temp[2*thid]; // write results to device memory
    g_odata[2*thid+1] = temp[2*thid+1];
}

int main(int argc, char *argv[]){

    int i;
    int *input = 0;
    int *output = 0;
    int *g_idata = 0;
    int *g_odata = 0;

    int numblocks = 1;
    int radix = 16;

    input = (int*)malloc(numblocks*radix*sizeof(int));
    output = (int*)malloc(numblocks*radix*sizeof(int));

    cudaMalloc((void**)&g_idata, numblocks*radix*sizeof(int));
    cudaMalloc((void**)&g_odata, numblocks*radix*sizeof(int));

    for(i=0; i<numblocks*radix; i++){
            input[i] = 1 + 2*i;
    }

    for(i=0; i<numblocks*radix; i++){
        printf("%d   ", input[i]);
    }

    cudaMemcpy(g_idata, input, numblocks*radix*sizeof(int), cudaMemcpyHostToDevice);

    prescan<<<1,8>>>(g_odata, g_idata, numblocks*radix);

    cudaThreadSynchronize();

    cudaMemcpy(output, g_odata, numblocks*radix*sizeof(int), cudaMemcpyDeviceToHost);

    for(i=0; i<numblocks*radix; i++){
        printf("%d   ", output[i]);
    }

    free(input);
    free(output);
    cudaFree(g_idata);
    cudaFree(g_odata);

    return 0;
}

The output is this: 1 3 5 7 9 11 13 15 17 19 21 23 25 27 29 31 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0.I want to have this output: 1 3 5 7 9 11 13 15 17 19 21 23 25 27 29 31 0 1 4 9 16 25 36 49 64 81 100 121 144 169 196 225

Answer 1

Just go through this code to implement scan in parallel environment. The algorithm which I implemented here is Hillis Steele exclusive scan.I implemented algorithm through shared memory, it will definitely improve the execution time for the large data set.

#include<stdio.h>
#include<math.h>
__global__ void scan(int *d_in,int *d_out,int n)
{
    extern __shared__ int sdata[];
    int i;
    int tid = threadIdx.x;
    sdata[tid] = d_in[tid];
    for (i = 1; i <n; i <<= 1)
     {

        if (tid>=i)
         {
            sdata[tid] +=sdata[tid-i];
        }
        __syncthreads();
    }
    d_out[tid] = sdata[tid];
     __syncthreads();
 }

int main()
 {

    int h_in[16],h_out[16];
     int i,j;
     for (i = 0; i < 16; i++)
        h_in[i] = 2*i+1;
    for (i = 0; i < 16; i++)
        printf("%d ", h_in[i]);
   int *d_in;
   int *d_out;
   cudaMalloc((void**)&d_in, sizeof(int)* 16);
   cudaMalloc((void**)&d_out, sizeof(int)* 16);
   cudaMemcpy(d_in, h_in, sizeof(int) * 16, cudaMemcpyHostToDevice);

   scan <<<1, 16, sizeof(int)*16 >>>(d_in,d_out, 16);

   cudaMemcpy(h_out, d_out, sizeof(int) * 16, cudaMemcpyDeviceToHost);
   for (i = 0; i < 16; i++)
      printf("%d ", h_out[i]);
  return 0;
}