バージョン 15 と 16 のコンパイラーのユーザーズガイドには、Offload を使用する場合の制限に、
変数をターゲットで使用できるようにするには、オフロードコード内から呼び出される関数で参照されるグローバル変数が、一致するターゲット属性で宣言されている必要があります。オフロードコードはホスト CPU のグローバル変数にアクセスできません。これはコンパイラーによる制約です。というのがあります。
デフォルトでは、ポインター変数は対応する型の単一要素を指すと仮定されます。オフロードコードはポインターを逆参照して単一要素にアクセスできます。ポインターが指すデータ要素はターゲットメモリーおよび調整されたポインター値に自動的にコピーされます。element-count-expr 式と in / out / inout パラメーターを使用して、可変長データをコピーできます。
#include <omp.h>
#include <time.h>
static long num_steps = 1000000000;
double step;
double start, end;
int main (){
int i;
double pi, x, sum = 0.0;
step = 1.0/(double) num_steps;
start = clock();
#pragma offload target(mic) in(x, num_steps, step) out(sum)
#pragma omp parallel for simd reduction(+:sum) private(x)
for (i=0;i< num_steps; i++){
x = (i+0.5) * step;
sum = sum + 4.0/(1.0+x*x);
}
pi = step * sum;
end = clock();
printf("Pi = %f Time = %f\n", pi, (double)((end-start)/CLOCKS_PER_SEC));
}
上記のソースを、V15 のコンパイラーでコンパイルすると、以下のエラーになります。
pi1_err.c(18): エラー: オフロード領域で使用される 変数 "num_steps" には互換性のある "target" 属性がなければなりません。
#pragma offload target(mic) optional in(x, num_steps, step) out(sum)
^
pi1_err.c(18): エラー: オフロード領域で使用される 変数 "step" には互換性のある "target" 属性がなければなりません。
#pragma offload target(mic) optional in(x, num_steps, step) out(sum)
^
しかし、V16 (16.0.0.079) でコンパイルすると、コンパイルおよび実行できてしまいます。V16 のマニュアル表記もほぼ同じですが、これはマニュアルのエラーのようです。
インテル® C/C++ コンパイラー V16 は、OpenMP* 4.1 のサポートに向けてオフロード構文のデータスコープにけるグローバルデータの扱いが変わりました。これに伴い、オフロード向けのインテル® コンパイラー言語拡張の振る舞いも変更されたようです。オフロードプラグマの構文内で参照されるグローバルデータは、ターゲット属性と一致する宣言を行う必要がなく、関連する変数の値はデフォルトでコプロセッサー環境に転送 (送受信) されます。
V15 向けには正式には次のように記述する必要があります。V16 は以下でも OK;
#include <omp.h>
#include <time.h>
__declspec(target(mic)) static long num_steps = 1000000000;
__declspec(target(mic)) double step;
もしくは
#pragma offload_attribute(push, target(mic))
static long num_steps = 1000000000;
double step;
#pragma offload_attribute(pop)
double start, end;
int main ()
{
int i;
double pi, x, sum = 0.0;
step = 1.0/(double) num_steps;
start = clock();
#pragma offload target(mic) in(x, num_steps, step) out(sum)
#pragma omp parallel for simd reduction(+:sum) private(x)
for (i=0;i< num_steps; i++){
x = (i+0.5)*step;
sum = sum + 4.0/(1.0+x*x);
}
pi = step * sum;
end = clock();
printf("Pi = %f Time = %f\n", pi, (double)((end-start)/CLOCKS_PER_SEC));
}
V16 のコンパイラーを使用して開発したオフロード構文を含むソースを、V15 のコンパイラーでコンパイルした場合、前述のようなエラーが発生する可能性が非常に高いと思われます。複数のバージョンのコンパイラーを使用して、オフロードを行うアプリケーションを開発する場合、V15 の旧方式に従ってソースを作成されることをお勧めします。
