【问题标题】:Nested data environment with different subparts of the same array具有同一数组的不同子部分的嵌套数据环境
【发布时间】:2014-02-05 06:00:56
【问题描述】:

这是我关于 openacc 的问题。 我阅读了 API(v1 和 v2),但我不清楚具有同一数组的不同子部分的嵌套数据环境的行为。

代码示例:

#pragma acc data pcopyin(a[0:20])
{
  #pragma acc data pcopyin(a[100:20])
  {
    #pragma acc parallel loop
    for(i=0; i<20; i++)
      a[i] = i;
      a[i+100] = i;
  }
}

我的理解是这应该可以工作(或至少两个 acc 数据部分):

  • 第一个编译指示检查 a[0,20] 是否在加速器上
  • 否 -> 数据在设备上分配并传输
  • 第二个编译指示检查 a[100,120] 是否在加速器上
  • 指针 a 在加速器上,但不是来自 a[100,120] 的数据
  • 数据在设备上分配和传输

我用 CAPS 编译器(v3.3.0,这是我的测试机器上目前唯一可用的)尝试了这种事情,第二个 pragma acc 数据返回一个错误(我的第二个子数组没有正确的形状)。 所以我的测试(我想)发生的事情是在加速器上找到了指针“a”,但与它相关的形状([0:20])在我的第二个编译指示([100:20] )。

这是 API 中计划的正常行为,还是我的示例应该有效?

此外,如果这应该起作用,同一数组的子部分之间是否存在某种连贯性(不知何故,它们将像在主机上一样定位,我将能够放置 a[i] += a [100+i] 在我的内核中)?

【问题讨论】:

    标签: openacc


    【解决方案1】:

    当前测试将查看设备上是否有“a”。因此,当遇到第二个数据区域时,“a”已经在设备上,但只是部分存在。相反,更好的方法是添加一个指向“a”的指针并在设备上引用该指针。比如:

    #include <stdio.h>
    
    int main () {
    
       int a[200];
       int *b;
       int i;
       for(i=0; i<200; i++) a[i] = 0;
       b=a+100;
    
    #pragma acc data pcopy(a[0:20])
    {
      #pragma acc data pcopy(b[0:20])
      {
        #pragma acc parallel loop
        for(i=0; i<20; i++) {
          a[i] = i;
          b[i] = i;
        }
      }
    }
       for(i=0; i<22; i++) printf("%d = %d \n", i, a[i]);
       for(i=100; i<122; i++) printf("%d = %d \n", i, a[i]);
      return 0;
     }
    

    如果您刚刚复制了“a[100:20]”,那么超出此范围的访问将被视为程序员错误。

    希望这会有所帮助, 垫子

    【讨论】:

    • 好的,我昨天测试了你的方法,结果正确。但我想知道我的示例是否应该根据 API 工作。根据 API,观察到的行为是正确的,还是仅仅是 openacc 的 CAPS 编译器的实现限制?
    • 就 API 而言,代码应该可以工作。但是,每个编译器都可能以不同的方式实现这一点。规范中唯一的保证是,当内核在加速器上运行时,两个子数组都可用。
    猜你喜欢
    • 2017-06-19
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 2021-03-10
    • 2018-11-24
    • 1970-01-01
    • 2018-08-13
    相关资源
    最近更新 更多