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

__global__ void full_coalesced_read(int *d_array,int size)
{
        int x=d_array[blockIdx.x*blockDim.x+threadIdx.x];
}

__global__ void partial_coalesced_read(int *d_array,int size)
{
                int x=d_array[blockIdx.x*blockDim.x+threadIdx.x+1];
}

int main(int argc, char** argv)
{
        CUT_DEVICE_INIT(argc,argv);

        int *d_array;
        int size=512;

        cudaMalloc((void **)&d_array,sizeof(int)*(size*2));

        unsigned int timer=0,timer1=0;
        CUT_SAFE_CALL( cutCreateTimer( &timer));
        CUT_SAFE_CALL( cutCreateTimer( &timer1));

        float avg_timer=0,avg_timer1=0;
        for(int i=0;i<100;i++ )
        {
                CUT_SAFE_CALL( cutStartTimer( timer));
                full_coalesced_read<<<size/512,512>>>(d_array,size);
                cudaThreadSynchronize();
                CUT_SAFE_CALL( cutStopTimer( timer));
                avg_timer += cutGetTimerValue( timer);
        }
        for(int i=0;i<100;i++ )
        {
                CUT_SAFE_CALL( cutStartTimer( timer1));
                partial_coalesced_read<<<size/512,512>>>(d_array,size);
                cudaThreadSynchronize();
                CUT_SAFE_CALL( cutStopTimer( timer1));
                avg_timer1 += cutGetTimerValue( timer1);
        }

        printf( "Time with coalesced reads : %f (ms)\n", avg_timer / 100);
        printf( "Time with uncoalesced reads : %f (ms)\n", avg_timer1 / 100);

}