OpenACC vs C++：致命的なエラー：デバイスに変数が部分的に存在します

Question

OpenACCを使用してC++アプリケーションをGPUに移植しようとしています。 が期待するように、C++コードには多くのカプセル化と抽象化があります。メモリは で、ベクターのようなクラスに割り当てられています。このクラスは、アプリケーションの周りの他の多くのクラスで再利用されます。そして、私は正しく問題を解決しようとしています OpenACCプラグマをコードに挿入します。ここで私は は働いているコードの簡単な例です：

#define DATASIZE 16 

class Data { 
    float *arr; 
public: 
    Data() {arr = new float[DATASIZE];} 
    ~Data() { delete [] arr; } 
    float &get(int i) { return arr[i]; } 
}; 

class DataKeeper { 
    Data a, b, c; 
public: 
    void init() { 
    for (int i = 0; i < DATASIZE; ++i) 
     a.get(i) = 0.0; 
    } 
}; 

int main() { 
    DataKeeper DK; 
    DK.init(); 
} 

を私はデバイスに必要なデータを送信し、このようなコードで を終了するには、いくつかのOpenACCプラグマを挿入します。

#define DATASIZE 16 

class Data { 
    float *arr; 

public: 
    Data() { 
    arr = new float[DATASIZE]; 
#pragma acc enter data copyin(this) 
#pragma acc enter data create(arr[:DATASIZE]) 
    } 

    ~Data() { 
#pragma acc exit data delete(arr) 
#pragma acc exit data delete(this) 
    delete [] arr; 
    } 

    float &get(int i) { return arr[i]; } 
}; 

class DataKeeper { 
    Data a, b, c; 

public: 
    DataKeeper() { 
#pragma acc enter data copyin(this) 
    } 

    ~DataKeeper() { 
#pragma acc exit data delete(this) 
    } 

    void init() { 
#pragma acc parallel loop 
    for (int i = 0; i < DATASIZE; ++i) { 
     a.get(i) = 0.0; 
    } 
    } 
}; 

int main() { 
    DataKeeper DK; 
    DK.init(); 
} 

コンパイルして実行すると、次のエラーが表示されます。

$ pgc++ test.cc -acc -g 

$ ./a.out 
_T24395416_101 lives at 0x7fff49e03070 size 24 partially present 
Present table dump for device[1]: NVIDIA Tesla GPU 0, compute capability 3.5, threadid=1 
host:0x1ae6eb0 device:0xc05ca0200 size:64 presentcount:0+1 line:11 name:(null) 
host:0x1f33620 device:0xc05ca0600 size:64 presentcount:0+1 line:11 name:(null) 
host:0x1f33d10 device:0xc05ca0a00 size:64 presentcount:0+1 line:11 name:(null) 
host:0x7fff49e03070 device:0xc05ca0000 size:8 presentcount:0+1 line:11 name:_T24395600_98 
host:0x7fff49e03078 device:0xc05ca0400 size:8 presentcount:0+1 line:11 name:_T24395600_98 
host:0x7fff49e03080 device:0xc05ca0800 size:8 presentcount:0+1 line:11 name:_T24395600_98 
allocated block device:0xc05ca0000 size:512 thread:1 
allocated block device:0xc05ca0200 size:512 thread:1 
allocated block device:0xc05ca0400 size:512 thread:1 
allocated block device:0xc05ca0600 size:512 thread:1 
allocated block device:0xc05ca0800 size:512 thread:1 
allocated block device:0xc05ca0a00 size:512 thread:1 

FATAL ERROR: variable in data clause is partially present on the device: name=_T24395416_101 
file:/home/bozhenovn/tst/test.cc _ZN10DataKeeperC1Ev line:27 

コードに何が間違っているかわかりません。コードを修正する方法や問題をさらに詳しく調べる方法を提案していただければ幸いです。ありがとうございました！

Answer 1

"a"のホストアドレスは "DK"の開始アドレスと同じです。したがって、コンパイラが変数のホストアドレスをデバイスアドレスにマッピングするために使用する現在のテーブル内のホストアドレスをルックアップすると、サイズが異なることがわかります。 「a」はサイズ8、「DK」はサイズ24です。

ここでは修正を表示しますが、ここで何が起こっているのかを理解してみましょう。ホスト上に「DK」が作成されると、最初にそのデータメンバーのそれぞれのストレージが作成され、各データメンバーのクラスコンストラクターが呼び出されます。それはそれ自身のコンストラクタを実行します。したがって、各データメンバーに対して、コードはデバイス上にこのポインタのクラスを作成し、デバイス上に「arr」配列を割り当てます。これが完了すると、デバイス上に各データメンバーのためのスペースを持つ「DK」が作成されます。ただし、 "DK"のデバイスコピーはデータメンバーの後に作成されるため、コンパイラは自動的に2つを関連付けることはできません。

以下、2つの可能な修正を掲載しました。

まず、 "データ"クラスで独自のデータを管理できますが、クラスデータのメンバーを動的に割り当てる必要があります。このようにして、データコンストラクタはDataKeeperコンストラクタの後に発生し、コンパイラはデバイスデータ（「アタッチ」とも呼ばれます）を関連付けることができます。

第2に、DataKeeperクラスでDataクラスのデータを管理することができます。ただし、これはDataのデータを公開することを要求します。

「OpenACCによる並列プログラミング」の第5章「高度なデータ管理」を書き、C++クラスのデータ管理についてのセクションが含まれています。私のサンプルコードはhttps://github.com/rmfarber/ParallelProgrammingWithOpenACC/tree/master/Chapter05 で見つけることができます。特に、「accList」ジェネリックコンテナクラスのやり方を見てください。

修正＃1：

#define DATASIZE 16 
#include <iostream> 
#ifdef _OPENACC 
#include <openacc.h> 
#endif 

class Data { 
    float *arr; 

public: 
    Data() { 
    arr = new float[DATASIZE]; 
#pragma acc enter data copyin(this) 
#pragma acc enter data create(arr[:DATASIZE]) 
    } 

    ~Data() { 
#pragma acc exit data delete(arr) 
#pragma acc exit data delete(this) 
    delete [] arr; 
    } 

    float &get(int i) { return arr[i]; } 
    void updatehost() { 
    #pragma acc update host(arr[0:DATASIZE]) 
    } 

}; 

class DataKeeper { 
    Data *a, *b, *c; 

public: 
    DataKeeper() { 
#pragma acc enter data copyin(this) 
    a = new Data; 
    b = new Data; 
    c = new Data; 
    } 

    ~DataKeeper() { 
#pragma acc exit data delete(this) 
    delete a; 
    delete b; 
    delete c; 
    } 

    void init() { 
#pragma acc parallel loop present(a,b,c) 
    for (int i = 0; i < DATASIZE; ++i) { 
     a->get(i) = i; 
    } 
    a->updatehost(); 
    std::cout << "a.arr[0]=" << a->get(0) << std::endl; 
    std::cout << "a.arr[end]=" << a->get(DATASIZE-1) << std::endl; 
    } 
}; 

int main() { 
    DataKeeper DK; 
    DK.init(); 
} 

修正＃2

#define DATASIZE 16 
#include <iostream> 
#ifdef _OPENACC 
#include <openacc.h> 
#endif 

class Data { 
public: 
    float *arr; 

    Data() { 
    arr = new float[DATASIZE]; 
    } 

    ~Data() { 
    delete [] arr; 
    } 

    float &get(int i) { return arr[i]; } 
}; 

class DataKeeper { 
    Data a, b, c; 

public: 
    DataKeeper() { 
#pragma acc enter data copyin(this) 
#pragma acc enter data create(a.arr[0:DATASIZE]) 
#pragma acc enter data create(b.arr[0:DATASIZE]) 
#pragma acc enter data create(c.arr[0:DATASIZE]) 
    } 

    ~DataKeeper() { 
#pragma acc exit data delete(this) 
#pragma acc exit data delete(a.arr) 
#pragma acc exit data delete(b.arr) 
#pragma acc exit data delete(c.arr) 
    } 

    void init() { 
#pragma acc parallel loop present(a,b,c) 
    for (int i = 0; i < DATASIZE; ++i) { 
     a.get(i) = i; 
    } 
#pragma acc update host(a.arr[0:DATASIZE]) 
    std::cout << "a.arr[0]=" << a.arr[0] << std::endl; 
    std::cout << "a.arr[end]=" << a.arr[DATASIZE-1] << std::endl; 
    } 
}; 

int main() { 
    DataKeeper DK; 
    DK.init(); 
}