Assignment 2 (CS427)

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