CUDA学习，使用shared memory实现Matrix Parallel Add

#include< stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include
#include
#include

using namespace std;

#define N (2048*2048)
#define THREADS_PER_BLOCK 256  //16*16
#define TILE_WIDTH 16
#define width 2048

__global__ void Add(int *dev_a, int *dev_b,int *dev_c)
{

int bx=blockIdx.x;
int by=blockIdx.y;

int tx=threadIdx.x;
int ty=threadIdx.y;

int i=bx*blockDim.x+by*blockDim.y*width+ty*width+tx;

__shared__ int A[TILE_WIDTH][TILE_WIDTH];
__shared__ int B[TILE_WIDTH][TILE_WIDTH];

A[ty][tx]=dev_a[i];
B[ty][tx]=dev_b[i];

__syncthreads();

dev_c[i]=A[ty][tx]+B[ty][tx];

/*
int i=bx*blockDim.x+by*blockDim.y*width+ty*width+tx;
dev_c[i]=dev_a[i]+dev_b[i];
*/
}

int main( void ) {

int *a, *b, *c; // host copies of a, b, c

int *dev_a, *dev_b, *dev_c; // device copies of a, b, c

int size = N * sizeof( int); // we need space for N integers

// allocate device copies of a, b, c
cudaMalloc( (void**)&dev_a, size );
cudaMalloc( (void**)&dev_b, size );
cudaMalloc( (void**)&dev_c, size );

a = (int*)malloc( size );
b = (int*)malloc( size );
c = (int*)malloc( size );

//random_ints( a, N );
//random_ints( b, N );

for (int i = 0; i < N; ++i)
{
a[i] = rand();
b[i] = rand();
}

// copy inputs to device
cudaMemcpy( dev_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy( dev_b, b, size, cudaMemcpyHostToDevice);

dim3 dimGrid(128,128,1);
dim3 dimBlock(16,16,1);

int sharesize = 16*16*sizeof(int);

cudaEvent_t timeStartEvent,timeEndEvent;
cudaEventCreate( &timeStartEvent, 0);
cudaEventCreate(&timeEndEvent, 0);
cudaEventRecord( timeStartEvent, 0);

// launch add() kernel with blocks and threads
Add<<< dimGrid, dimBlock,sharesize >>>( dev_a, dev_b, dev_c);

// copy device result back to host copy of c
cudaMemcpy( c, dev_c, size, cudaMemcpyDeviceToHost);

// verify the data returned to the host is correct

for (int i = 0; i < N; i++)
{
assert(c[i] == a[i]+b[i] );
}

free( a );
free( b );
free( c );
cudaFree( dev_a);
cudaFree( dev_b);
cudaFree( dev_c);

cudaEventRecord( timeEndEvent, 0) ;
cudaEventSynchronize( timeEndEvent ) ;
float elapsedTime = 0 ;
cudaEventElapsedTime( & elapsedTime, timeStartEvent, timeEndEvent ) ;

std::cout<< "elapsedTime is  " << elapsedTime << " ms. ";
cudaEventDestroy( timeStartEvent ) ;
cudaEventDestroy( timeEndEvent ) ;

return 0;
}