インテル® C++ コンパイラー
インテル® Fortran コンパイラー
インテル® oneAPI
C/C++
Fortran
OpenMP

ベータ版インテル® C++ コンパイラーおよびベータ版インテル® Fortran コンパイラー向けの GPU への OpenMP* オフロード導入

同カテゴリーの次の記事

GPU 構成を調査してパフォーマンスを向上する方法

この記事は、インテル® デベロッパー・ゾーンに公開されている「Get Started with OpenMP* Offload to GPU for the Intel® C++ Compiler and Intel® Fortran Compiler (Beta)」の日本語参考訳です。

この記事の PDF 版はこちらからご利用になれます。

インテル® C++ コンパイラーおよびインテル® Fortran コンパイラーの GPU への OpenMP* オフロード機能は、広範囲のアクセラレーター向けに OpenMP* ソースファイルをコンパイルできます。

はじめに

既知の問題と最新情報については、リリースノートをご覧ください。

GPU 向けの OpenMP* 4.5 サブセット

oneAPI 向けインテル® C++ コンパイラー (英語) は、ユーザー定義のリダクション (UDR) を除く、CPU 向けのすべての OpenMP* 4.5 機能をサポートします。GPU とオフロードは、OpenMP* 4.5 および OpenMP* 5.0 の declare variant のサブセットをサポートします。この機能は、ハイパフォーマンス・コンピューティング (HPC) アプリケーションで使用される OpenMP* 構造の解析とコレクションに基づいています。サポートの詳細については、『oneAPI 向けインテル® C++ コンパイラー・デベロッパー・ガイドおよびリファレンス』の「OpenMP* のサポート」 (英語) を参照してください。

OpenMP* 4.5 のサブセットに含まれる GPU ターゲットのデバイスサポート

以下は、GPU および CPU 向けの次世代コンパイラーでサポートされる OpenMP* プラグマです。

  • declare_target
  • declare_simd
  • target
  • target_teams
  • target_teams_distribute
  • target_teams_distribute_simd
  • target_parallel
  • target_parallel_for
  • target_parallel_for_simd
  • target_teams_distribute_parallel_for
  • target_teams_distribute_parallel_for_simd
  • target_variant_dispatch
  • simd
  • master
  • atomic
  • barrier
  • parallel_for
  • teams
  • teams_distribute
  • teams_distribute_parallel_for
  • teams_distribute_parallel_for_simd
  • target_data)
  • target_update
  • target_enter_data
  • target_exit_data
  • parallel
  • single
  • for
  • for_simd
  • parallel_for_simd
  • distribute_parallel_for
  • parallel_sections)
  • sections

以下は、CPU および GPU 向けの次世代コンパイラーでサポートされる OpenMP* 指示句です。

  • if
  • final
  • num_threads
  • safelen
  • simdlen
  • collapse
  • default
  • private
  • firstprivate
  • lastprivate
  • shared
  • reduction
  • linear
  • aligned
  • proc_bind
  • schedule
  • ordered
  • nowait
  • untied
  • mergeable
  • flush
  • read
  • write
  • update
  • capture
  • seq_cst
  • depend
  • device
  • simd
  • map
  • num_teams
  • thread_limit
  • num_tasks
  • ist_schedule
  • defaultmap
  • to
  • from
  • use_device_ptr
  • is_device_ptr

以下は、CPU ホストで利用可能なランタイム・サポート・ルーチンです。

  • EXTERN int omp_get_num_devices(void);
  • EXTERN int omp_get_initial_device(void);
  • EXTERN void *omp_target_alloc(size_t size, int device_num);
  • EXTERN void omp_target_free(void *device_ptr, int device_num);
  • EXTERN int omp_target_is_present(void *ptr, int device_num);
  • EXTERN int omp_target_memcpy(void *dst, void *src, size_t length, size_t dst_offset, size_t src_offset, int dst_device, int src_device);
  • EXTERN int omp_target_memcpy_rect(void *dst, void *src, size_t element_size, int num_dims, const size_t *volume, const size_t *dst_offsets, const size_t *src_offsets, const size_t *dst_dimensions, const size_t *src_dimensions, int dst_device, int src_device);
  • EXTERN int omp_target_associate_ptr(void *host_ptr, void *device_ptr, size_t size, size_t device_offset, int device_num);
  • EXTERN int omp_target_disassociate_ptr(void *host_ptr, int device_num);
  • EXTERN int omp_is_initial_device(void);
  • EXTERN int omp_get_initial_device(void);
  • EXTERN void kmp_global_barrier_init(void); // Intel externsion
  • EXTERN void kmp_global_barrier(void); // Intel externsion
  • EXTERN void omp_set_default_device(int dev_num )
  • EXTERN int omp_get_default_device(void)

以下は GPU 向けのデバイス・ランタイム・ルーチンです。

  • EXTERN int omp_get_team_num(void);
  • EXTERN int omp_get_num_teams(void);
  • EXTERN int omp_get_team_size(int);
  • EXTERN int omp_get_thread_num(void);
  • EXTERN int omp_get_num_threads(void);
  • EXTERN int omp_in_parallel(void);
  • EXTERN int omp_get_max_threads(void);
  • EXTERN int omp_get_device_num(void);
  • EXTERN int omp_get_num_devices(void);

以下は環境変数です。

  • OMP_DEFAULT_DEVICE で検出された制御デフォルトデバイス:
    • o 負ではない整数値を指定。
  • export OMP_TARGET_OFFLOAD={"MANDATORY" | "DISABLED" | "DEFAULT" }:
    • MANDATORY: GPU またはアクセラレーターで実行されるターゲット領域のコード。
    • DISABLED: CPU で実行されるターゲット領域のコード。
    • DEFAULT: デバイスが利用可能な場合に GPU で実行され、利用できない場合は CPU にフォールバックされるターゲット領域のコード。

OpenMP* 4.5 (Gen9 以降のターゲットを含む) サポート

OpenMP* オフロードランタイムは、GPU ターゲット固有のマッピングを行います。OpenMP* 実行モデルから GPU ハードウェアへのマッピングを以下に示します。

OpenMP* GPU ハードウェア チーム数/スレッド数/SIMD レーン数
チーム サブスライス (SS) DSS ごとに 2+ チーム
スレッド EU スレッド DSS ごとに 64 EU スレッドを利用可能
SIMD レーン (または “チャネル”) SIMD1、SIMD4、SIMD8、SIMD16、SIMD32

GPU サポートでは、複数レベルの並列処理が複数の OpenMP* チーム、スレッド、SIMD レーンを介して有効になり、すべてのレベルでハードウェア・リソースを完全に活用します。インテル® GPU では複数のチームをサブスライスにマッピングできるため、アプリケーションは DSS ごとに 64 以上のスレッドを実行できます。

  • チーム内のスレッド間でハードウェア・バリアの使用が許可されます。
  • OpenMP* スレッドのセマンティクスを許可します (独立した分岐など)。
  • チーム間の同期が可能です。
  • チームを使用してマシン全体を利用できます (より多くのマイグレーションが必要)。
  • OpenMP* simd またはコンパイラーのベクトル化により SIMD を活用します。インテル® oneAPI HPC ツールキットでは、OpenMP* simd が CPU でもサポートされます。GPC では、カーネルレベルでコンパイラーのベクトル化を利用し、SIMT のような SIMD 実行モデルを活用します。MS66 では、インテル® oneAPI HPC ツールキットは GPU 上の OpenMP* simd もサポートします。明示的な SIMD ベクトル化モードでは、afelen(n) 句は安全性をチェックするためにのみ使用され、コンパイラーは simdlen < n を選択できます。simdlen(n) 句は SIMD 命令レーン/チャンネルにマッピングされます。

OpenMP* 3.0 レガシー・アプリケーションで Gen9 以降をターゲットにするオプションのサポート (以前の調査に基づきます)

2 つの新しいオプションがあります。

  • -fiopenmp
  • -fopenmp-targets=spir64

CPU および GPU での OpenMP* とオフロード実行をサポートします。-fiopenmp オプションは、LLVM* での OpenMP* 変換をサポートするミドルエンドを有効にします (clang フロントエンドではサポートされません)。-fopenmp-targets=spir64 オプションを使用すると、コンパイラーは GPU デバイス向けのバイナリーとして x86 + SPIR64 ファットバイナリーを生成できます。

C++ および Fortran の Gen9 以降のディレクティブのサポート強化

OpenMP* オフロードモデルを使用してライブラリー関数を記述し、オフロード領域をさらに高速なバージョンの関数にすることもできます。この利用モデルをサポートするため、コンパイラーは target variant dispatch 構造と拡張機能をサポートし、関数呼び出しの周辺でディスパッチ・コードを発行するようコンパイラーに指示します。構造はオプションとして device 句を受け入れます。構文は次のとおりです。

#pragma omp target variant dispatch [device( n )]
関数呼び出し

ディスパッチ・コードは、関数の基本バージョンと異形バージョンのどちらを呼び出すのかを決定するランタイム・チェック・ルーチンです。GPU が利用可能であれば異形バージョンが呼び出され、それ以外は基本バージョンが呼び出されます。device (n) 句が指定されると、デバイス n が利用可能な場合にのみ異形バージョンが呼び出されます。

関数の基本バージョンと異形バージョンの名前を指定するため、このコンパイラーのリリースでは次のように OpenMP* 5.0 declare variant 構造のサブセットをサポートしています。

#pragma omp declare variant ( variant_func )match(construct={target variant dispatch}, device={arch(gen)})
ベース関数

以下はこの機能を使用した例です。

#include &lt;stdio.h&gt;

#define N 1024

float __attribute__((nothrow, noinline)) vecadd_gpu_offload() {
  float result = 0.0;
  float a[N], b[N];

  #pragma omp target parallel for reduction(+: result) map(to: a, b)
  for (int k=0; k&lt;N; k++) {
    a[k] = k;
    b[k] = k + 1;
    result = result + a[k] + b[k];  
  }
  printf(&quot;GPU version was called. &quot;);
  return result;
}

#pragma omp declare variant(vecadd_gpu_offload) \
            match(construct={target variant dispatch}, device={arch(gen)})
float __attribute__((nothrow, noinline)) vecadd_base() {
  float result = 0.0;              
  float a[N], b[N];

  #pragma omp parallel for reduction(+: result)
  for (int k=0; k&lt;N; k++) {
    a[k] = k;
    b[k] = k + 1;
    result = result + a[k] + b[k];
  }
  printf(&quot;CPU version was called. &quot;);
  return result;
}

int main() {
  float result=0.0;

  #pragma omp target variant dispatch
  {
     result = vecadd_base();
  }

  if (result == 1048576.0) {
    printf(&quot;PASSED: correct results\n&quot;);
    return 0;
  }

  printf(&quot;FAILED: incorrect results\n&quot;);
  return -1;
}

Gen9 以降のターゲット領域の制限に関するドキュメント

OpenMP* オフロードは GPU 向けの OpenCL* ランタイムスタック上にあるため、OpenCL* カーネル関数に適用される次の制限が OpenMP* オフロード領域のコードにもあてはまります。

  • 再帰関数呼び出し (コンパイル時の定数式を除く)
  • プレースメント以外の new と delete
  • goto 文の制限
  • register と thread_local ストレージ修飾子
  • 仮想関数修飾子
  • 関数ポインター (コンパイル時の定数式を除く)
  • 仮想関数
  • 例外処理
  • C++ 標準ライブラリー (GPU では printf のみサポート)
  • ラムダ式から関数への暗黙的なポインター変換
  • 可変関数
  • 可変長配列 (VLA) (タスクモデルと非同期オフロードではサポートされない)

OpenMP* オフロードの利用例

次のコードは、OpenMP* target、teams、distribute および parallel for を組み合わせた簡単な行列乗算の例です。

// matmul.cpp: OpenMP* のオフロードを使用した行列乗算の例
#include &lt;stdio.h&gt;
#include &lt;math.h&gt;
#include &lt;stdlib.h&gt;

#define MAX 128
int A[MAX][MAX], B[MAX][MAX], C[MAX][MAX], C_SERIAL[MAX][MAX];

typedef int BOOL;
typedef int TYPE;

BOOL check_result(TYPE *actual, TYPE *expected, unsigned n) {
    for (unsigned i = 0; i &lt; n; i++) {
        if(actual[i] != expected[i]) {
            printf(&quot;Value mismatch at index = %d. Expected: %d&quot;
                   &quot;, Actual: %d.\n&quot;, i, expected[i], actual[i]);
            return 0;
        }
    }
    return 1;
}

void __attribute__ ((noinline)) Compute()
{
  #pragma omp target teams distribute parallel for map(to: A, B) map(tofrom: C) \
                                                   thread_limit(128)
  {
    for (int i = 0; i &lt; MAX; i++)
    for (int j = 0; j &lt; MAX; j++)
    for (int k = 0; k &lt; MAX; k++)
         C[i][j] += A[i][k] * B[k][j];
  }
}

int main() {
  for (int i = 0; i &lt; MAX; i++)
    for (int j = 0; j &lt; MAX; j++) {
      A[i][j] = i + j - 1;
      B[i][j] = i - j + 1;
    }

  for (int i = 0; i &lt; MAX; i++)
  for (int j = 0; j &lt; MAX; j++)
  for (int k = 0; k &lt; MAX; k++)
       C_SERIAL[i][j] += A[i][k] * B[k][j];

  Compute();

  if (!check_result((int*) &amp;C[0][0], (int*) &amp;C_SERIAL[0][0], MAX * MAX)) {
    printf(&quot;FAILED\n&quot;);
    return 1;
  }

  printf(&quot;PASSED\n&quot;);
  return 0;
}

コンパイルと実行コマンド

注: Linux* では、ホストコードをコンパイルするため GCC 4.8.5 以降をインストールする必要があります。これは、C++ アプリケーション・バイナリー・インターフェイス (ABI) の変更による非互換性を回避するためです。

  1. コンパイル: GPU オフロードを起動するため、icx、icpx、または ifx コンパイル・オプションを使用してソースコードをコンパイルします。

    icx -fiopenmp –fopenmp-targets=spir64 file.cpp matmul_offload.cpp -o matmul

    または

    icpx -fiopenmp –fopenmp-targets=spir64 file.cpp matmul_offload.cpp –o matmul

    または

    ifx -fiopenmp –fopenmp-targets=spir64 file.f90 matmul_offload.f90 –o matmul

  2. 実行: 環境変数を設定します: export OMP_TARGET_OFFLOAD="MANDATORY"。デフォルト値は DEFAULT で、CPU と GPU で実行できることを意味します。以下に例を示します。

    sh-4.2$ ./matmul 
** Program Scope patch lists **
** Kernel Patch Lists : Kernel Name = __omp_offloading_811_600f59f6__Z7Computev_l14 **
BINDING_TABLE_STATE
        Entry[   0 ]        : 00000000 (0)
        Entry[   1 ]        : 00000040 (2)
        Entry[   2 ]        : 00000080 (4)
INTERFACE_DESCRIPTOR_DATA = { 00000000, 00000000, 00000000, 00000000, 000000c3,
00000000, 00000000, 00000003 }
KernelStartPointer =               : 0
Kernel64bitStartPointer =          : 0
SoftwareExceptionEnable =          : 0
MaskStackExceptionEnable =         : 0
IllegalOpcodeExceptionEnable =     : 0
FloatingPointMode =                : 0
ThreadPriority =                   : 0
SingleProgramFlow =                : 0
DenormMode =                       : 0
SamplerCount =                     : 0
SamplerStatePointer =              : 0
BindingTableEntryCount =           : 3
BindingTablePointer =              : 6
ConstantURBEntryReadOffset =       : 0
ConstantURBEntryReadLength =       : 0
NumberOfThreadsInThreadGroup =     : 0
GlobalBarrierEnable =              : 0
SharedLocalMemorySize =            : 0
BarrierEnable =                    : 0
RoundingMode =                     : 0
CrossThreadConstantDataReadLength  : 3
Kernel Name: __omp_offloading_811_600f59f6__Z7Computev_l14 PASSED

GPU 向けに最適化された LIBM 関数のコンパイラー統合の強化

場合によっては、数学関数には精度とパフォーマンスのトレードオフが異なる複数の異形があります。コンパイラーは、オプションによって適切な異形を選択する手段を提供します。コンパイラーでサポートされるすべての fp-model は、インテル® C++ コンパイラーとインテル® Fortran コンパイラーの OpenMP* の GPU へのオフロード機能でサポートされます。インテル® C++ コンパイラーでサポートされる fp-model も移行されます。以下は、OpenCL* 組込み数学関数に基づく Gen9 以降でサポートされる数学関数のリストです。

std::unordered_map&lt;std::string, std::string&gt; llvm::vpo::OCLBuiltin = {
    // float:
    {&quot;sinf&quot;,                  &quot;_Z3sinf&quot;},
    {&quot;cosf&quot;,                  &quot;_Z3cosf&quot;},
    {&quot;tanf&quot;,                  &quot;_Z3tanf&quot;},
    {&quot;erff&quot;,                  &quot;_Z3erff&quot;},
    {&quot;expf&quot;,                  &quot;_Z3expf&quot;},
    {&quot;logf&quot;,                  &quot;_Z3logf&quot;},
    {&quot;log2f&quot;,                 &quot;_Z4log2f&quot;},
    {&quot;powf&quot;,                  &quot;_Z3powff&quot;},
    {&quot;sqrtf&quot;,                 &quot;_Z4sqrtf&quot;},
    {&quot;fmaxf&quot;,                 &quot;_Z4fmaxff&quot;},
    {&quot;llvm.maxnum.f32&quot;,       &quot;_Z4fmaxff&quot;},
    {&quot;fminf&quot;,                 &quot;_Z4fminff&quot;},
    {&quot;llvm.minnum.f32&quot;,       &quot;_Z4fminff&quot;},
    {&quot;fabsf&quot;,                 &quot;_Z4fabsf&quot;},
    {&quot;llvm.fabs.f32&quot;,         &quot;_Z4fabsf&quot;},
    {&quot;ceilf&quot;,                 &quot;_Z4ceilf&quot;},
    {&quot;llvm.ceil.f32&quot;,         &quot;_Z4ceilf&quot;},
    {&quot;floorf&quot;,                &quot;_Z5floorf&quot;},
   {&quot;llvm.floor.f32&quot;,        &quot;_Z5floorf&quot;},

   // double:
   {&quot;sin&quot;,                   &quot;_Z3sind&quot;},
   {&quot;cos&quot;,                   &quot;_Z3cosd&quot;},
   {&quot;tan&quot;,                   &quot;_Z3tand&quot;},
   {&quot;erf&quot;,                   &quot;_Z3erfd&quot;},
   {&quot;exp&quot;,                   &quot;_Z3expd&quot;},
   {&quot;log&quot;,                   &quot;_Z3logd&quot;},
   {&quot;log2&quot;,                  &quot;_Z4log2d&quot;},
   {&quot;pow&quot;,                   &quot;_Z3powdd&quot;},
   {&quot;sqrt&quot;,                  &quot;_Z4sqrtd&quot;},
   {&quot;fmax&quot;,                  &quot;_Z4fmaxdd&quot;},
   {&quot;llvm.maxnum.f64&quot;,       &quot;_Z4fmaxdd&quot;},
   {&quot;fmin&quot;,                  &quot;_Z4fmindd&quot;},
   {&quot;llvm.minnum.f64&quot;,       &quot;_Z4fmindd&quot;},
   {&quot;fabs&quot;,                  &quot;_Z4fabsd&quot;},
   {&quot;llvm.fabs.f64&quot;,         &quot;_Z4fabsd&quot;},
   {&quot;ceil&quot;,                  &quot;_Z4ceild&quot;},
   {&quot;llvm.ceil.f64&quot;,         &quot;_Z4ceild&quot;},
   {&quot;floor&quot;,                 &quot;_Z5floord&quot;},
   {&quot;llvm.floor.f64&quot;,        &quot;_Z5floord&quot;},
   {&quot;invsqrtf&quot;,              &quot;_Z5rsqrtf&quot;},
   {&quot;invsqrt&quot;,               &quot;_Z5rsqrtd&quot;}};

高速逆平方根関数

注: libomptarget ランタイム・ライブラリーには、GPU カーネルの開始/終了時間とデータ転送時間を追跡するパフォーマンス・プロファイル機能が実装されています。この機能を有効にするには、環境変数 LIBOMPTARGET_PROFILE=T を設定します。結果は次のようになります。

GPU Performance (Gen9, export LIBOMPTARGET_PROFILE=T,usec) 
… … 
Kernel Name: __omp_offloading_811_29cbc383__ZN12BlackScholesIdE12execute_partEiii_l368
iteration #0 ...
calling validate ... ok
calling close ...
execution finished in 1134.914ms, total time 0.045min
passed
LIBOMPTARGET_PROFILE:
-- DATA-READ: 16585.256 usec
-- DATA-WRITE: 9980.499 usec
-- EXEC-__omp_offloading_811_29cbc383__ZN12BlackScholesIfE12execute_partEiii_l368: 24048.503 usec

データの読み取りとそのコストは、OpenCL* SVMMap/SVMUnmap とデータコピーを測定しました。以下に例を示します。

INVOKE_CL_RET_FAIL(clEnqueueSVMMap, queue, CL_TRUE, CL_MAP_WRITE, tgt_ptr,
size, 0, nullptr, nullptr);
memcpy(tgt_ptr, hst_ptr, size);
INVOKE_CL_RET_FAIL(clEnqueueSVMUnmap, queue, tgt_ptr, 0, nullptr, nullptr);

サポートされる USM は、L0 API のサポート・スケジュールに合わせて OpenMP* 5.0 の USM 機能をサポートする予定です。

GPU 固有のデバッグ情報の事前統合

GPU 固有のデバッグを可能にするため、export LIBOMPTARGET_DEBUG=1 環境変数がサポートされました。これにより、オフロードのランタイムデバッグ情報を取得できます。デフォルト値は 0 であり、オフロードのランタイムデバッグ情報を出力しないことを意味します。以下の例をご覧ください。

sh-4.2$ export LIBOMPTARGET_DEBUG=1 
sh-4.2$ ./matmul 

Libomptarget --&gt; Loading RTLs...
Libomptarget --&gt; Loading library 'libomptarget.rtl.nios2.so'...
Libomptarget --&gt; Unable to load library 'libomptarget.rtl.nios2.so': libomptarget.rtl.nios2.so: cannot open shared object file: No such file or directory!
Libomptarget --&gt; Loading library 'libomptarget.rtl.x86_64_mic.so'...
Libomptarget --&gt; Successfully loaded library 'libomptarget.rtl.x86_64_mic.so'!
Libomptarget --&gt; No devices supported in this RTL
Libomptarget --&gt; Loading library 'libomptarget.rtl.opencl.so'...
Target OPENCL RTL --&gt; Start initializing OpenCL
Target OPENCL RTL --&gt; cl platform version is OpenCL 2.1 LINUX
Target OPENCL RTL --&gt; Found 1 OpenCL devices
Target OPENCL RTL --&gt; Device#0: Genuine Intel(R) CPU 0000 @ 3.00GHz
Target OPENCL RTL --&gt; max WGs is: 8
Target OPENCL RTL --&gt; max WG size is: 8192
Target OPENCL RTL --&gt; addressing mode is 64 bit
Target OPENCL RTL --&gt; cl platform version is OpenCL 2.1
Target OPENCL RTL --&gt; Found 1 OpenCL devices
Target OPENCL RTL --&gt; Device#0: Intel(R) Gen9 HD Graphics NEO
Target OPENCL RTL --&gt; max WGs is: 24
Target OPENCL RTL --&gt; max WG size is: 256
Target OPENCL RTL --&gt; addressing mode is 64 bit
Libomptarget --&gt; Successfully loaded library 'libomptarget.rtl.opencl.so'!
Libomptarget --&gt; Optional interface: __tgt_rtl_run_target_team_nd_region
Libomptarget --&gt; Optional interface: __tgt_rtl_run_target_region_nowait
Libomptarget --&gt; Optional interface: __tgt_rtl_run_target_team_region_nowait
Libomptarget --&gt; Optional interface: __tgt_rtl_run_target_team_nd_region_nowait
Libomptarget --&gt; Registering RTL libomptarget.rtl.opencl.so supporting 1 devices!
Libomptarget --&gt; Loading library 'libomptarget.rtl.ppc64.so'...
Libomptarget --&gt; Unable to load library 'libomptarget.rtl.ppc64.so': libomptarget.rtl.ppc64.so: cannot open shared object file: No such file or directory!
Libomptarget --&gt; Loading library 'libomptarget.rtl.x86_64.so'...
Libomptarget --&gt; Unable to load library 'libomptarget.rtl.x86_64.so': libomptarget.rtl.x86_64.so: cannot open shared object file: No such file or directory!
Libomptarget --&gt; Loading library 'libomptarget.rtl.cuda.so'...
Libomptarget --&gt; Unable to load library 'libomptarget.rtl.cuda.so': libomptarget.rtl.cuda.so: cannot open shared object file: No such file or directory!
Libomptarget --&gt; Loading library 'libomptarget.rtl.aarch64.so'...
Libomptarget --&gt; Unable to load library 'libomptarget.rtl.aarch64.so': libomptar
get.rtl.aarch64.so: cannot open shared object file: No such file or directory!
Libomptarget --&gt; RTLs loaded!
Target OPENCL RTL --&gt; Target binary is VALID
Libomptarget --&gt; Image 0x000000000060d0a0 is compatible with RTL libomptarget.rtl.opencl.so!
Libomptarget --&gt; RTL 0x0000000002012250 has index 0!
Libomptarget --&gt; Registering image 0x000000000060d0a0 with RTL libomptarget.rtl.opencl.so!
Libomptarget --&gt; Done registering entries!
Libomptarget --&gt; Entering target region with entry point 0x000000000040a2a0 anddevice Id -1
Libomptarget --&gt; Checking whether device 0 is ready.
Libomptarget --&gt; Is the device 0 (local ID 0) initialized? 0
Target OPENCL RTL --&gt; Initialize OpenCL device
Libomptarget --&gt; Device 0 is ready to use.
Target OPENCL RTL --&gt; Dev 0: load binary from 0x000000000060d0a0 image
Target OPENCL RTL --&gt; Expecting to have 1 entries defined.
Target OPENCL RTL --&gt; Found device RTL: /home/users/xtian/cmplr/dev_xmain/builds/xmainoffloadlinuxefi2_debug/llvm/lib/libomptarget-opencl.a

** Program Scope patch lists **
Libomptarget --&gt; Creating new map entry: HstBase=0x00000000006224f0, HstBegin=0x
00000000006224f0, HstEnd=0x000000000062c130, TgtBegin=0x0000000002100ad0
Libomptarget --&gt; There are 40000 bytes allocated at target address 0x00000000021
00ad0 - is new
… … … … 
Libomptarget --&gt; Launching target execution __omp_offloading_811_600f59f6__Z7Com
putev_l14 with pointer 0x0000000002a994a0 (index=0).
Target OPENCL RTL --&gt; OpenCL: Kernel Arg 0 set successfully
Target OPENCL RTL --&gt; OpenCL: Kernel Arg 1 set successfully
Target OPENCL RTL --&gt; OpenCL: Kernel Arg 2 set successfully
… … … 
Libomptarget --&gt; Image 0x000000000060d0a0 is compatible with RTL 0x0000000002012250!
Libomptarget --&gt; Unregistered image 0x000000000060d0a0 from RTL 0x0000000002012250!
Libomptarget --&gt; Done unregistering images!
Libomptarget --&gt; Removing translation table for descriptor 0x000000000060ea50
Libomptarget --&gt; Done unregistering library!

注: インテル® GPU を使用するプログラミングは、ほかの GPU 向けのプログラミングに似ています。異なる GPU (マイクロ) アーキテクチャーは異なる動作になります。新しい (マイクロ) アーキテクチャー向けにコードをチューニングするのは、機能の移行よりも困難です。インテルは、後者の負担を軽減するためコンパイラー、ライブラリー、およびツールの提供に取り組んでいますが、これによってパフォーマンス最適化の要求がなくなるわけではありません。

さらに詳しく

ドキュメント 説明とリンク
OpenMP* 5.0 仕様 (PDF) OpenMP* 仕様 (英語) では OpenMP* オフロードをデバイス向けに使用する方法を説明しています。
GNU* C/C++ ライブラリー 2 つの ABI の使用 (英語) についての説明です。
OpenMP* のサポートに関する SC’16 と SC’17 LLVM-HPC ワークショップの文書 明示的な並列化と SIMD ベクトル化のための LLVM コンパイラーの実装。
LLVM-HPC@SC 2017: 4:1-4:11

並列化、SIMD ベクトル化、オフロードのための LLVM フレームワークと IR 拡張。
LLVM-HPC@SC 2016: 21-31

法務上の注意書きと最適化に関する注意事項

インテル® テクノロジーの機能と利点はシステム構成によって異なり、対応するハードウェアやソフトウェア、またはサービスの有効化が必要となる場合があります。実際の性能はシステム構成によって異なります。

絶対的なセキュリティーを提供できるコンピューター・システムはありません。

本資料は、(明示されているか否かにかかわらず、また禁反言によるとよらずにかかわらず) いかなる知的財産権のライセンスも許諾するものではありません。

本資料で説明されている製品およびサービスには、不具合が含まれている可能性があり、公表されている仕様とは異なる動作をする場合があります。現在確認済みのエラッタについては、インテルまでお問い合わせください。

インテルは、明示されているか否かにかかわらず、いかなる保証もいたしません。ここにいう保証には、商品適格性、特定目的への適合性、および非侵害性の黙示の保証、ならびに履行の過程、取引の過程、または取引での使用から生じるあらゆる保証を含みますが、これらに限定されるわけではありません。

Intel、インテル、Intel ロゴは、アメリカ合衆国および / またはその他の国における Intel Corporation またはその子会社の商標です。

* その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

製品とパフォーマンス情報

1 インテル® コンパイラーでは、インテル® マイクロプロセッサーに限定されない最適化に関して、他社製マイクロプロセッサー用に同等の最適化を行えないことがあります。これには、インテル® ストリーミング SIMD 拡張命令 2、インテル® ストリーミング SIMD 拡張命令 3、インテル® ストリーミング SIMD 拡張命令 3 補足命令などの最適化が該当します。インテルは、他社製マイクロプロセッサーに関して、いかなる最適化の利用、機能、または効果も保証いたしません。本製品のマイクロプロセッサー依存の最適化は、インテル® マイクロプロセッサーでの使用を前提としています。インテル® マイクロアーキテクチャーに限定されない最適化のなかにも、インテル® マイクロプロセッサー用のものがあります。この注意事項で言及した命令セットの詳細については、該当する製品のユーザー・リファレンス・ガイドを参照してください。

注意事項の改訂 #20110804

関連記事