OpenACC 计算构建内的自定义函数

▶ 使用 routine 构件创建的自定义函数，在并行调用上的差别

● 代码，自定义一个 sqab 函数，使用内建函数 fabsf 和 sqrtf 计算一个矩阵所有元素绝对值的平方根

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

#define ROW 8
#define COL 64

#pragma acc routine vector
void sqab(float *a, const int m)
{    
#pragma acc loop
    for (int idx = 0; idx < m; idx++)
        a[idx] = sqrtf(fabsf(a[idx]));    
}

int main()
{
    float x[ROW][COL];
    int row, col;
    for (row = 0; row < ROW; row++)
    {
        for (col = 0; col < COL; col++)
            x[row][col] = row * 10 + col;
    }
    printf("
x[1][1] = %f
", x[1][1]);

#pragma acc parallel loop vector pcopy(x[0:ROW][0:COL]) // 之后在这里分别添加 gang，worker，vector
    for (row = 0; row < ROW; row++)
        sqab(&x[row][0], COL);
    printf("
x[1][1] = %f
", x[1][1]);

    //getchar();
    return 0;
}

 ● 输出结果，第 28 行不添加并行级别子句（默认使用 gang）

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc main.c -acc -Minfo -o main_acc.exe
sqab:
     11, Generating Tesla code
         13, #pragma acc loop vector /* threadIdx.x */
     13, Loop is parallelizable
main:
     28, Generating copy(x[:][:])
         Accelerator kernel generated
         Generating Tesla code
         29, #pragma acc loop gang /* blockIdx.x */

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe

x[1][1] = 11.000000
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=28 device=0 threadid=1 num_gangs=8 num_workers=1 vector_length=32 grid=8 block=32      // 8 个 gang 在 blockIdx.x 层级，1 个 worker，vector 在 threadIdx.x 层级

x[1][1] = 3.316625
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): 9
    28: compute region reached 1 time
        28: kernel launched 1 time
            grid: [8]  block: [32]
            elapsed time(us): total=1000 max=1000 min=1000 avg=1000
    28: data region reached 2 times
        28: data copyin transfers: 1
             device time(us): total=4 max=4 min=4 avg=4
        31: data copyout transfers: 1
             device time(us): total=5 max=5 min=5 avg=5

 ● 输出结果，第 28 行添加并行级别子句 worker

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc main.c -acc -Minfo -o main_acc.exe
sqab:
     11, Generating Tesla code
         13, #pragma acc loop vector /* threadIdx.x */
     13, Loop is parallelizable
main:
     28, Generating copy(x[:][:])
         Accelerator kernel generated
         Generating Tesla code
         29, #pragma acc loop worker(4) /* threadIdx.y */
     29, Loop is parallelizable

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe

x[1][1] = 11.000000
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main
line=28 device=0 threadid=1 num_gangs=1 num_workers=4 vector_length=32 grid=1 block=32x4    // 1 个 gang，4 个 worker 在 threadIdx.y 层级，使用 2 维线程网格

x[1][1] = 3.316625
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): 10
    28: compute region reached 1 time
        28: kernel launched 1 time
            grid: [1]  block: [32x4]
             device time(us): total=0 max=0 min=0 avg=0
    28: data region reached 2 times
        28: data copyin transfers: 1
             device time(us): total=5 max=5 min=5 avg=5
        31: data copyout transfers: 1
             device time(us): total=5 max=5 min=5 avg=5

 ● 输出结果，第 28 行添加并行级别子句 vector

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc main.c -acc -Minfo -o main_acc.exe
sqab:
     11, Generating Tesla code
         13, #pragma acc loop vector /* threadIdx.x */
     13, Loop is parallelizable
main:
     28, Generating copy(x[:][:])
         Accelerator kernel generated
         Generating Tesla code
         29, #pragma acc loop seq
     29, Loop is parallelizable

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe

x[1][1] = 11.000000
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main 
line=28 device=0 threadid=1 num_gangs=1 num_workers=1 vector_length=32 grid=1 block=32      // 1 个 gang，1 个 worker，并行全都堆在 threadIdx.x 层级上

x[1][1] = 3.316625
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): 10
    28: compute region reached 1 time
        28: kernel launched 1 time
            grid: [1]  block: [32]
            elapsed time(us): total=1000 max=1000 min=1000 avg=1000
    28: data region reached 2 times
        28: data copyin transfers: 1
             device time(us): total=5 max=5 min=5 avg=5
        31: data copyout transfers: 1
             device time(us): total=5 max=5 min=5 avg=5

● 如果自定义函数并行子句等级高于主调函数，则主调函数并行子句会变成 seq；如果自定义函数并行子句等级低于内部并行子句等级，则会报 warning，忽略掉内部并行子句：

#pragma acc routine vector
void sqab(float *a, const int m)
{    
#pragma acc loop worker
    for (int idx = 0; idx < m; idx++)
        a[idx] = sqrtf(fabsf(a[idx]));    
}

● 编译结果（运行结果通上面的 worker，不写）

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc main.c -acc -Minfo -o main_acc.exe
PGC-W-0155-acc loop worker clause ignored in acc routine vector procedure  (main.c: 13)
sqab:
     11, Generating Tesla code
         13, #pragma acc loop vector /* threadIdx.x */
     13, Loop is parallelizable