Archive

Posts Tagged ‘CUDA’

Assignment 3 (CS427) – Game of Life

February 3, 2010 Leave a comment

1. โปรแกรม Game of Life ที่เป็นการทำงานแบบ Sequential จำเป็นจะต้องมีอาร์เรย์จำนวน 2 ชุดในการเก็บข้อมูลของเวลาที่ต่างกันในการทำงาน และใช้ Pointer ในการสลับอาร์เรย์เพื่อประมวลผล

#include <stdio.h>
#include <time.h>
#define MAX_WIDTH 1000

FILE *f;
char board1[MAX_WIDTH * MAX_WIDTH];
char board2[MAX_WIDTH * MAX_WIDTH];
char *currentBoard;
char *pastBoard;
int maxx, maxy;

void init();                           	        /* initial board */
void rand_init_pop(int rand_size);      	/* initial population */
int in_bound(int i, int j);
int neighbor_num(int i, int j);
void process();
void print_board();
void write_to_file();

int main(){
     int i, t = 10;
     int randnum = 0;

     printf("Enter board width : ");
     scanf("%d", &maxx);
     printf("Enter board height : ");
     scanf("%d", &maxy);
     printf("Enter number of random population : ");
     scanf("%d", &randnum);
     printf("Enter time step : ");
     scanf("%d", &t);

     pastBoard = board1;
     currentBoard = board2;

     init();
     rand_init_pop(randnum);

     f = fopen("sequential_board.txt", "w");

     for(i=0;i<t;i++){
          fprintf(f, "t = %d\n", i);
          write_to_file();
          process();
     }
     fprintf(f, "t = %d\n", i);
     write_to_file();

     fclose(f);

     return 0;
}
void init(){
     int i, j;
     for(i=0;i<maxy;i++) for(j=0;j<maxx;j++)
          board1[i * maxx + j] = board2[i * maxx + j] = 'O';
}
void rand_init_pop(int rand_size){
     int i, x, y;
     srand( time(0));

     for(i=0;i<rand_size;i++){
          y = rand() % maxy;
          x = rand() % maxx;
          if(pastBoard[y * maxx + x] == 'X') i--;
          else pastBoard[y * maxx + x] = 'X';
     }
}
int in_bound(int i, int j){
     if(i<0 || j=maxy || j>=maxy)    return 0;
     else                                    return 1;
}
int neighbor_num(int i, int j){
     int num = 0;
     if( in_bound(i-1,j-1) && pastBoard[(i-1) * maxx + (j-1)]=='X' ) num++;
     if( in_bound(i-1, j ) && pastBoard[(i-1) * maxx + ( j )]=='X' ) num++;
     if( in_bound( i ,j-1) && pastBoard[( i ) * maxx + (j-1)]=='X' ) num++;
     if( in_bound(i+1,j+1) && pastBoard[(i+1) * maxx + (j+1)]=='X' ) num++;
     if( in_bound(i+1, j ) && pastBoard[(i+1) * maxx + ( j )]=='X' ) num++;
     if( in_bound( i ,j+1) && pastBoard[( i ) * maxx + (j+1)]=='X' ) num++;
     if( in_bound(i+1,j-1) && pastBoard[(i+1) * maxx + (j-1)]=='X' ) num++;
     if( in_bound(i-1,j+1) && pastBoard[(i-1) * maxx + (j+1)]=='X' ) num++;
     return num;
}
void process(){
     int i,j;
     char* tmp;

     for(i=0;i<maxy;i++){
          for(j=0;j<maxx;j++){
               //For a space that is 'populated'
               if(pastBoard[i * maxx + j]=='X'){
                    //Each cell with one or no neighbors dies, as if by loneliness.
                    if(neighbor_num(i,j) <= 1) currentBoard[i * maxx + j] = 'O';
                    //Each cell with two or three neighbors survives.
                    else if(neighbor_num(i,j) <= 3) currentBoard[i * maxx + j] = 'X';
                    // Each cell with four or more neighbors dies, as if by overpopulation.
                    else currentBoard[i * maxx + j] = 'O';
               }
               //For a space that is 'empty' or 'unpopulated'
               else{
                    // Each cell with three neighbors becomes populated.
                    if(neighbor_num(i,j) == 3) currentBoard[i * maxx + j] = 'X';
                    else currentBoard[i * maxx + j] = 'O';
               }
          }
     }
     tmp = pastBoard;
     pastBoard = currentBoard;
     currentBoard = tmp;
}
void write_to_file(){
     int i, j;
     for(i=0;i<maxy;i++){
          for(j=0;j<maxx;j++){
               if(pastBoard[i * maxx + j]=='X')    fprintf(f, "1”);
               else                                fprintf(f, "0”);
          }
          fprintf(f, "\n");
     }
     fprintf(f, "\n");
}

เมื่อเริ่มต้นโปรแกรม โปรแกรมจะให้ใส่ข้อมูลขนาดความกว้าง และความสูงของตารางที่จะใช้ จากนั้นจะถามถึงจำนวน Cell ที่มีชีวิตเริ่มต้นที่ต้องการกำหนด โดยจะมีตำแหน่งแบบสุ่ม และสุดท้ายจะถามถึง จำนวนหน่วยเวลาที่ต้องการทำการ Simulation เมื่อใส่ข้อมูลครบแล้วโปรแกรมจะทำงาน และเขียนข้อมูลของตารางในแต่ละช่วงเวลาไปยังไฟล์ที่ชื่อ “sequential_board.txt”


2. การทำงานของโปรแกรม Game of Life บนจีพียู ใน 1 time step จะต้องเข้าถึงหน่วยความจำเฉลี่ยประมาณ 8 ครั้งต่อ 1 เทรด ดังนั้นเพื่อให้มีความรวดเร็วในการทำงาน จะต้องทำการคัดลอดข้อมูลจาก Global Memory มายัง Shared Memory ก่อน เพื่อให้เข้าถึงหน่วยความจำได้เร็วขึ้น โดยวิธีที่ใช้คือการแบ่งการทำงานออกเป็นเทรดบล็อคโดยที่ในตอนเริ่มต้นการทำงาน แต่ละเทรดบล็อคจะต้องคัดลอกข้อมูลในส่วนของ cell ที่มีพื้นที่ติดกันกับพื้นที่ที่เทรดบล็อคนั้นรับผิดชอบมาด้วย แต่ในการเขียนข้อมูลจะเขียนข้อมูลลงใน Global Memory ในช่องที่เทรดบล็อคนั้นรับผิดชอบเท่านั้น

#include <stdio.h>
#include <math.h>
#include <time.h>
#define MAX_WIDTH 500
#define TILE_WIDTH 16
#define TY threadIdx.y
#define TX threadIdx.x

void rand_init_pop(int *board, int maxx, int maxy, int rand_size);
void write_to_file(FILE *f, int *board, int maxx, int maxy);

__global__ void process(int *board, int maxx, int maxy);

int main(){
     int *board, *d_board, i, t;
     int maxx, maxy, size, randnum;
     FILE *f;

     printf("Enter board width : ");
     scanf("%d", &maxx);
     printf("Enter board height : ");
     scanf("%d", &maxy);
     printf("Enter number of random population : ");
     scanf("%d", &randnum);
     printf("Enter time step : ");
     scanf("%d", &t);
     size = sizeof(int) * maxx * maxy;

     f = fopen("cuda_board.txt", "w");

     board = (int*)malloc(size);
     cudaMalloc((void**)&d_board, size);
     rand_init_pop(board, maxx, maxy, randnum);

     dim3 blockDim(TILE_WIDTH+2, TILE_WIDTH+2);
     dim3 gridDim(ceil(maxy*1.0/TILE_WIDTH), ceil(maxx*1.0/TILE_WIDTH));

     cudaMemcpy(d_board, board, size, cudaMemcpyHostToDevice);

     for(i=0;i<t;i++){
          fprintf(f, "t = %d\n", i);
          write_to_file(f, board, maxx, maxy);

          process<<>>(d_board, maxx, maxy);
          cudaMemcpy(board, d_board, size, cudaMemcpyDeviceToHost);
     }
     fprintf(f, "t = %d\n", i);
     write_to_file(f, board, maxx, maxy);

     fclose(f);
     free(board);
     cudaFree(d_board);

     return 0;
}

void init(int *board, int maxx, int maxy){
     int i, j;
     for(i=0;i<maxy;i++) for(j=0;j<maxx;j++)
          board[i * maxx + j] = 0;
}
void rand_init_pop(int *board, int maxx, int maxy, int rand_size){
     int i, x, y;
     srand( time(0));
     init(board, maxx, maxy);

     for(i=0;i<rand_size;i++){
          y = rand() % maxy;
          x = rand() % maxx;
          if(board[y * maxx + x] == 1) i--;
          else board[y * maxx + x] = 1;
     }
}
void write_to_file(FILE *f, int *board, int maxx, int maxy){
     for(int i=0;i<maxy;i++){
          for(int j=0;j<maxx;j++)
               fprintf(f, "%d", board[i * maxx + j]);
          fprintf(f, "\n");
     }
     fprintf(f, "\n");
}

__global__ void process(int *board, int maxx, int maxy){
     __shared__ int sboard[TILE_WIDTH+2][TILE_WIDTH+2];

     int i = (blockIdx.x * TILE_WIDTH) + TX - 1;
     int j = (blockIdx.y * TILE_WIDTH) + TY - 1;

     if(i<=maxy && j=0 && j>=0 && i<maxy && j0 && TX0 && TY<TILE_WIDTH){
          int n = 0;
          n += sboard[TY+1][TX+1];
          n += sboard[TY+1][TX-1];
          n += sboard[TY-1][TX-1];
          n += sboard[TY-1][TX+1];
          n += sboard[TY+1][ TX ];
          n += sboard[ TY ][TX+1];
          n += sboard[ TY ][TX-1];
          n += sboard[TY-1][ TX ];

          if(sboard[TY][TX]==1 && n!=2 && n!=3)
               board[boardidx] = 0;
          else if(sboard[TY][TX]==0 && n==3)
               board[boardidx] = 1;
          }
     }
}

เมื่อรันโปรแกรม โปรแกรมจะถามข้อมูลเช่นเดียวกับโปรแกรม Sequential แต่ข้อมูลจะเขียนลงไปที่ไฟล์ที่มีชื่อว่า “cuda_board.txt”

สามารถดาวน์โหลดไฟล์โปรแกรมได้ที่ด้านล่าง
GameOfLife.c
GameOfLifeCUDA.cu

Advertisements

Assignment 2 (CS427)

January 10, 2010 Leave a comment

1. การทำให้โปรแกรมมีลักษณะการทำงานแบบ Parallelism มากที่สุด จะต้องทำการวิเคราะห์ความสัมพันธ์ของโปรแกรม ในส่วนต่างๆว่ามีส่วนใดที่สามารถทำงานพร้อมกันได้โดยดูจากกราฟ

จากโปรแกรมตัวอย่างเมื่อเขียนความสัมพันธ์ในรูปแบบกราฟจะได้กราฟลักษณะข้างต้น ซึ่งหมายความว่าลูปที่ 2 และ 3 ที่เป็นส่วนที่คำนวนค่าของ c และ d สามารถทำงานพร้อมกันได้ ดังนั้นจึงสามารถใช้คำสั่ง #pragma omp sections ในส่วนนี้ และภายในลูปแต่ละลูปก็มีไม่ความสัมพันธ์กันระหว่างแต่ละการวนรอบสังเกตุได้จากการอ้างอิงถึง index เดียวกันในแต่ละการวนรอบเท่านั้น ดังนั้นทุกลูปสามารถใช้คำสั่ง #pragma omp for เพื่อให้เกิดการทำงานแบบขนานขึ้นในแต่ละลูป

โดยสามารถเขียนเป็นโปรแกรมได้ดังดังนี้

#include <omp.h>
#include <stdio.h>
#include <stdlib.h>
#define N 50

int main (int argc, char *argv[])
{
     int i, nthreads, tid;
     float a[N], b[N], c[N], d[N];
     float pa[N], pb[N], pc[N], pd[N];

     /* Sequential Program */
     for (i=0; i<N; i++) {
          a[i] = i * 1.5;
          b[i] = i + 22.35;
          c[i] = d[i] = 0.0;
     }

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

     /* Parallel Program with OpenMP */
     #pragma omp parallel for
     for (i=0; i<N; i++) {
          pa[i] = i * 1.5;
          pb[i] = i + 22.35;
          pc[i] = pd[i] = 0.0;
     }

     #pragma omp parallel sections
     {
          #pragma omp section
          #pragma omp parallel for
          for (i=0; i<N; i++){
               pc[i] = pa[i] + pb[i];
          }

          #pragma omp section
          #pragma omp parallel for
          for (i=0; i<N; i++){
               pd[i] = pa[i] * pb[i];
          }
     }

     /* check answers */
     for(i=0;i<N;i++){
          if(c[i]!=pc[i] || d[i]!=pd[i]){
               printf("Compute Error\n");
               break;
          }
     }
     if(i==N) printf("Corrected Answer\n");

     return 0;
}

โปรแกรมดังกล่าวจะทำงานทั้งหมดสองครั้งคือ คำนวนเชิงลำดับ (Sequential) และคำนวนแบบขนาน (Parallel) ด้วยการใช้ OpenMP หลังจากนั้นจะนำคำตอบมาตรวจสอบว่าการคำนวนแบบขนานให้ผลลัพท์เช่นเดียวกับการทำงานเชิงลำดับหรือไม่ ซึ่งหากคำนวนได้ถูกต้องจะปรากฎผลลัพท์ขึ้นมาเป็น Corrected Answer และจะปรากฎผลลัพท์ว่า Compute Error หากทำงานผิดพลาด ในที่นี้จะทำการคอมไพล์ด้วยโปรแกรม gcc เวอร์ชัน 4 ด้วยคำสั่ง

$gcc-4 –fopenmp program1.c –o program1

และเมื่อนำโปรแกรมไปรันทดสอบ ผลลัพท์ที่ได้คือ Corrected Answer หรือโปรแกรมทำงานได้ถูกต้อง ถึงแม้โปรแกรมนี้จะมีความเป็น Parallelism สูง แต่ในการทำงานจริงจะทำงานได้ช้า เนื่องจากโปรแกรมจะเสียเวลาในการสร้างเทรดขึ้นมาจำนวนมาก ซึ่งไม่คุ้มกับการคำนวนข้อมูลที่มีจำนวนน้อย ในการทำงานจริงเพื่อให้มีประสิทธิภาพสูงสุดอาจจะกำหนดให้มีการทำงานแบบขนานเพียงบางส่วนเท่านั้น เพื่อลด overhead ในการสร้างเทรด


2. ในข้อนี้จะทำการเปลี่ยนโปรแกรมเชิงลำดับให้เป็นโปรแกรมแบบขนานที่ทำงานได้บนจีพียู ซึ่งจะเขียนโปรแกรมที่คำนวนบนซีพียูและจีพียูจากนั้นจะนำคำตอบที่ได้มาตรวจสอบว่าได้ค่าที่เหมือนกัน ซึ่งการคำนวนค่าของ Floating-point บนซีพียูและจีพียูนั้นให้ค่าที่แตกต่างกันเล็กน้อย ดังนั้นเทียบค่าของ Floating-point ที่คำนวนได้ด้วยฟังก์ชัน floatIsEqual ในโปรแกรม

#include <stdio.h>
#include <time.h>
#define DELTA 0.001

int floatIsEqual(float a, float b){
     if( (b<a+DELTA) && (b>a-DELTA) ) return 1;
     else  return 0;
}

__global__ void GPUfunction(float *A, float *B, float *C){
     int i = blockIdx.x * blockDim.x + threadIdx.x;

     if(i<1024)       C[i] = A[i] + B[i];
     else if(i<2048) C[i] = A[i] * B[i];
     else if(i<3072) C[i] = A[i] - B[i];
     else if(i<4096) C[i] = A[i] / B[i];
}

void CPUfunction(float *A, float *B, float *C){
     int i;

     for(i=0; i<1024; i++)              C[i] = A[i] + B[i];
     for(i=1024; i<2048; i++)         C[i] = A[i] * B[i];
     for(i=2048; i<3072; i++)         C[i] = A[i] - B[i];
     for(i=3072; i<4096; i++)         C[i] = A[i] / B[i];
}

int main(){
     int i;
     float A[4096], B[4096], C[4096];
     float *hA, *hB, *hC;
     float *dA, *dB, *dC;
     int size = 4096 * sizeof(float);

     hA = (float*)malloc(size);
     hB = (float*)malloc(size);
     hC = (float*)malloc(size);

     cudaMalloc((void**)&dA, size);
     cudaMalloc((void**)&dB, size);
     cudaMalloc((void**)&dC, size);

     /* initialize */
     srand(time(0));
     for(i=0;i<4096;i++){
          A[i] = hA[i] = rand() % 40;
          // make sure that B is not zero
          B[i] = hB[i] = rand() % 39 + 1;
     }

     /* GPU code */
     cudaMemcpy(dA, hA, size, cudaMemcpyHostToDevice);
     cudaMemcpy(dB, hB, size, cudaMemcpyHostToDevice);
     GPUfunction<<<8, 512>>>(dA, dB, dC);
     cudaMemcpy(hC, dC, size, cudaMemcpyDeviceToHost);

     /* CPU code */
     CPUfunction(A, B, C);

     /* check answers from CPU and GPU are same */
     for(i=0;i<4096;i++){
          if(!floatIsEqual(C[i], hC[i])){
               printf("Compute Error\n");
               break;
          }
     }
     if(i==4096) printf("Corrected Answer\n");

     free(hA);
     free(hB);
     free(hC);

     cudaFree(dA);
     cudaFree(dB);
     cudaFree(dC);

     return 0;
}

ทำการคอมไพล์โปรแกรมด้วยคำสั่ง

$nvcc  program2.cu –o program2

ได้ผลลัพท์ออกมาเป็น Corrected Answer ซึ่งหมายถึงโปรแกรมทำงานได้ถูกต้อง การคำนวนแบบขนานจากจีพียูให้ผลลัพท์ตรงตามการคำนวนจากซีพียู


3. โปรแกรมนี้สามารถเพิ่มประสิทธิภาพได้ด้วยการใช้ Shared Memory และ Parallel Reduction ในการรวมข้อมูลหลังจากนั้นจึงทำการคำนวน โค้ดของโปรแกรมเหมือนกับข้อ 2 ยกเว้นในส่วนของฟังก์ชัน GPUfunction และ CPUfunction ซึ่งทำการเปลี่ยนแปลงดังนี้

GPUfunction :

__global__ void GPUfunction(float *A, float *B, float *C){
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     __shared__ float ssuma[512];
     float sum;
     int s, tid = threadIdx.x;

     /* copy global mem to shared mem synchonously */
     ssuma[tid] = A[i];
     __syncthreads();

     /* parallel reduction */
     for(s=256; s>0; s/=2){
          if(tid<s)  ssuma[tid] = ssuma[tid] + ssuma[tid+s];
          __syncthreads();
     }
     sum = ssuma[0];

     if(i<1024)        C[i] = sum + B[i];
     else if(i<2048) C[i] = sum * B[i];
     else if(i<3072) C[i] = sum - B[i];
     else if(i<4096) C[i] = sum / B[i];
}

CPUfunction :

void CPUfunction(float *A, float *B, float *C){
     int i;
     float SUMA[8];

     /* initialize SUMA */
     for(i=0;i<8;i++)           SUMA[i] = 0.0;
     for(i=0;i<4096;i++)     SUMA[i/512] += A[i];

     for(i=0; i<1024; i++)
          C[i] = SUMA[i/512] + B[i];
     for(i=1024; i<2048; i++)
          C[i] = SUMA[i/512] * B[i];
     for(i=2048; i<3072; i++)
          C[i] = SUMA[i/512] - B[i];
     for(i=3072; i<4096; i++)
          C[i] = SUMA[i/512] / B[i];

}

ทำการคอมไพล์โปรแกรมด้วย nvcc โดยใช้คำสั่ง

$nvcc  program3.cu –o program3

ซึ่งเมื่อรันโปรแกรมแล้วให้ผลลัพท์ถูกต้อง และแสดงข้อความว่า Corrected Answer เช่นเดียวกันในข้อ 2

ดาวน์โหลดโปรแกรมทั้ง 3 โปรแกรมได้ที่ด้านล่าง
program 1
program 2
program 3

ขั้นตอนการติดตั้ง CUDA Toolkit และ NVIDIA CUDA SDK

December 13, 2009 Leave a comment

1. หลังจากที่ Log in เข้าไปที่เครื่องเซิฟเวอร์แล้ว ทำการสร้างไดเรคเทอรี่ส่วนตัวขึ้นมา (ในที่นี้จะเก็บไว้ที่ /tmp/4909611727) ทำการคัดลอกไฟล์ cudatoolkit_2.3_linux_64_ubuntu9.04.run  และ cudasdk_2.3_linux.run จาก /home/kasidit มาไว้ยังไดเรคเทอรี่ที่สร้างขึ้น

2. ติดตั้ง CUDA Toolkit ด้วยคำสั่ง sh cudatoolkit_2.3_linux_64_ubuntu9.04.run จากนั้นโปรแกรมจะถามถึง path ที่ต้องการจะใช้ติดตั้ง พิมพ์ /tmp/4909611727 ลงไป

3. คอมไพเลอร์ของ CUDA หรือ nvcc นั้นจะเก็บไว้ที่ /tmp/4909611727/cuda/bin ให้ตั้งค่าของ $PATH และ $LD_LIBRARY_PATH ดังนี้

      export PATH=/tmp/4909611727/cuda/bin:$PATH
      export LD_LIBRARY_PATH=/tmp/4909611727/cuda/lib64:$LD_LIBRARY_PATH

4. ทดสอบว่าสามารถใช้งาน nvcc ได้ด้วยการพิมพ์คำสั่ง nvcc –version ซึ่งจะได้ผลลัพท์ดังนี้

5. ติดตั้ง NVIDIA CUDA SDK ด้วยคำสั่ง sh cudasdk_2.3_linux.run โปรแกรมจะถามถึง path ที่จะใช้ติดตั้ง ซึ่งถ้าไม่ใส่จะถูกกำหนดให้เป็น NVIDIA_GPU_Computing_SDK แต่ในที่นี้จะกำหนดให้เป็น /tmp/4909611727/NVIDAI_CUDA_SDK

6. คอมไพล์ไฟล์ตัวอย่างที่มากับ NVIDIA CUDA SDK ด้วยการเข้าไปที่ /tmp/4909611727/NVIDIA_CUDA_SDK/C แล้วใช้คำสั่ง make

7. ทดสอบโปรแกรมด้วยการรันโปรแกรมตัวอย่างที่อยู่ที่ /tmp/4909611727/NVIDIA_CUDA_SDK/C/bin/linux/release ซึ่งโปรแกรมตัวอย่างที่จะทดสอบคือ deviceQuery และ bandwidthTest

deviceQuery

โปรแกรม deviceQuery เป็นโปรแกรมที่ใช้แสดงคุณลักษณะของกราฟิกการ์ดที่ใช้ ซึ่งผลลัพธ์จากการรันโปรแกรม deviceQuery เป็นดังต่อไปนี้

bandwidthTest

โปรแกรม bandwidthTest จะแสดงให้เห็นถึงความสามารถในการส่งข้อมูลแบบต่างๆระหว่างหน่วยความจำหลัก และหน่วยความจำ Global ของ GPU ทดสอบโดยการส่งข้อมูลขนาด 32 MB ผลได้เป็นดังโปรแกรม

Vector Addition

เขียนโปรแกรมบวกเวกเตอร์ เซฟด้วยชื่อไฟล์ vecadd.cu แล้วทำการคัดลอกไปยังเครื่องเซิฟเวอร์ และคอมไพล์ด้วยคำสั่ง nvcc vecadd.cu และรันโปรแกรมจะได้ผลลัพท์ออกมาดังนี้

โค้ดของโปรแกรมนี้คือ


#include <stdio.h>
__global__ void VecAdd(float *A, float *B, float *C, int N){
	int idx = (blockIdx.x * blockDim.x) + threadIdx.x;

	if(idx<N){
		C[idx] = A[idx] + B[idx];
	}
}
int main(){
	float *h_A, *h_B, *h_C;
	float *d_A, *d_B, *d_C;
	int N = 200000;
	int size = sizeof(float) * N;
	int threadPerBlock = 512;
	int blockPerGrid = (N + threadPerBlock -1) / threadPerBlock;

	// Allocate host and device memory
	h_A = (float*)malloc(size); h_B = (float*)malloc(size); h_C = (float*)malloc(size);
	cudaMalloc((void**)&d_A, size); cudaMalloc((void**)&d_B, size); cudaMalloc((void**)&d_C, size); 

	// Initial value on vector A and B
	for(int i=0;i<N;i++)
		h_A[i] = rand()/200, h_B[i] = rand()/200;

	// Copy vector A and B to device
	cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
	cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

	// Call kernel
	VecAdd<<< blockPerGrid, threadPerBlock >>>(d_A, d_B, d_C, N);

	// Copy back vector C
	cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

	printf("%s\n", cudaGetErrorString( cudaGetLastError() ) );
	printf("Threads per block = %d\nBlocks per grid = %d\n", threadPerBlock, blockPerGrid);
	printf("Vector size = %d\nMemory used = %f MB\n", N, 3 * size / 1024.0 / 1024);

	// Deallocate host and device memory
	free(h_A); free(h_B); free(h_C);
	cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);

	return 0;
}

Categories: Parallel Programming Tags: , , ,