OpenACC 绘制曼德勃罗集

▶ 书上第四章，用一系列步骤优化曼德勃罗集的计算过程。

● 代码

 // constants.h

 const unsigned int WIDTH=;

 const unsigned int HEIGHT=;

 const unsigned int MAX_ITERS=;

 const unsigned int MAX_COLOR=;

 const double xmin=-1.7;

 const double xmax=.;

 const double ymin=-1.2;

 const double ymax=1.2;

 const double dx = (xmax - xmin) / WIDTH;

 const double dy = (ymax - ymin) / HEIGHT;

 // mandelbrot.h

 #pragma acc routine seq

 unsigned char mandelbrot(int Px, int Py);

 // mandelbrot.cpp

 #include <cstdio>

 #include <cstdlib>

 #include <fstream>

 #include "mandelbrot.h"

 #include "constants.h"

 using namespace std;

 unsigned char mandelbrot(int Px, int Py)

 {

     const double x0 = xmin + Px * dx, y0 = ymin + Py * dy;

     double x = 0.0, y = 0.0;

     int i;

     for(i=; x * x + y * y < 4.0 && i < MAX_ITERS; i++)

     {

         double xtemp = x * x - y * y + x0;

         y =  * x * y + y0;

         x = xtemp;

     }

     return (double)MAX_COLOR * i / MAX_ITERS;

 }

 // main.cpp

 #include <cstdio>

 #include <cstdlib>

 #include <fstream>

 #include <cstring>

 #include <omp.h>

 #include <openacc.h>

 #include "mandelbrot.h"

 #include "constants.h"

 using namespace std;

 int main()

 {

     unsigned char *image = (unsigned char*)malloc(sizeof(unsigned int) * WIDTH * HEIGHT);

     FILE *fp=fopen("image.pgm","wb");

     fprintf(fp,"P5\n\"#comment\"\n%d %d\n%d\n",WIDTH, HEIGHT, MAX_COLOR);

     acc_init(acc_device_nvidia);

 #pragma acc parallel num_gangs(1)

     {

         image[] = ;

     }

     double st = omp_get_wtime();

 #pragma acc parallel loop

     for(int y = ; y < HEIGHT; y++)

     {

         for(int x = ; x < WIDTH; x++)

             image[y * WIDTH + x] = mandelbrot(x, y);

     }

     double et = omp_get_wtime();

     printf("Time: %lf seconds.\n", (et-st));

     fwrite(image,sizeof(unsigned char),WIDTH * HEIGHT, fp);

     fclose(fp);

     free(image);

     return ;

 }

● 输出结果

// Ubuntu:

cuan@CUAN:/media/cuan/02FCDA52FCDA4019/Code/ParallelProgrammingWithOpenACC-master/Chapter04/cpp$ pgc++ -std=c++ -acc -mp -fast -Minfo -c mandelbrot.cpp

mandelbrot(int, int):

      , Generating acc routine seq

         Generating Tesla code

     , FMA (fused multiply-add) instruction(s) generated

     , Loop not vectorized/parallelized: potential early exits

     , FMA (fused multiply-add) instruction(s) generated

cuan@CUAN:/media/cuan/02FCDA52FCDA4019/Code/ParallelProgrammingWithOpenACC-master/Chapter04/cpp$ pgc++ -std=c++ -acc -mp -fast -Minfo main.cpp mandelbrot.o -o acc1.exe

main.cpp:

main:

     , Accelerator kernel generated

         Generating Tesla code

         Generating implicit copyout(image[])

     , Accelerator kernel generated

         Generating Tesla code

         , #pragma acc loop gang /* blockIdx.x */

         , #pragma acc loop vector(128) /* threadIdx.x */

     , Generating implicit copy(image[:])

     , Loop is parallelizable

         Loop not vectorized/parallelized: contains call

cuan@CUAN:/media/cuan/02FCDA52FCDA4019/Code/ParallelProgrammingWithOpenACC-master/Chapter04/cpp$ ./acc1.exe

Time: 0.646578 seconds.

● 优化 03，变化仅在 main.cpp 中

 // main.cpp

 #include <cstdio>

 #include <cstdlib>

 #include <fstream>

 #include <cstring>

 #include <omp.h>

 #include <openacc.h>

 #include "mandelbrot.h"

 #include "constants.h"

 using namespace std;

 int main()

 {

     const int num_blocks = , block_size = HEIGHT / num_blocks * WIDTH;

     unsigned char *image=(unsigned char*)malloc(sizeof(unsigned int) * WIDTH * HEIGHT);

     FILE *fp=fopen("image.pgm","wb");

     fprintf(fp,"P5\n\"#comment\"\n%d %d\n%d\n",WIDTH, HEIGHT, MAX_COLOR);

     acc_init(acc_device_nvidia);

 #pragma acc parallel num_gangs(1)

     {

         image[] = ;

     }

     double st = omp_get_wtime();

 #pragma acc data create(image[WIDTH*HEIGHT])

     {

         for(int block = ; block < num_blocks; block++)

         {

             const int start = block * (HEIGHT/num_blocks), end   = start + (HEIGHT/num_blocks);

 #pragma acc parallel loop async(block)

             for(int y=start;y<end;y++)

             {

                 for(int x=;x<WIDTH;x++)

                     image[y*WIDTH+x]=mandelbrot(x,y);

             }

 #pragma acc update self(image[block*block_size:block_size]) async(block)

         }

     }

 #pragma acc wait

     double et = omp_get_wtime();

     printf("Time: %lf seconds.\n", (et-st));

     fwrite(image,sizeof(unsigned char), WIDTH * HEIGHT, fp);

     fclose(fp);

     free(image);

     return ;

 }

● 输出结果

// Ubuntu:

cuan@CUAN:/media/cuan/02FCDA52FCDA4019/Code/ParallelProgrammingWithOpenACC-master/Chapter04/cpp/task3$ pgc++ -std=c++ -acc -mp -fast -Minfo -c mandelbrot.cpp

mandelbrot(int, int):

     , Generating acc routine seq

         Generating Tesla code

     , FMA (fused multiply-add) instruction(s) generated

     , Loop not vectorized/parallelized: potential early exits

     , FMA (fused multiply-add) instruction(s) generated

cuan@CUAN:/media/cuan/02FCDA52FCDA4019/Code/ParallelProgrammingWithOpenACC-master/Chapter04/cpp/task3$ pgc++ -std=c++ -acc -mp -fast -Minfo main.cpp mandelbrot.o -o acc2.exe

main.cpp:

main:

     , Accelerator kernel generated

         Generating Tesla code

         Generating implicit copyout(image[])

     , Generating create(image[:])

     , Accelerator kernel generated

         Generating Tesla code

         , #pragma acc loop gang /* blockIdx.x */

         , #pragma acc loop vector(128) /* threadIdx.x */

     , Loop is parallelizable

         Loop not vectorized/parallelized: contains call

     , Generating update self(image[block*:])

cuan@CUAN:/media/cuan/02FCDA52FCDA4019/Code/ParallelProgrammingWithOpenACC-master/Chapter04/cpp/task3$ ./acc2.exe

Time: 0.577263 seconds.

● 优化 05，添加异步计算

 // main.cpp

 #include <cstdio>

 #include <cstdlib>

 #include <fstream>

 #include <cstring>

 #include <omp.h>

 #include <openacc.h>

 #include "mandelbrot.h"

 #include "constants.h"

 using namespace std;

 int main()

 {

   const int num_blocks=, block_size = HEIGHT / num_blocks * WIDTH;

   unsigned char *image=(unsigned char*)malloc(sizeof(unsigned int) * WIDTH * HEIGHT);

   FILE *fp = fopen("image.pgm", "wb");

   fprintf(fp,"P5\n\"#comment\"\n%d %d\n%d\n",WIDTH, HEIGHT, MAX_COLOR);

   const int num_gpus = acc_get_num_devices(acc_device_nvidia);

 #pragma omp parallel num_threads(num_gpus)

     {

         acc_init(acc_device_nvidia);

         acc_set_device_num(omp_get_thread_num(),acc_device_nvidia);

     }

         printf("Found %d NVIDIA GPUs.\n", num_gpus);

     double st = omp_get_wtime();

 #pragma omp parallel num_threads(num_gpus)

     {

         int queue = ;

         int my_gpu = omp_get_thread_num();

         acc_set_device_num(my_gpu,acc_device_nvidia);

         printf("Thread %d is using GPU %d\n", my_gpu, acc_get_device_num(acc_device_nvidia));

 #pragma acc data create(image[WIDTH*HEIGHT])

         {

 #pragma omp for schedule(static, 1)

             for(int block = ; block < num_blocks; block++)

             {

                 const int start = block * (HEIGHT/num_blocks), end   = start + (HEIGHT/num_blocks);

 #pragma acc parallel loop async(queue)

                 for(int y=start;y<end;y++)

                 {

                     for(int x=;x<WIDTH;x++)

                         image[y*WIDTH+x]=mandelbrot(x,y);

                 }

 #pragma acc update self(image[block*block_size:block_size]) async(queue)

         queue = (queue + ) % ;

             }

         }

 #pragma acc wait

     }

     double et = omp_get_wtime();

     printf("Time: %lf seconds.\n", (et-st));

     fwrite(image,sizeof(unsigned char), WIDTH * HEIGHT, fp);

     fclose(fp);

     free(image);

     return ;

 }

● 输出结果

// Ubuntu:

cuan@CUAN:/media/cuan/02FCDA52FCDA4019/Code/ParallelProgrammingWithOpenACC-master/Chapter04/cpp/task5.multithread$ pgc++ -std=c++ -acc -mp -fast -Minfo -c mandelbrot.cpp

mandelbrot(int, int):

     , Generating acc routine seq

         Generating Tesla code

     , FMA (fused multiply-add) instruction(s) generated

     , Loop not vectorized/parallelized: potential early exits

     , FMA (fused multiply-add) instruction(s) generated

cuan@CUAN:/media/cuan/02FCDA52FCDA4019/Code/ParallelProgrammingWithOpenACC-master/Chapter04/cpp/task5.multithread$ pgc++ -std=c++ -acc -mp -fast -Minfo main.cpp mandelbrot.o -o acc3.exe

main.cpp:

main:

     , Parallel region activated

     , Parallel region terminated

     , Parallel region activated

     , Generating create(image[:])

     , Parallel loop activated with static cyclic schedule

     , Accelerator kernel generated

         Generating Tesla code

         , #pragma acc loop gang /* blockIdx.x */

         , #pragma acc loop vector(128) /* threadIdx.x */

     , Loop is parallelizable

         Loop not vectorized/parallelized: contains call

     , Generating update self(image[block*:])

     , Barrier

     , Parallel region terminated

cuan@CUAN:/media/cuan/02FCDA52FCDA4019/Code/ParallelProgrammingWithOpenACC-master/Chapter04/cpp/task5.multithread$ ./acc3.exe

Found  NVIDIA GPUs.

Thread  is using GPU

Time: 0.497450 seconds.

● nvprof 的结果汇总，三张图分别为 “并行和数据优化”，“优化 03（分块分流）” 和 “优化 05（分块调度）”