OpenACCとGNU Scientific Library - gsl_matrixのデータ移動

Question

記録されたopenaccの概要のコースビデオを、データの動きを表現するための講義3まで見てきました。 copy_in()を使用して、どのようにgsl_matrix*をCPUからGPUに移動しますか？例えば、CPU上では、

gsl_matrix *Z = gsl_matrix_calloc(100, 100), 

のような何かをすることができます。これは私に100x100の0の行列を与えます。今、Zは、私がcopyin()を使用してGPUにCPUから（ポインタである）Zのデータ移動を表現する方法を

typedef struct{ 
    size_t size1; 
    size_t size2; 
    size_t tda; 
    double * data; 
    gsl_block * block; 
    int owner; 
    } gsl_matrix; 

、のように見えるgsl_matrix構造体へのポインタですか？

Answer 1

私はOpenACCデータと計算領域内でGSLを使用することについて直接話すことはできませんが、ダイナミックデータメンバーを使用して集約タイプに関する一般的な回答を得ることができます。

PGIコンパイラと新しいNVIDIAデバイスを使用していると仮定すると、まずCUDA Unified Memory（UVM）が必要です。フラグ "-ta = tesla：managed"でコンパイルすると、動的に割り当てられたすべてのデータがCUDAランタイムによって管理されるため、データの移動を自分で管理する必要はありません。オーバーヘッドがあり、警告がありますが、作業を開始しやすくなります。 PGI 16.9以降に同梱されているCUDA 8.0は、UVMの性能を向上させることに注意してください。

UVMを使用しない場合、データの手動ディープコピーを実行する必要があります。以下は、最初にデバイス上に親構造を作成し、浅いコピーを実行する基本的な考え方です。次に、デバイス上の "data"という動的配列を作成し、初期値を配列にコピーしてから、デバイス構造のデータポインタにデータのデバイスポインタを付加します。 「ブロック」自体は動的データメンバーを持つ構造体の配列なので、配列をループしてデバイス上にデータ配列を作成する必要があります。

matrix * mat = (matrix*) malloc(sizeof(matrix)); 
    #pragma acc enter data copyin(mat) 
// Change this to the correct size of "data" and blocks 
    #pragma acc enter data copyin(mat.data[0:dataSize]); 
    #pragma acc enter data copyin(mat.block[0:blockSize]); 
    for (i=0; i < blockSize; ++i) { 
     #pragma acc enter data copyin(mat.block[i].data[0:mat.block[i].size]) 
    } 

は、削除ボトムアップ

for (i=0; i < blockSize; ++i) { 
    #pragma acc exit data delete(mat.block[i].data) 
} 
#pragma acc exit data delete(mat.block); 
#pragma acc exit data delete(mat.data); 
#pragma acc exit data delete(mat); 

アップデート

から削除、再構造を歩く、スカラーのみまたは基本的なデータ型の配列を必ず更新してください。すなわち、「データ」を更新するが「ブロック」は更新しない。 Updateは浅いコピーを行い、 "ブロック"を更新することで不正なアドレスにつながるホストまたはデバイスのポインタが更新されます。

最後に、行列変数を計算領域で使用する場合は、必ず「present」節に入れてください。

#pragma acc parallel loop present(mat)