OpenACC Julia 图形

▶ 书上的代码，逐步优化绘制 Julia 图形的代码

● 无并行优化（手动优化了变量等）

#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>

#define N   (1024 * 8)

int julia(const float cre, const float cim, float zre, float zim, const int maxIter)// 计算单点迭代次数
{
    float zre2 = 0.0f, zim2 = 0.0f;
    for (int iter = 1; iter < maxIter; iter += 2)   // 一个迭代里计算两次
    {
        zre2 = zre * zre - zim * zim + cre, zim2 = 2 * zre * zim + cim;
        if (zre2 * zre2 + zim2 * zim2 > 4.0f)
            return iter;

        zre = zre2 * zre2 - zim2 * zim2 + cre, zim = 2 * zre2 * zim2 + cim;
        if (zre * zre + zim * zim > 4.0f)
            return iter;
    }
    return maxIter + 1 + (maxIter % 2);
}

int main()
{
    const int maxIter = 128;                                // 最大迭代次数
    const float cre = -0.8350, cim = -0.2321, h = 4.0f / N; // 迭代常数和画幅步长
    int *image = (int *)malloc(sizeof(int) * N * N);
    FILE *pf = fopen("R:/output.txt", "w");

    for (int i = 0; i < N; i++)
    {
        for (int j = 0; j < N; j++)
            fprintf(pf, "%d ", julia(cre, cim, i * h - 2.0f, j * h - 2.0f, maxIter));
        fprintf(pf, "
");
    }

    fclose(pf);
    free(image);
    //getchar();
    return 0;
}

● 输出结果（后面所有代码的输出都相同，不再写了）

● 改进 1，计算并行化

#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>

#define N   (1024 * 8)

#pragma acc routine seq
int julia(const float cre, const float cim, float zre, float zim, const int maxIter)
{
    float zre2 = 0.0f, zim2 = 0.0f;
    for (int iter = 1; iter < maxIter; iter += 2)
    {
        zre2 = zre * zre - zim * zim + cre, zim2 = 2 * zre * zim + cim;
        if (zre2 * zre2 + zim2 * zim2 > 4.0f)
            return iter;

        zre = zre2 * zre2 - zim2 * zim2 + cre, zim = 2 * zre2 * zim2 + cim;
        if (zre * zre + zim * zim > 4.0f)
            return iter;
    }
    return maxIter + 1 + (maxIter % 2);
}

int main()
{
    const int maxIter = 128;
    const float cre = -0.8350, cim = -0.2321, h = 4.0f / N;
    int *image = (int *)malloc(sizeof(int) * N * N);    

#pragma acc data copyout(image[0:N * N])    // 数据域
    {
#pragma acc kernels loop independent        // loop 并行化，强制独立
        for (int i = 0; i < N; i++)
        {
            for (int j = 0; j < N; j++)
                image[i * N + j] = julia(cre, cim, i * h - 2.0f, j * h - 2.0f, maxIter);
        }
    }
    /*// 注释掉写入文件的部分，防止 Nvvp 加入分析
    FILE *pf = fopen("R:/output.txt", "w");
    for (int i = 0; i < N; i++)
    {
        for (int j = 0; j < N; j++)
            fprintf(pf, "%d ", image[i * N + j]);
        fprintf(pf, "
");
    }
    fclose(pf);
    */    
    free(image);
    //getchar();
    return 0;
}

● 输出结果

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc -acc -Minfo main.c -o main_acc.exe
julia:
      9, Generating acc routine seq
         Generating Tesla code
     11, FMA (fused multiply-add) instruction(s) generated
     17, FMA (fused multiply-add) instruction(s) generated
main:
     30, Generating copyout(image[:67108864])
     33, Loop is parallelizable
         FMA (fused multiply-add) instruction(s) generated
     35, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         33, #pragma acc loop gang, vector(4) /* blockIdx.y threadIdx.y */
         35, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
     35, FMA (fused multiply-add) instruction(s) generated

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=35 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
PGI: "acc_shutdown" not detected, performance results might be incomplete.
 Please add the call "acc_shutdown(acc_device_nvidia)" to the end of your application to ensure that the performance results are complete.

Accelerator Kernel Timing data
D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c
  main  NVIDIA  devicenum=0
    time(us): 20,486
    30: data region reached 2 times
        49: data copyout transfers: 17
             device time(us): total=20,486 max=1,288 min=5 avg=1,205
    32: compute region reached 1 time
        35: kernel launched 1 time
            grid: [256x128]  block: [32x4]
             device time(us): total=0 max=0 min=0 avg=0

● 改进 2，分块计算，没有明显性能提升，为异步做准备

#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>

#define N   (1024 * 8)

#pragma acc routine seq
int julia(const float cre, const float cim, float zre, float zim, const int maxIter)
{
    float zre2 = 0.0f, zim2 = 0.0f;
    for (int iter = 1; iter < maxIter; iter += 2)
    {
        zre2 = zre * zre - zim * zim + cre, zim2 = 2 * zre * zim + cim;
        if (zre2 * zre2 + zim2 * zim2 > 4.0f)
            return iter;

        zre = zre2 * zre2 - zim2 * zim2 + cre, zim = 2 * zre2 * zim2 + cim;
        if (zre * zre + zim * zim > 4.0f)
            return iter;
    }
    return maxIter + 1 + (maxIter % 2);
}

int main()
{
    const int maxIter = 128;
    const float cre = -0.8350, cim = -0.2321, h = 4.0f / N;
    int *image = (int *)malloc(sizeof(int) * N * N);   

#pragma acc data copyout(image[0:N * N])
    {
        const int numblock = 4;                         // 指定分块数量
        for (int block = 0; block < numblock; block++)  // 每次计算一块
        {
            const int start = block * (N / numblock), end = start + N / numblock;   // 每块的始末下标
#pragma acc kernels loop independent
            for (int i = start; i < end; i++)
            {
                for (int j = 0; j < N; j++)
                    image[i * N + j] = julia(cre, cim, i * h - 2.0f, j * h - 2.0f, maxIter);
            }
        }
    }
    /*
    FILE *pf = fopen("R:/output.txt", "w");
    for (int i = 0; i < N; i++)
    {
        for (int j = 0; j < N; j++)
            fprintf(pf, "%d ", image[i * N + j]);
        fprintf(pf, "
");
    }
    fclose(pf);
    */    
    free(image);
    //getchar();
    return 0;
}

● 输出结果

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc -acc -Minfo main.c -o main_acc.exe
julia:
      9, Generating acc routine seq
         Generating Tesla code
     11, FMA (fused multiply-add) instruction(s) generated
     17, FMA (fused multiply-add) instruction(s) generated
main:
     30, Generating copyout(image[:67108864])
     37, Loop is parallelizable
         FMA (fused multiply-add) instruction(s) generated
     39, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         37, #pragma acc loop gang, vector(4) /* blockIdx.y threadIdx.y */
         39, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
     39, FMA (fused multiply-add) instruction(s) generated

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main 
=39 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
PGI: "acc_shutdown" not detected, performance results might be incomplete.
 Please add the call "acc_shutdown(acc_device_nvidia)" to the end of your application to ensure that the performance results are complete.

Accelerator Kernel Timing data
D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c
  main  NVIDIA  devicenum=0
    time(us): 20,456
    30: data region reached 2 times
        54: data copyout transfers: 17
             device time(us): total=20,456 max=1,297 min=11 avg=1,203
    36: compute region reached 4 times
        39: kernel launched 4 times
            grid: [256x128]  block: [32x4]
             device time(us): total=0 max=0 min=0 avg=0

● 改进 3，分块传输，没有明显性能提升，为异步做准备

#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>

#define N   (1024 * 8)

#pragma acc routine seq
int julia(const float cre, const float cim, float zre, float zim, const int maxIter)
{
    float zre2 = 0.0f, zim2 = 0.0f;
    for (int iter = 1; iter < maxIter; iter += 2)
    {
        zre2 = zre * zre - zim * zim + cre, zim2 = 2 * zre * zim + cim;
        if (zre2 * zre2 + zim2 * zim2 > 4.0f)
            return iter;

        zre = zre2 * zre2 - zim2 * zim2 + cre, zim = 2 * zre2 * zim2 + cim;
        if (zre * zre + zim * zim > 4.0f)
            return iter;
    }
    return maxIter + 1 + (maxIter % 2);
}

int main()
{
    const int maxIter = 128;
    const float cre = -0.8350, cim = -0.2321, h = 4.0f / N;
    int *image = (int *)malloc(sizeof(int) * N * N);
    
#pragma acc data create(image[0:N * N])                         // 改成 create，不需要从主机拷贝初始数据
    {
        const int numBlock = 4, blockSize = N * N / numBlock;   // 仍然分块计算
        for (int block = 0; block < numBlock; block++)
        {
            const int start = block * (N / numBlock), end = start + N / numBlock;
#pragma acc kernels loop independent
            for (int i = start; i < end; i++)
            {
                for (int j = 0; j < N; j++)
                    image[i * N + j] = julia(cre, cim, i * h - 2.0f, j * h - 2.0f, maxIter);
            }
#pragma acc update host(image[block * blockSize : blockSize])   // 每计算完一块就向主机回传数据
        }
    }
    /*
    FILE *pf = fopen("R:/output.txt", "w");
    for (int i = 0; i < N; i++)
    {
        for (int j = 0; j < N; j++)
            fprintf(pf, "%d ", image[i * N + j]);
        fprintf(pf, "
");
    }
    fclose(pf);
    */    
    free(image);
    //getchar();
    return 0;
}

● 输出结果

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc -acc -Minfo main.c -o main_acc.exe
julia:
      9, Generating acc routine seq
         Generating Tesla code
     11, FMA (fused multiply-add) instruction(s) generated
     17, FMA (fused multiply-add) instruction(s) generated
main:
     30, Generating create(image[:67108864])
     37, Loop is parallelizable
         FMA (fused multiply-add) instruction(s) generated
     39, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         37, #pragma acc loop gang, vector(4) /* blockIdx.y threadIdx.y */
         39, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
     39, FMA (fused multiply-add) instruction(s) generated
     43, Generating update self(image[block*blockSize:blockSize])

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
PGI: "acc_shutdown" not detected, performance results might be incomplete.
 Please add the call "acc_shutdown(acc_device_nvidia)" to the end of your application to ensure that the performance results are complete.

Accelerator Kernel Timing data
D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c
  main  NVIDIA  devicenum=0
    time(us): 20,474
    30: data region reached 2 times
    36: compute region reached 4 times
        39: kernel launched 4 times
            grid: [256x128]  block: [32x4]
            elapsed time(us): total=16,000 max=16,000 min=0 avg=4,000
    43: update directive reached 4 times
        43: data copyout transfers: 20
             device time(us): total=20,474 max=1,287 min=5 avg=1,023

● 改进 4，异步计算 - 双向传输

#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>

#define N   (1024 * 8)

#pragma acc routine seq
int julia(const float cre, const float cim, float zre, float zim, const int maxIter)
{
    float zre2 = 0.0f, zim2 = 0.0f;
    for (int iter = 1; iter < maxIter; iter += 2)
    {
        zre2 = zre * zre - zim * zim + cre, zim2 = 2 * zre * zim + cim;
        if (zre2 * zre2 + zim2 * zim2 > 4.0f)
            return iter;

        zre = zre2 * zre2 - zim2 * zim2 + cre, zim = 2 * zre2 * zim2 + cim;
        if (zre * zre + zim * zim > 4.0f)
            return iter;
    }
    return maxIter + 1 + (maxIter % 2);
}

int main()
{
    const int maxIter = 128;
    const float cre = -0.8350, cim = -0.2321, h = 4.0f / N;
    int *image = (int *)malloc(sizeof(int) * N * N);    

#pragma acc data create(image[0:N * N])
    {
        const int numBlock = 4, blockSize = N / numBlock * N;
        for (int block = 0; block < numBlock; block++)
        {
            const int start = block * (N / numBlock), end = start + N / numBlock;
#pragma acc kernels loop independent async(block + 1)                           // 异步计算，用块编号作标记
            for (int i = start; i < end; i++)
            {
                for (int j = 0; j < N; j++)
                    image[i * N + j] = julia(cre, cim, i * h - 2.0f, j * h - 2.0f, maxIter);
            }
#pragma acc update host(image[block * blockSize : blockSize]) async(block + 1)  // 计算完一块就异步传输
        }
#pragma acc wait
    }
    /*
    FILE *pf = fopen("R:/output.txt", "w");
    for (int i = 0; i < N; i++)
    {
        for (int j = 0; j < N; j++)
            fprintf(pf, "%d ", image[i * N + j]);
        fprintf(pf, "
");
    }
    fclose(pf);
    */
    free(image);
    //getchar();
    return 0;
}

● 输出结果

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc -acc -Minfo main.c -o main_acc.exe
julia:
      9, Generating acc routine seq
         Generating Tesla code
     11, FMA (fused multiply-add) instruction(s) generated
     17, FMA (fused multiply-add) instruction(s) generated
main:
     30, Generating create(image[:67108864])
     37, Loop is parallelizable
         FMA (fused multiply-add) instruction(s) generated
     39, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         37, #pragma acc loop gang, vector(4) /* blockIdx.y threadIdx.y */
         39, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
     39, FMA (fused multiply-add) instruction(s) generated
     43, Generating update self(image[block*blockSize:blockSize])

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 queue=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 queue=2 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 queue=3 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=39 device=0 threadid=1 queue=4 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
PGI: "acc_shutdown" not detected, performance results might be incomplete.
 Please add the call "acc_shutdown(acc_device_nvidia)" to the end of your application to ensure that the performance results are complete.

Accelerator Kernel Timing data
    Timing may be affected by asynchronous behavior
    set PGI_ACC_SYNCHRONOUS to 1 to disable async() clauses
D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c
  main  NVIDIA  devicenum=0
    time(us): 20,503
    30: data region reached 2 times
    36: compute region reached 4 times
        39: kernel launched 4 times
            grid: [256x128]  block: [32x4]
             device time(us): total=0 max=0 min=0 avg=0
    43: update directive reached 4 times
        43: data copyout transfers: 20
             device time(us): total=20,503 max=1,297 min=5 avg=1,025

● 使用统一内存访址（Ubuntu，win64 不支持）

● 改进 4，多设备版本 1，使用 OpenMP

#include <stdio.h>
#include <stdlib.h>
#include <omp.h>
#include <openacc.h>

#define N   (1024 * 8)

#pragma acc routine seq
int julia(const float cre, const float cim, float zre, float zim, const int maxIter)
{
    float zre2 = 0.0f, zim2 = 0.0f;
    for (int iter = 1; iter < maxIter; iter += 2)
    {
        zre2 = zre * zre - zim * zim + cre, zim2 = 2 * zre * zim + cim;
        if (zre2 * zre2 + zim2 * zim2 > 4.0f)
            return iter;

        zre = zre2 * zre2 - zim2 * zim2 + cre, zim = 2 * zre2 * zim2 + cim;
        if (zre * zre + zim * zim > 4.0f)
            return iter;
    }
    return maxIter + 1 + (maxIter % 2);
}

int main()
{
    const int maxIter = 128;
    const int numBlock = acc_get_num_devices(acc_device_nvidia), blockSize = N / numBlock * N;  // 使用 OpenMP 检测目标设备数量，以此作为分块数
    const float cre = -0.8350, cim = -0.2321, h = 4.0f / N;

    int *image = (int *)malloc(sizeof(int) * N * N);
    acc_init(acc_device_nvidia);                                    // 一次性初始化全部目标设备

#pragma omp parallel num_threads(numBlock)                          // 使用多个线程，分别向目标设备发送任务
    {
        acc_set_device_num(omp_get_thread_num(), acc_device_nvidia);// 标记目标设备
#pragma omp for
        for (int block = 0; block < numBlock; block++)
        {
            const int start = block * (N / numBlock), end = start + N / numBlock;
#pragma acc data copyout(image[block * blockSize : blockSize])
            {
#pragma acc kernels loop independent
                for (int i = start; i < end; i++)
                {
                    for (int j = 0; j < N; j++)
                        image[i * N + j] = julia(cre, cim, i * h - 2.0f, j * h - 2.0f, maxIter);
                }
            }
        }
    }
    /*
    FILE *pf = fopen("R:/output.txt", "w");
    for (int i = 0; i < N; i++)
    {
        for (int j = 0; j < N; j++)
            fprintf(pf, "%d ", image[i * N + j]);
        fprintf(pf, "
");
    }
    fclose(pf);
    */
    free(image);
    //getchar();
    return 0;
}

● 输出结果

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc -acc -mp -Minfo main.c -o main_acc.exe
julia:
     10, Generating acc routine seq
         Generating Tesla code
     12, FMA (fused multiply-add) instruction(s) generated
     18, FMA (fused multiply-add) instruction(s) generated
main:
     35, Parallel region activated
     38, Parallel loop activated with static block schedule
     41, Generating copyout(image[block*blockSize:blockSize])
     44, Loop is parallelizable
         FMA (fused multiply-add) instruction(s) generated
     46, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         44, #pragma acc loop gang, vector(4) /* blockIdx.y threadIdx.y */
         46, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
     46, FMA (fused multiply-add) instruction(s) generated
     51, Barrier
     62, Parallel region terminated

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=46 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
PGI: "acc_shutdown" not detected, performance results might be incomplete.
 Please add the call "acc_shutdown(acc_device_nvidia)" to the end of your application to ensure that the performance results are complete.

Accelerator Kernel Timing data
D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c
  main  NVIDIA  devicenum=0
    time(us): 20,462
    41: data region reached 2 times
        50: data copyout transfers: 17
             device time(us): total=20,462 max=1,297 min=5 avg=1,203
    43: compute region reached 1 time
        46: kernel launched 1 time
            grid: [256x128]  block: [32x4]
             device time(us): total=0 max=0 min=0 avg=0

 ● 改进 5，多设备版本 2，调整 OpenMP

#include <stdio.h>
#include <stdlib.h>
#include <omp.h>
#include <openacc.h>

#define N   (1024 * 8)

#pragma acc routine seq
int julia(const float cre, const float cim, float zre, float zim, const int maxIter)
{
    float zre2 = 0.0f, zim2 = 0.0f;
    for (int iter = 1; iter < maxIter; iter += 2)
    {
        zre2 = zre * zre - zim * zim + cre, zim2 = 2 * zre * zim + cim;
        if (zre2 * zre2 + zim2 * zim2 > 4.0f)
            return iter;

        zre = zre2 * zre2 - zim2 * zim2 + cre, zim = 2 * zre2 * zim2 + cim;
        if (zre * zre + zim * zim > 4.0f)
            return iter;
    }
    return maxIter + 1 + (maxIter % 2);
}

int main()
{
    const int maxIter = 128;
    const int numBlock = acc_get_num_devices(acc_device_nvidia), blockSize = N / numBlock * N;
    const float cre = -0.8350, cim = -0.2321, h = 4.0f / N;   

    int *image = (int *)malloc(sizeof(int) * N * N);
    acc_init(acc_device_nvidia);

#pragma omp parallel for num_threads(numBlock)          // 把函数 acc_set_device_num 单独放在一起
    for(int block = 0;block<numBlock;block++)
        acc_set_device_num(block, acc_device_nvidia);

#pragma omp for num_threads(numBlock)
    for (int block = 0; block < numBlock; block++)
    {
        const int start = block * (N / numBlock), end = start + N / numBlock;
#pragma acc data copyout(image[block * blockSize : blockSize])
        {
#pragma acc kernels loop independent
            for (int i = start; i < end; i++)
            {
                for (int j = 0; j < N; j++)
                    image[i * N + j] = julia(cre, cim, i * h - 2.0f, j * h - 2.0f, maxIter);
            }
        }
    }
    /*
    FILE *pf = fopen("R:/output.txt", "w");
    for (int i = 0; i < N; i++)
    {
        for (int j = 0; j < N; j++)
            fprintf(pf, "%d ", image[i * N + j]);
        fprintf(pf, "
");
    }    
    fclose(pf);
    */
    free(image);
    //getchar();
    return 0;
}

● 输出结果

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc main.c -acc -Minfo -o main_acc.exe
julia:
     10, Generating acc routine seq
         Generating Tesla code
     12, FMA (fused multiply-add) instruction(s) generated
     18, FMA (fused multiply-add) instruction(s) generated
main:
     42, Generating copyout(image[block*blockSize:blockSize])
     45, Loop is parallelizable
         FMA (fused multiply-add) instruction(s) generated
     47, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         45, #pragma acc loop gang, vector(4) /* blockIdx.y threadIdx.y */
         47, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
     47, FMA (fused multiply-add) instruction(s) generated

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=47 device=0 threadid=1 num_gangs=32768 num_workers=4 vector_length=32 grid=256x128 block=32x4
PGI: "acc_shutdown" not detected, performance results might be incomplete.
 Please add the call "acc_shutdown(acc_device_nvidia)" to the end of your application to ensure that the performance results are complete.

Accelerator Kernel Timing data
D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c
  main  NVIDIA  devicenum=0
    time(us): 20,571
    42: data region reached 2 times
        51: data copyout transfers: 17
             device time(us): total=20,571 max=1,336 min=11 avg=1,210
    44: compute region reached 1 time
        47: kernel launched 1 time
            grid: [256x128]  block: [32x4]
            elapsed time(us): total=3,000 max=3,000 min=3,000 avg=3,000