/*#ifndef CUDACC
#define CUDACC
#endif
*/
#include <time.h>
#include <stdio.h>
#include “cuda_runtime.h”
#include “device_launch_parameters.h”
//#include <cuda.h>
//#include <cuda_runtime_api.h>
//#include <device_functions.h>
#define N 32
#define TILE_WIDTH 2
double r2()
{
return (double)rand() / (double)RAND_MAX ;
}
__global__ void test(double *a, double *b, double *c,int num)
{
/* shared int A_tile[ TILE_WIDTH ][ TILE_WIDTH ];
shared int B_tile[ TILE_WIDTH ][ TILE_WIDTH ];
double accu=0;
// target element coordinates
int row = blockIdx.y * TILE_WIDTH + threadIdx.y;
int column = blockIdx.x * TILE_WIDTH + threadIdx.x;
// compute target element value
for(int i=0;i<(N/(int)TILE_WIDTH);i++){
// move the tiles and update shared memory value for new tile positions
if(row < N && (i*TILE_WIDTH + threadIdx.x)<N)
A_tile[threadIdx.y][threadIdx.x] = a[row*N + i*TILE_WIDTH + threadIdx.x];
// else
// A_tile[threadIdx.y][threadIdx.x] = 0;
if(column < N && (iTILE_WIDTH + threadIdx.y)<N)
B_tile[threadIdx.x][threadIdx.y] = b[(iTILE_WIDTH + threadIdx.y)*N + column];
//else
//B_tile[threadIdx.y][threadIdx.x] = 0;
// after the entire tile's values are available, proceed
__syncthreads();
for(int j=0;j<TILE_WIDTH;j++)
accu += A_tile[threadIdx.y][j] * B_tile[j][threadIdx.x];
// after the entire tile's values have been used, proceed
__syncthreads();
}
if(row < N && column < N)
c[row*N+column] = accu;
*/
__shared__ int A_tile[ TILE_WIDTH ][ TILE_WIDTH ];
__shared__ int B_tile[ TILE_WIDTH ][ TILE_WIDTH ];
double accu=0;
for(int tileIdx = 0; tileIdx<(num/ blockDim.x-1) ; tileIdx++)
{
int i = blockIdx.y * blockDim.y + threadIdx.y ;
int j = tileIdx * blockDim.x + threadIdx.x ;
A_tile[threadIdx.y][threadIdx.x] = a[i*(num/ blockDim.x-1) +tileIdx];
B_tile[threadIdx.x][threadIdx.y] = b[tileIdx*(num/ blockDim.x-1)+j];
__syncthreads();
for(int k=0 ; k<2 ; k++)
accu = accu + A_tile[threadIdx.y][k] * B_tile[k][threadIdx.x];
__syncthreads();
}
int i = blockIdx.y * blockDim.y + threadIdx.y ;
int j = blockIdx.x * blockDim.x + threadIdx.x ;
c[i*num+j]=accu;
__syncthreads();
/*
// correct used code
int i = blockIdx.y * blockDim.y + threadIdx.y;
int j = blockIdx.x * blockDim.x + threadIdx.x;
double accu=0;
if(i<num && j<num)
{
accu = accu + a[i*num+n] * b[n*num+j];
}
c[i*num+j]=accu;
// correct used code
*/
}
int main()
{
float ms;
cudaEvent_t start_g,stop_g;
clock_t start, end;
double cpu_time_used;
//declare host and device arrays;
double a[N][N], *a_device, b[N][N], *b_device, c[N][N], *c_device;
//declaring size of device array
const size_t a_size = sizeof(double) * size_t(N*N);
cudaMalloc((void **)&a_device, a_size);
cudaMalloc((void **)&b_device, a_size);
cudaMalloc((void **)&c_device, a_size);
/*
double d_A[4][4];
double d_B[4][4];
double d_C[4][4];
double C[4][4];
double check[4][4]; */
// initializing 2d arrays d_A and d_B between 0 and 1 as well as d_C with 0;
for(int i=0;i<N;i++)
for(int j=0;j<N;j++)
{
a[i][j]=r2();
b[i][j]=r2();
// c[i][j] = 0;
}
cudaMemcpy(a_device, a, a_size, cudaMemcpyHostToDevice);
cudaMemcpy(b_device, b, a_size, cudaMemcpyHostToDevice);
dim3 dimBlock(16,16); //32 threads in one block in x and y directions;
// dim3 dimGrid(1, 1); // 2 blocks in one grid in both directions;
dim3 dimGrid(N/dimBlock.x, N/dimBlock.y);// 16 blocks in one grid in both directions;
// cudaDeviceSynchronize() ;
cudaEventCreate(&start_g);
cudaEventCreate(&stop_g);
cudaEventRecord(start_g, 0);
cudaEventRecord(stop_g, 0);
//GPU timer starts;
test<<<dimGrid, dimBlock>>>(a_device,b_device,c_device,N);
// cudaDeviceSynchronize();
cudaMemcpy(&c, c_device, a_size, cudaMemcpyDeviceToHost);
cudaEventSynchronize(stop_g);//GPU timer ends;
cudaEventElapsedTime(&ms, start_g, stop_g);
// cudaMemcpy(&C,d_C,4*4 , cudaMemcpyDeviceToHost);
cudaEventDestroy(start_g);
cudaEventDestroy(stop_g);
printf("GPU: %f ms\n",ms);
// cudaDeviceSynchronize() ;
start = clock();
//CPU matrix multiplication
double check[N][N];
int i, j, k;
for (i = 0; i < N; i++)
{
for (j = 0; j < N; j++)
{
check[i][j] = 0;
for (k = 0; k < N; k++)
check[i][j] += a[i][k]*b[k][j];
}
}
end = clock();
//printf("start:=%f/n"start);
//printf("end:=%f/n"end);
cpu_time_used = ((double) (end - start)) / CLOCKS_PER_SEC;
cpu_time_used*=1000;
printf (" cpu_time_used=%fms\n",cpu_time_used);
//CPU-GPU matrix matching
for(int i =0;i<N;i++)
{
for(int j=0;j<N;j++)
{
if(check[i][j]==c[i][j])
{ // printf("correct\n");
//printf("check = %f, c= %f\n",check[i][j],c[i][j]);
continue;
}
else
{
printf("not equal at i= %d and at j= %d\n",i,j);
printf("check = %f, c= %f",check[i][j],c[i][j]);
// return 0;
}
}
}
cudaFree(a_device); cudaFree(b_device); cudaFree(c_device);
return 0;
}
/*
#include “cuda_runtime.h”
#include “device_launch_parameters.h”
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);
global void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}
int main()
{
const int arraySize = 5;
const int a[arraySize] = { 1, 2, 3, 4, 5 };
const int b[arraySize] = { 10, 20, 30, 40, 50 };
int c[arraySize] = { 0 };
// Add vectors in parallel.
cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addWithCuda failed!");
return 1;
}
printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
c[0], c[1], c[2], c[3], c[4]);
// cudaDeviceReset must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaDeviceReset();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceReset failed!");
return 1;
}
return 0;
}
// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
int *dev_a = 0;
int *dev_b = 0;
int *dev_c = 0;
cudaError_t cudaStatus;
// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// Allocate GPU buffers for three vectors (two input, one output) .
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
// Launch a kernel on the GPU with one thread for each element.
addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
// Check for any errors launching the kernel
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
goto Error;
}
// cudaDeviceSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
goto Error;
}
// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
*/