OpenACC的對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

這是怎麼回事是主機地址「一」是一樣的「DK」的起始地址。因此，當編譯器在當前表中查找主機地址時，它使用該主機地址將變量的主機地址映射到設備地址，從而發現大小不同。 「一」是大小爲8，而「DK」的大小爲24

我會告訴修復之下，但讓我們反向跟蹤和了解這裏發生了什麼。當在主機上創建「DK」時，它首先爲其每個數據成員創建存儲，然後調用每個數據成員的類構造函數。然後執行它自己的構造函數。因此，對於每個數據成員，您的代碼將在設備上創建類指針，然後在設備上分配「arr」數組。完成此操作後，將在設備上爲每個數據成員創建「DK」空間。但是，由於在數據成員之後創建了「DK」的設備副本，因此編譯器無法自動將兩者關聯。

下面，我已經發布了兩個可能的修復。

首先，你可以有「數據」類管理它自己的數據，但你需要動態分配類的數據成員。這樣Data構造函數將在DataKeeper構造函數之後出現，以便編譯器可以關聯設備數據（也稱爲「attach」）。

其次，你可以有DataKeeper的班級管理數據類的數據。但是，這將會要求Data的數據公開。

注意，我寫的第5章「高級數據管理」一書的「並行編程與OpenACC的」，其中包括一個關於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(); 
}