【问题标题】:OpenACC: How to select an array on device from a pointer to corresponding array on hostOpenACC:如何从指向主机上相应数组的指针中选择设备上的数组
【发布时间】:2020-04-24 18:37:03
【问题描述】:

我正在尝试使用 OpenACC 将现有的 C 代码卸载到 GPU。在原始 CPU 代码中,很多时候需要根据某个参数的值来选择一个数据数组。下面给出了一个示例 CPU 代码:

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


void selectArray (int **F, int a);

#define NN 1000
int *C, *D, *E;

int main(void)
{
    int *F, a = 10; // a is the parameter used to select the array

    C = (int *)malloc(NN * sizeof(int));
    D = (int *)malloc(NN * sizeof(int));
    E = (int *)malloc(NN * sizeof(int));

    for (int i = 0; i < NN; i++)
    {
        C[i] = 10;
        D[i] = 20;
    }

    selectArray(&F, a);

    for (int i = 0; i < NN; i++)
    {
       E[i] = 2 * F[i];
    }

    for (int i = 0; i < 200; i++)
       printf("%d %d \n", i, E[i]);

    return 0;

}

void selectArray(int **F, int a)
{
    if (a <= 15)
    {
        (*F) = C;
    }
    else
    {
        (*F) = D;
    }
}

对于 OpenACC 版本的代码,数组 C 和 D 已经存在于 GPU 上,需要对根据参数 a 选择的数组进行进一步的计算。

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

void selectArray(int **F, int a);

#define NN 1000
int *C, *D, *E;

int main(void)
{
    int *F, a = 10; // a is the parameter used to select the array

    C = (int *)malloc(NN * sizeof(int));
    D = (int *)malloc(NN * sizeof(int));
    E = (int *)malloc(NN * sizeof(int));

#pragma acc enter data create(C[:NN], D[:NN])
#pragma acc parallel loop present(C[:NN], D[:NN])
    for (int i = 0; i < NN; i++)
    {
        C[i] = 10;
        D[i] = 20;
    }

    selectArray(&F, a);

#pragma acc enter data copyin(F[:1]) create(E[:NN])

// Here, I cannot figure out how to point F to a selected array (C or D) on the device

#pragma acc parallel loop
        for (int i = 0; i < NN; i++)
        {
            E[i] = 2 * F[i]; //further calculations on selected array on GPU
        }
    }

#pragma acc exit data delete (C[:NN], D[:NN], F)copyout(E[:200])
    for (int i = 0; i < 200; i++)
        printf("%d %d \n", i, E[i]);

    return 0;
}

void selectArray(int **F, int a)
{
    if (a <= 15)
    {
        (*F) = C;
    }
    else
    {
        (*F) = D;
    }
}

在实际代码中,数组 C 和 D 是在不同的函数中计算的,而不是在 main 函数中。我曾尝试在互联网上搜索以解决此问题,但找不到任何相关示例。我在 Windows 10 上使用 PGI 19.10 编译器。需要这方面的帮助。 提前致谢

【问题讨论】:

  • 您的输出是什么?除了在设置E 的循环之后的额外} 之外,我得到了我认为正确的答案(所有20 的书面答案,如果我修改a,则所有40 的答案)。我在带有 GPU 目标的 Linux 上使用 19.10。
  • 先生,我在 Windows 编译器上也得到了相同的结果,我应该在问题中提到。但是,在实际代码中,删除“F”会引起问题。 Mat Colgrove 对此问题提供了出色的解释和解决方案。

标签: c arrays openacc pgcc


【解决方案1】:

您只需要在并行循环中添加一个“present(F)”,而不是在数据区域中包含“F”。由于 acc present 表查找是通过主机地址完成的,如果“F”与设备上存在的现有主机地址匹配,则它将“F”与相同的设备地址相关联。但是,不要将“F”放在它自己的数据区域中,特别是不要删除它,因为它会导致同一设备阵列上的多个空闲。

我稍微修改了您的代码,使“F”在一种情况下指向“C”,在第二种情况下指向“D”。

% cat test.c
#include <stdio.h>
#include <stdlib.h>

void selectArray(int **F, int a);

#define NN 1000
int *C, *D, *E;

int main(void)
{
    int *F, a = 10; // a is the parameter used to select the array

    C = (int *)malloc(NN * sizeof(int));
    D = (int *)malloc(NN * sizeof(int));
    E = (int *)malloc(NN * sizeof(int));

#pragma acc enter data create(C[:NN], D[:NN], E[:NN])
#pragma acc parallel loop present(C[:NN], D[:NN])
    for (int i = 0; i < NN; i++)
    {
        C[i] = 10;
        D[i] = 20;
    }

    for (a=10;a<=20;a+=10) {
       selectArray(&F, a);

#pragma acc parallel loop present(E,F)
       for (int i = 0; i < NN; i++)
       {
          E[i] = 2 * F[i]; //further calculations on selected array on GPU
       }
#pragma acc update host(E[:20])
       for (int i = 0; i < 20; i++)
       {
          printf("a=%d E[%d]=%d \n", a, i, E[i]);
       }
     }
#pragma acc exit data delete(C, D, E)
    return 0;
}

void selectArray(int **F, int a)
{
    if (a <= 15)
    {
        (*F) = C;
    }
    else
    {
        (*F) = D;
    }
}
% pgcc -ta=tesla -Minfo=accel test.c; a.out
main:
     17, Generating enter data create(D[:1000],E[:1000],C[:1000])
     18, Generating present(D[:1000],C[:1000])
         Generating Tesla code
         19, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     28, Generating present(E[:],F[:])
         Generating Tesla code
         29, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     34, Generating update self(E[:20])
     39, Generating exit data delete(E[:1],D[:1],C[:1])
a=10 E[0]=20
a=10 E[1]=20
a=10 E[2]=20
a=10 E[3]=20
a=10 E[4]=20
a=10 E[5]=20
a=10 E[6]=20
a=10 E[7]=20
a=10 E[8]=20
a=10 E[9]=20
a=10 E[10]=20
a=10 E[11]=20
a=10 E[12]=20
a=10 E[13]=20
a=10 E[14]=20
a=10 E[15]=20
a=10 E[16]=20
a=10 E[17]=20
a=10 E[18]=20
a=10 E[19]=20
a=20 E[0]=40
a=20 E[1]=40
a=20 E[2]=40
a=20 E[3]=40
a=20 E[4]=40
a=20 E[5]=40
a=20 E[6]=40
a=20 E[7]=40
a=20 E[8]=40
a=20 E[9]=40
a=20 E[10]=40
a=20 E[11]=40
a=20 E[12]=40
a=20 E[13]=40
a=20 E[14]=40
a=20 E[15]=40
a=20 E[16]=40
a=20 E[17]=40
a=20 E[18]=40
a=20 E[19]=40

【讨论】:

    猜你喜欢
    • 2014-11-07
    • 1970-01-01
    • 2012-12-25
    • 2020-08-25
    • 1970-01-01
    • 2021-02-07
    • 2021-11-12
    • 1970-01-01
    • 1970-01-01
    相关资源
    最近更新 更多