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

2022年10月8日

oneAPI DPC++ 導入ガイド

この記事は、インテル® デベロッパー・ゾーンに公開されている「Get Started with OpenMP* Offload to GPU for the Intel® oneAPI DPC/C++ Compiler and Intel® Fortran Compiler」の日本語参考訳です。原文は更新される可能性があります。原文と翻訳文の内容が異なる場合は原文を優先してください。

Version: 2022.1
原文最終更新日: 04/11/2022

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

はじめに

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

icx と ifx の OpenMP* 5.0/5.1 サポート

icx および ifx コンパイラーは、GPU ターゲットのデバイスサポートを含む OpenMP* 5.0 および 5.1 のほとんどの機能をサポートします。サポートの詳細については、『インテル® oneAPI DPC++/C++ コンパイラー・デベロッパー・ガイドおよびリファレンス』の「OpenMP* サポート」 (英語) を参照してください。

icx および ifx コンパイラーで GPU と CPU 向けにサポートされる OpenMP* ディレクティブ:

C++ と Fortran で共通のキーワードリストを示します。

allocate
atomic
barrier
critical
declare simd
declare target
declare variant
dispatch
distribute
distribute parallel for/do
distribute parallel for/do simd
distribute simd
for/do
for/do simd
loop
master
parallel
parallel for/do
parallel for/do simd
parallel loop
parallel sections
sections
simd
single
target
target data
target enter data
target exit data
target parallel
target parallel for/do
target parallel for/do simd
target parallel loop
target simd
target teams
target teams distribute
target teams distribute parallel for/do
target teams distribute parallel for/do simd
target teams distribute simd
target teams loop
target variant dispatch (インテル拡張)
target update
teams
teams distribute
teams distribute parallel for/do
teams distribute parallel for/do simd
teams distribute simd
teams loop

GPU ランタイムルーチン:

デバイスルーチンを CPU から呼び出してデバイス情報を取得し、GPU にデバイスメモリーを割り当てることができます。

omp_get_initial_device
omp_get_interop_int
omp_get_interop_name
omp_get_interop_ptr
omp_get_interop_rc_desc
omp_get_interop_str
omp_get_interop_type_desc
omp_get_mapped_ptr
omp_get_num_devices
omp_get_num_interop_properties
omp_is_initial_device
omp_target_alloc
omp_target_alloc_device (インテル拡張)
omp_target_alloc_host (インテル拡張)
omp_target_alloc_shared (インテル拡張)
omp_target_associate_ptr
omp_target_disassociate_ptr
omp_target_free
omp_target_is_accessible
omp_target_is_present
omp_target_memcpy
omp_target_memcpy_rect

GPU で呼び出し可能なデバイス・ランタイム・ルーチン:

omp_get_device_num
omp_get_max_threads
omp_get_num_devices
omp_get_num_procs
omp_get_num_teams
omp_get_num_threads
omp_get_team_num
omp_get_team_size
omp_get_thread_limit
omp_get_thread_num
omp_in_parallel
omp_is_initial_device

環境変数:

OMP_DEFAULT_DEVICE: デフォルトデバイスを設定します。
OMP_TARGET_OFFLOAD: デバイスまたはホストへのオフロードを制御します。
LIBOMPTARGET_PLUGIN: OpenCL* またはレベルゼロを選択します。
LIBOMPTARGET_DEBUG: デバッグ情報の生成を有効にします。

Gen9 ターゲットをサポートするオプション

2 つの新しいオプションがサポートされます。

-qopenmp
-fopenmp-targets=spir64

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

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 <stdio.h>
#include <math.h>
#include <stdlib.h>

#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 < n; i++) {
        if(actual[i] != expected[i]) {
            printf("Value mismatch at index = %d. Expected: %d"
                   ", Actual: %d.\n", 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 < MAX; i++)
    for (int j = 0; j < MAX; j++)
    for (int k = 0; k < MAX; k++)
         C[i][j] += A[i][k] * B[k][j];
  }
}

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

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

  Compute();

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

  printf("PASSED\n");
  return 0;
}

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

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

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

$ icx -qopenmp -fopenmp-targets=spir64 matmul_offload.c -o matmul

または

$ icpx -qopenmp -fopenmp-targets=spir64 matmul_offload.cpp -o matmul

または

$ ifx -qopenmp -fopenmp-targets=spir64 matmul_offload.f90 -o matmul

実行: OMP_TARGET_OFFLOAD 環境変数を MANDATORY に設定して、オフロードを強制します。
```
$ export OMP_TARGET_OFFLOAD=MANDATORY
```
デフォルト値は DEFAULT で、デバイスが利用可能な場合は GPU で実行され、利用できない場合は CPU にフォールバックされます。以下に例を示します。
```
$ ./matmul
PASSED
```

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

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

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

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

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

================================================================================
LIBOMPTARGET_PLUGIN_PROFILE(LEVEL0) for OMP DEVICE(0) Intel(R) UHD Graphics 630 [0x3e92], Thread 0
--------------------------------------------------------------------------------
-- Kernel 0                  : __omp_offloading_3a_dca6bdd1_MAIN___l30
--------------------------------------------------------------------------------
-- Name                      :     Host Time (msec)   Device Time (msec)
-- Compiling                 :             2632.679             2632.679
-- DataAlloc                 :               11.350               11.350
-- DataRead (Device to Host) :                2.633                2.633
-- DataWrite (Host to Device):                8.398                8.398
-- Kernel 0                  :             5279.800             5258.379
-- OffloadEntriesInit        :                2.731                2.731
-- Total                     :             7937.591             7916.170
================================================================================

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

GPU 固有のデバッグを可能にするため、LIBOMPTARGET_DEBUG 環境変数がサポートされました。この環境変数を 1 に設定すると、オフロードのランタイムデバッグ情報を取得できます。デフォルト値は 0 であり、オフロードのランタイムデバッグ情報は出力されません。次の例では、前述の行列乗算を実行して GPU 固有のデバッグ情報を取得しています。

$ icx -v
Intel(R) oneAPI DPC++/C++ Compiler 2021.4.0 …
$ icx -qopenmp -fopenmp-targets=spir64 matmul.cpp -o matmul
$ export LIBOMPTARGET_DEBUG=1 
$ ./matmul 
Libomptarget --> Init target library!
Libomptarget --> Initialized OMPT
Libomptarget --> Loading RTLs...
Libomptarget --> Loading library 'libomptarget.rtl.level0.so'...
Target LEVEL0 RTL --> Init Level0 plugin!
Target LEVEL0 RTL --> omp_get_thread_limit() returned 2147483647
Target LEVEL0 RTL --> omp_get_max_teams() returned 0
Libomptarget --> Successfully loaded library 'libomptarget.rtl.level0.so'!
Target LEVEL0 RTL --> Looking for Level0 devices...
Target LEVEL0 RTL --> Initialized L0, API 10002
Target LEVEL0 RTL --> Found 1 driver(s)!
Target LEVEL0 RTL --> Found copy command queue for device 0x0000000000e4bc00, ordinal = 1
Target LEVEL0 RTL --> Found a GPU device, Name = Intel(R) Iris(R) Xe MAX Graphics [0x4905]
Target LEVEL0 RTL --> Found a GPU device, Name = Intel(R) UHD Graphics 630 [0x3e98]
Target LEVEL0 RTL --> No subdevices are found for device 0x0000000000e4bc00 at level 0
Target LEVEL0 RTL --> Could not find multi-context command queue group for device 0x0000000000e4bc00
Target LEVEL0 RTL --> No subdevices are found for device 0x0000000000ea31e0 at level 0
Target LEVEL0 RTL --> Could not find multi-context command queue group for device 0x0000000000ea31e0
Target LEVEL0 RTL --> Found 2 root devices, 2 total devices.
Target LEVEL0 RTL --> List of devices (DeviceID[.SubDeviceLevel.SubDeviceID])
Target LEVEL0 RTL --> -- 0
Target LEVEL0 RTL --> -- 1
Target LEVEL0 RTL --> Driver API version is 10001
Target LEVEL0 RTL --> Interop property IDs, Names, Descriptions
Target LEVEL0 RTL --> -- 0, device_num_eus, intptr_t, total number of EUs
Target LEVEL0 RTL --> -- 1, device_num_threads_per_eu, intptr_t, number of threads per EU
Target LEVEL0 RTL --> -- 2, device_eu_simd_width, intptr_t, physical EU simd width
Target LEVEL0 RTL --> -- 3, device_num_eus_per_subslice, intptr_t, number of EUs per sub-slice
Target LEVEL0 RTL --> -- 4, device_num_subslices_per_slice, intptr_t, number of sub-slices per slice
Target LEVEL0 RTL --> -- 5, device_num_slices, intptr_t, number of slices
Target LEVEL0 RTL --> Returning 2 top-level devices
Libomptarget --> Registering RTL libomptarget.rtl.level0.so supporting 2 devices!
Libomptarget --> Optional interface: __tgt_rtl_data_alloc_base
Libomptarget --> Optional interface: __tgt_rtl_data_alloc_user
Libomptarget --> Optional interface: __tgt_rtl_data_alloc_explicit
Libomptarget --> Optional interface: __tgt_rtl_data_alloc_managed
Libomptarget --> Optional interface: __tgt_rtl_data_submit_nowait
Libomptarget --> Optional interface: __tgt_rtl_data_retrieve_nowait
Libomptarget --> Optional interface: __tgt_rtl_create_offload_queue
Libomptarget --> Optional interface: __tgt_rtl_release_offload_queue
Libomptarget --> Optional interface: __tgt_rtl_get_platform_handle
Libomptarget --> Optional interface: __tgt_rtl_set_device_handle
Libomptarget --> Optional interface: __tgt_rtl_get_context_handle
Libomptarget --> Optional interface: __tgt_rtl_init_ompt
Libomptarget --> Optional interface: __tgt_rtl_is_device_accessible_ptr
Libomptarget --> Optional interface: __tgt_rtl_manifest_data_for_region
Libomptarget --> Optional interface: __tgt_rtl_push_subdevice
Libomptarget --> Optional interface: __tgt_rtl_pop_subdevice
Libomptarget --> Optional interface: __tgt_rtl_add_build_options
Libomptarget --> Optional interface: __tgt_rtl_is_supported_device
Libomptarget --> Optional interface: __tgt_rtl_deinit
Libomptarget --> Optional interface: __tgt_rtl_create_interop
Libomptarget --> Optional interface: __tgt_rtl_release_interop
Libomptarget --> Optional interface: __tgt_rtl_use_interop
Libomptarget --> Optional interface: __tgt_rtl_get_num_interop_properties
Libomptarget --> Optional interface: __tgt_rtl_get_interop_property_value
Libomptarget --> Optional interface: __tgt_rtl_get_interop_property_info
Libomptarget --> Optional interface: __tgt_rtl_get_interop_rc_desc
Libomptarget --> Optional interface: __tgt_rtl_get_num_sub_devices
Libomptarget --> Optional interface: __tgt_rtl_is_accessible_addr_range
Libomptarget --> Optional interface: __tgt_rtl_run_target_team_nd_region
Libomptarget --> Optional interface: __tgt_rtl_run_target_region_nowait
Libomptarget --> Optional interface: __tgt_rtl_run_target_team_region_nowait
Libomptarget --> Optional interface: __tgt_rtl_run_target_team_nd_region_nowait
Target LEVEL0 RTL --> Initialized OMPT
Libomptarget --> Loading library 'libomptarget.rtl.opencl.so'...
Target OPENCL RTL --> Init OpenCL plugin!
Target OPENCL RTL --> omp_get_thread_limit() returned 2147483647
Target OPENCL RTL --> omp_get_max_teams() returned 0
Target OPENCL RTL --> Target device type is set to GPU
Libomptarget --> Successfully loaded library 'libomptarget.rtl.opencl.so'!
Target OPENCL RTL --> Start initializing OpenCL
Target OPENCL RTL --> Platform OpenCL 3.0  has 1 Devices
Target OPENCL RTL --> Extension clGetMemAllocInfoINTEL is found.
Target OPENCL RTL --> Extension clHostMemAllocINTEL is found.
Target OPENCL RTL --> Extension clDeviceMemAllocINTEL is found.
Target OPENCL RTL --> Extension clSharedMemAllocINTEL is found.
Target OPENCL RTL --> Extension clMemFreeINTEL is found.
Target OPENCL RTL --> Extension clSetKernelArgMemPointerINTEL is found.
Target OPENCL RTL --> Extension clEnqueueMemcpyINTEL is found.
Target OPENCL RTL --> Extension clGetDeviceGlobalVariablePointerINTEL is found.
Target OPENCL RTL --> Extension clGetKernelSuggestedLocalWorkSizeINTEL is found.
Target OPENCL RTL --> Extension clSetProgramSpecializationConstant is found.
Target OPENCL RTL --> Platform OpenCL 3.0  has 1 Devices
Target OPENCL RTL --> Extension clGetMemAllocInfoINTEL is found.
Target OPENCL RTL --> Extension clHostMemAllocINTEL is found.
Target OPENCL RTL --> Extension clDeviceMemAllocINTEL is found.
Target OPENCL RTL --> Extension clSharedMemAllocINTEL is found.
Target OPENCL RTL --> Extension clMemFreeINTEL is found.
Target OPENCL RTL --> Extension clSetKernelArgMemPointerINTEL is found.
Target OPENCL RTL --> Extension clEnqueueMemcpyINTEL is found.
Target OPENCL RTL --> Extension clGetDeviceGlobalVariablePointerINTEL is found.
Target OPENCL RTL --> Extension clGetKernelSuggestedLocalWorkSizeINTEL is found.
Target OPENCL RTL --> Extension clSetProgramSpecializationConstant is found.
Target OPENCL RTL --> Device 0: Intel(R) Iris(R) Xe MAX Graphics [0x4905]
Target OPENCL RTL --> Number of execution units on the device is 96
Target OPENCL RTL --> Maximum work group size for the device is 512
Target OPENCL RTL --> Maximum memory allocation size is 4044357632
Target OPENCL RTL --> Addressing mode is 64 bit
Target OPENCL RTL --> Device local mem size: 65536
Target OPENCL RTL --> Device 1: Intel(R) UHD Graphics 630 [0x3e98]
Target OPENCL RTL --> Number of execution units on the device is 24
Target OPENCL RTL --> Maximum work group size for the device is 256
Target OPENCL RTL --> Maximum memory allocation size is 4294959104
Target OPENCL RTL --> Addressing mode is 64 bit
Target OPENCL RTL --> Device local mem size: 65536
Libomptarget --> Registering RTL libomptarget.rtl.opencl.so supporting 2 devices!
Libomptarget --> Optional interface: __tgt_rtl_data_alloc_base
Libomptarget --> Optional interface: __tgt_rtl_data_alloc_user
Libomptarget --> Optional interface: __tgt_rtl_data_alloc_explicit
Libomptarget --> Optional interface: __tgt_rtl_data_alloc_managed
Libomptarget --> Optional interface: __tgt_rtl_data_submit_nowait
Libomptarget --> Optional interface: __tgt_rtl_data_retrieve_nowait
Libomptarget --> Optional interface: __tgt_rtl_create_offload_queue
Libomptarget --> Optional interface: __tgt_rtl_release_offload_queue
Libomptarget --> Optional interface: __tgt_rtl_get_device_name
Libomptarget --> Optional interface: __tgt_rtl_get_platform_handle
Libomptarget --> Optional interface: __tgt_rtl_set_device_handle
Libomptarget --> Optional interface: __tgt_rtl_get_context_handle
Libomptarget --> Optional interface: __tgt_rtl_get_data_alloc_info
Libomptarget --> Optional interface: __tgt_rtl_init_ompt
Libomptarget --> Optional interface: __tgt_rtl_is_device_accessible_ptr
Libomptarget --> Optional interface: __tgt_rtl_manifest_data_for_region
Libomptarget --> Optional interface: __tgt_rtl_add_build_options
Libomptarget --> Optional interface: __tgt_rtl_is_supported_device
Libomptarget --> Optional interface: __tgt_rtl_deinit
Libomptarget --> Optional interface: __tgt_rtl_create_interop
Libomptarget --> Optional interface: __tgt_rtl_release_interop
Libomptarget --> Optional interface: __tgt_rtl_use_interop
Libomptarget --> Optional interface: __tgt_rtl_get_num_interop_properties
Libomptarget --> Optional interface: __tgt_rtl_get_interop_property_value
Libomptarget --> Optional interface: __tgt_rtl_get_interop_property_info
Libomptarget --> Optional interface: __tgt_rtl_get_interop_rc_desc
Libomptarget --> Optional interface: __tgt_rtl_is_accessible_addr_range
Libomptarget --> Optional interface: __tgt_rtl_run_target_team_nd_region
Libomptarget --> Optional interface: __tgt_rtl_run_target_region_nowait
Libomptarget --> Optional interface: __tgt_rtl_run_target_team_region_nowait
Libomptarget --> Optional interface: __tgt_rtl_run_target_team_nd_region_nowait
Target OPENCL RTL --> Initialized OMPT
Libomptarget --> Loading library 'libomptarget.rtl.ppc64.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.ppc64.so': libomptarget.rtl.ppc64.so: cannot open shared object file: No such file or directory!
Libomptarget --> Loading library 'libomptarget.rtl.x86_64.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.x86_64.so': libffi.so.6: cannot open shared object file: No such file or directory!
Libomptarget --> Loading library 'libomptarget.rtl.cuda.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.cuda.so': libomptarget.rtl.cuda.so: cannot open shared object file: No such file or directory!
Libomptarget --> Loading library 'libomptarget.rtl.aarch64.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.aarch64.so': libomptarget.rtl.aarch64.so: cannot open shared object file: No such file or directory!
Libomptarget --> Loading library 'libomptarget.rtl.ve.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.ve.so': libomptarget.rtl.ve.so: cannot open shared object file: No such file or directory!
Libomptarget --> Loading library 'libomptarget.rtl.amdgpu.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.amdgpu.so': libomptarget.rtl.amdgpu.so: cannot open shared object file: No such file or directory!
Libomptarget --> Loading library 'libomptarget.rtl.rpc.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.rpc.so': libomptarget.rtl.rpc.so: cannot open shared object file: No such file or directory!
Libomptarget --> RTLs loaded!
Target LEVEL0 RTL --> Target binary is a valid oneAPI OpenMP image.
Libomptarget --> Image 0x00000000004021d0 is compatible with RTL libomptarget.rtl.level0.so!
Libomptarget --> RTL 0x0000000000cd6f20 has index 0!
Libomptarget --> Registering image 0x00000000004021d0 with RTL libomptarget.rtl.level0.so!
Libomptarget --> Done registering entries!
Libomptarget --> Entering target region with entry point 0x00000000004020b0 and device Id 0
Libomptarget --> Call to omp_get_num_devices returning 2
Libomptarget --> Default TARGET OFFLOAD policy is now mandatory (devices were found)
Libomptarget --> Call to omp_get_num_devices returning 2
Libomptarget --> Call to omp_get_num_devices returning 2
Libomptarget --> Call to omp_get_initial_device returning 2
Libomptarget --> Checking whether device 0 is ready.
Libomptarget --> Is the device 0 (local ID 0) initialized? 0
Target LEVEL0 RTL --> Initialize requires flags to 1
Target LEVEL0 RTL --> Allocated a host memory object 0x0000149189791000
Target LEVEL0 RTL --> Initialized host memory pool for device 0x0000000000000000: AllocMax = 1048576, Capacity = 4, PoolSizeMax = 268435456
Target LEVEL0 RTL --> Allocated a shared memory object 0x0000149189791000
Target LEVEL0 RTL --> Initialized shared memory pool for device 0x0000000000e4bc00: AllocMax = 65536, Capacity = 1, PoolSizeMax = 268435456
Target LEVEL0 RTL --> Allocated a shared memory object 0x0000149189791000
Target LEVEL0 RTL --> Initialized shared memory pool for device 0x0000000000ea31e0: AllocMax = 1048576, Capacity = 4, PoolSizeMax = 268435456
Target LEVEL0 RTL --> Allocated a device memory object 0x00003b7070bc0000
Target LEVEL0 RTL --> Initialized device memory pool for device 0x0000000000e4bc00: AllocMax = 1048576, Capacity = 4, PoolSizeMax = 268435456
Target LEVEL0 RTL --> Allocated a device memory object 0xffffd556aa7e0000
Target LEVEL0 RTL --> Initialized device memory pool for device 0x0000000000ea31e0: AllocMax = 1048576, Capacity = 4, PoolSizeMax = 268435456
Target LEVEL0 RTL --> Initialized Level0 device 0
Libomptarget --> Device 0 is ready to use.
Target LEVEL0 RTL --> Device 0: Loading binary from 0x00000000004021d0
Target LEVEL0 RTL --> Expecting to have 1 entries defined
Target LEVEL0 RTL --> Base L0 module compilation options: -cl-std=CL2.0
Target LEVEL0 RTL --> Created module from image #0.
Target LEVEL0 RTL --> Kernel 0: Entry = 0x00000000004020b0, Name = __omp_offloading_42_d74c0b80__Z7Computev_l25, NumArgs = 5, Handle = 0x000000000186dff0
Target LEVEL0 RTL --> Looking up device global variable '__omp_spirv_program_data' of size 48 bytes on device 0.
Target LEVEL0 RTL --> Global variable lookup succeeded.
Target LEVEL0 RTL --> Created a command list 0x0000000001c94a50 for device 0.
Target LEVEL0 RTL --> Created a command queue 0x00000000016a2260 for device 0.
Libomptarget --> Entry  0: Base=0x00000000004651a0, Begin=0x00000000004651a0, Size=65536, Type=0x23, Name=unknown
Libomptarget --> Entry  1: Base=0x00000000004451a0, Begin=0x00000000004451a0, Size=65536, Type=0x21, Name=unknown
Libomptarget --> Entry  2: Base=0x00000000004551a0, Begin=0x00000000004551a0, Size=65536, Type=0x21, Name=unknown
Libomptarget --> Entry  3: Base=0x0000000000000000, Begin=0x0000000000000000, Size=0, Type=0x120, Name=unknown
Libomptarget --> Entry  4: Base=0x000000000000007f, Begin=0x000000000000007f, Size=0, Type=0x120, Name=unknown
Libomptarget --> Entry  5: Base=0x00007ffc42aa36a0, Begin=0x00007ffc42aa36a0, Size=32, Type=0x800, Name=unknown
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004651a0, Size=65536)...
Target LEVEL0 RTL --> Ptr 0x00000000004651a0 is not a device accessible memory pointer.
Target LEVEL0 RTL --> Allocated a device memory object 0x00003b7070b70000
Target LEVEL0 RTL --> New block allocation for device memory pool: base = 0x00003b7070b70000, size = 262144, pool size = 262144
Target LEVEL0 RTL --> Allocated target memory 0x00003b7070b70000 (Base: 0x00003b7070b70000, Size: 65536) from memory pool for host ptr 0x00000000004651a0
Libomptarget --> Creating new map entry with HstPtrBegin=0x00000000004651a0, TgtPtrBegin=0x00003b7070b70000, Size=65536, RefCount=1, Name=unknown
Libomptarget --> There are 65536 bytes allocated at target address 0x00003b7070b70000 - is new
Libomptarget --> Moving 65536 bytes (hst:0x00000000004651a0) -> (tgt:0x00003b7070b70000)
Target LEVEL0 RTL --> Copy Engine is used for data transfer
Target LEVEL0 RTL --> Copied 65536 bytes (hst:0x00000000004651a0) -> (tgt:0x00003b7070b70000)
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004451a0, Size=65536)...
Target LEVEL0 RTL --> Ptr 0x00000000004451a0 is not a device accessible memory pointer.
Target LEVEL0 RTL --> Allocated target memory 0x00003b7070b80000 (Base: 0x00003b7070b80000, Size: 65536) from memory pool for host ptr 0x00000000004451a0
Libomptarget --> Creating new map entry with HstPtrBegin=0x00000000004451a0, TgtPtrBegin=0x00003b7070b80000, Size=65536, RefCount=1, Name=unknown
Libomptarget --> There are 65536 bytes allocated at target address 0x00003b7070b80000 - is new
Libomptarget --> Moving 65536 bytes (hst:0x00000000004451a0) -> (tgt:0x00003b7070b80000)
Target LEVEL0 RTL --> Copy Engine is used for data transfer
Target LEVEL0 RTL --> Copied 65536 bytes (hst:0x00000000004451a0) -> (tgt:0x00003b7070b80000)
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004551a0, Size=65536)...
Target LEVEL0 RTL --> Ptr 0x00000000004551a0 is not a device accessible memory pointer.
Target LEVEL0 RTL --> Allocated target memory 0x00003b7070b90000 (Base: 0x00003b7070b90000, Size: 65536) from memory pool for host ptr 0x00000000004551a0
Libomptarget --> Creating new map entry with HstPtrBegin=0x00000000004551a0, TgtPtrBegin=0x00003b7070b90000, Size=65536, RefCount=1, Name=unknown
Libomptarget --> There are 65536 bytes allocated at target address 0x00003b7070b90000 - is new
Libomptarget --> Moving 65536 bytes (hst:0x00000000004551a0) -> (tgt:0x00003b7070b90000)
Target LEVEL0 RTL --> Copy Engine is used for data transfer
Target LEVEL0 RTL --> Copied 65536 bytes (hst:0x00000000004551a0) -> (tgt:0x00003b7070b90000)
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004651a0, Size=65536)...
Libomptarget --> Mapping exists with HstPtrBegin=0x00000000004651a0, TgtPtrBegin=0x00003b7070b70000, Size=65536, RefCount=1 (update suppressed)
Libomptarget --> Obtained target argument (Begin: 0x00003b7070b70000, Offset: 0) from host pointer 0x00000000004651a0
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004451a0, Size=65536)...
Libomptarget --> Mapping exists with HstPtrBegin=0x00000000004451a0, TgtPtrBegin=0x00003b7070b80000, Size=65536, RefCount=1 (update suppressed)
Libomptarget --> Obtained target argument (Begin: 0x00003b7070b80000, Offset: 0) from host pointer 0x00000000004451a0
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004551a0, Size=65536)...
Libomptarget --> Mapping exists with HstPtrBegin=0x00000000004551a0, TgtPtrBegin=0x00003b7070b90000, Size=65536, RefCount=1 (update suppressed)
Libomptarget --> Obtained target argument (Begin: 0x00003b7070b90000, Offset: 0) from host pointer 0x00000000004551a0
Libomptarget --> Forwarding first-private value 0x0000000000000000 to the target construct
Libomptarget --> Forwarding first-private value 0x000000000000007f to the target construct
Libomptarget --> Launching target execution __omp_offloading_42_d74c0b80__Z7Computev_l25 with pointer 0x0000000001091720 (index=0).
Libomptarget --> Manifesting used target pointers:
Target LEVEL0 RTL --> Executing a kernel 0x0000000001091720...
Target LEVEL0 RTL --> Kernel argument 0 (value: 0x00003b7070b70000) was set successfully
Target LEVEL0 RTL --> Kernel argument 1 (value: 0x00003b7070b80000) was set successfully
Target LEVEL0 RTL --> Kernel argument 2 (value: 0x00003b7070b90000) was set successfully
Target LEVEL0 RTL --> Kernel argument 3 (value: 0x0000000000000000) was set successfully
Target LEVEL0 RTL --> Kernel argument 4 (value: 0x000000000000007f) was set successfully
Target LEVEL0 RTL --> Setting indirect access flags 0x0000000000000000
Target LEVEL0 RTL --> Assumed kernel SIMD width is 16
Target LEVEL0 RTL --> Preferred group size is multiple of 32
Target LEVEL0 RTL --> Max group size is set to 128 (thread_limit clause)
Target LEVEL0 RTL --> Level 0: Lb = 0, Ub = 127, Stride = 1
Target LEVEL0 RTL --> Group sizes = {128, 1, 1}
Target LEVEL0 RTL --> Group counts = {1, 1, 1}
Target LEVEL0 RTL --> Created a command list 0x0000000001d040c0 for device 0.
Target LEVEL0 RTL --> Created a command queue 0x00000000011179e0 for device 0.
Target LEVEL0 RTL --> Executed a kernel 0x0000000001091720
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004551a0, Size=65536)...
Libomptarget --> Mapping exists with HstPtrBegin=0x00000000004551a0, TgtPtrBegin=0x00003b7070b90000, Size=65536, RefCount=1 (deferred final decrement)
Libomptarget --> There are 65536 bytes allocated at target address 0x00003b7070b90000 - is last
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004451a0, Size=65536)...
Libomptarget --> Mapping exists with HstPtrBegin=0x00000000004451a0, TgtPtrBegin=0x00003b7070b80000, Size=65536, RefCount=1 (deferred final decrement)
Libomptarget --> There are 65536 bytes allocated at target address 0x00003b7070b80000 - is last
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004651a0, Size=65536)...
Libomptarget --> Mapping exists with HstPtrBegin=0x00000000004651a0, TgtPtrBegin=0x00003b7070b70000, Size=65536, RefCount=1 (deferred final decrement)
Libomptarget --> There are 65536 bytes allocated at target address 0x00003b7070b70000 - is last
Libomptarget --> Moving 65536 bytes (tgt:0x00003b7070b70000) -> (hst:0x00000000004651a0)
Target LEVEL0 RTL --> Copy Engine is used for data transfer
Target LEVEL0 RTL --> Copied 65536 bytes (tgt:0x00003b7070b70000) -> (hst:0x00000000004651a0)
Target LEVEL0 RTL --> Ptr 0x00000000004551a0 is not a device accessible memory pointer.
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004551a0, Size=65536)...
Libomptarget --> Deleting tgt data 0x00003b7070b90000 of size 65536
Target LEVEL0 RTL --> Returned device memory 0x00003b7070b90000 to memory pool
Libomptarget --> Removing map entry with HstPtrBegin=0x00000000004551a0, TgtPtrBegin=0x00003b7070b90000, Size=65536, Name=unknown
Target LEVEL0 RTL --> Ptr 0x00000000004451a0 is not a device accessible memory pointer.
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004451a0, Size=65536)...
Libomptarget --> Deleting tgt data 0x00003b7070b80000 of size 65536
Target LEVEL0 RTL --> Returned device memory 0x00003b7070b80000 to memory pool
Libomptarget --> Removing map entry with HstPtrBegin=0x00000000004451a0, TgtPtrBegin=0x00003b7070b80000, Size=65536, Name=unknown
Target LEVEL0 RTL --> Ptr 0x00000000004651a0 is not a device accessible memory pointer.
Libomptarget --> Looking up mapping(HstPtrBegin=0x00000000004651a0, Size=65536)...
Libomptarget --> Deleting tgt data 0x00003b7070b70000 of size 65536
Target LEVEL0 RTL --> Returned device memory 0x00003b7070b70000 to memory pool
Libomptarget --> Removing map entry with HstPtrBegin=0x00000000004651a0, TgtPtrBegin=0x00003b7070b70000, Size=65536, Name=unknown
Target OPENCL RTL --> Closed RTL successfully
Libomptarget --> Unloading target library!
Target LEVEL0 RTL --> Target binary is a valid oneAPI OpenMP image.
Libomptarget --> Image 0x00000000004021d0 is compatible with RTL 0x0000000000cd6f20!
Libomptarget --> Unregistered image 0x00000000004021d0 from RTL 0x0000000000cd6f20!
Libomptarget --> Done unregistering images!
Libomptarget --> Removing translation table for descriptor 0x00000000004021b0
Libomptarget --> Done unregistering library!
Libomptarget --> Deinit target library!
Target LEVEL0 RTL --> Deinit Level0 plugin!
Target LEVEL0 RTL --> Memory usage for host memory, device 0:
Target LEVEL0 RTL --> -- Allocator:       Native,         Pool
Target LEVEL0 RTL --> -- Requested:            0,            0
Target LEVEL0 RTL --> -- Allocated:            0,            0
Target LEVEL0 RTL --> -- Freed    :            0,            0
Target LEVEL0 RTL --> -- InUse    :            0,            0
Target LEVEL0 RTL --> -- PeakUse  :            0,            0
Target LEVEL0 RTL --> -- NumAllocs:            0,            0
Target LEVEL0 RTL --> Memory usage for shared memory, device 0:
Target LEVEL0 RTL --> -- Allocator:       Native,         Pool
Target LEVEL0 RTL --> -- Requested:            0,            0
Target LEVEL0 RTL --> -- Allocated:            0,            0
Target LEVEL0 RTL --> -- Freed    :            0,            0
Target LEVEL0 RTL --> -- InUse    :            0,            0
Target LEVEL0 RTL --> -- PeakUse  :            0,            0
Target LEVEL0 RTL --> -- NumAllocs:            0,            0
Target LEVEL0 RTL --> Memory usage for shared memory, device 1:
Target LEVEL0 RTL --> -- Allocator:       Native,         Pool
Target LEVEL0 RTL --> -- Requested:            0,            0
Target LEVEL0 RTL --> -- Allocated:            0,            0
Target LEVEL0 RTL --> -- Freed    :            0,            0
Target LEVEL0 RTL --> -- InUse    :            0,            0
Target LEVEL0 RTL --> -- PeakUse  :            0,            0
Target LEVEL0 RTL --> -- NumAllocs:            0,            0
Target LEVEL0 RTL --> Memory usage for device memory, device 0:
Target LEVEL0 RTL --> -- Allocator:       Native,         Pool
Target LEVEL0 RTL --> -- Requested:       262144,       196608
Target LEVEL0 RTL --> -- Allocated:       262144,       196608
Target LEVEL0 RTL --> -- Freed    :       262144,       196608
Target LEVEL0 RTL --> -- InUse    :            0,            0
Target LEVEL0 RTL --> -- PeakUse  :       262144,       196608
Target LEVEL0 RTL --> -- NumAllocs:            1,            3
Target LEVEL0 RTL --> Memory usage for device memory, device 1:
Target LEVEL0 RTL --> -- Allocator:       Native,         Pool
Target LEVEL0 RTL --> -- Requested:            0,            0
Target LEVEL0 RTL --> -- Allocated:            0,            0
Target LEVEL0 RTL --> -- Freed    :            0,            0
Target LEVEL0 RTL --> -- InUse    :            0,            0
Target LEVEL0 RTL --> -- PeakUse  :            0,            0
Target LEVEL0 RTL --> -- NumAllocs:            0,            0
Target LEVEL0 RTL --> Closed RTL successfully
Target OPENCL RTL --> Deinit OpenCL plugin!
PASSED

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

ドキュメント	説明とリンク
OpenMP* 5.0 および 5.1 仕様 (PDF)	The OpenMP* API 5.0 仕様 (英語) および OpenMP* API 5.1 仕様 (英語) では、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
Yocto* プロジェクト向けのレイヤー (英語)	meta-intel レイヤーを使用して、oneAPI コンポーネントを Yocto* プロジェクト・ビルドに追加します。

法務上の注意書き

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

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

実際の費用と結果はシステム構成によって異なります。

© Intel Corporation. Intel、インテル、Intel ロゴ、その他のインテルの名称やロゴは、アメリカ合衆国および / またはその他の国における Intel Corporation またはその子会社の商標です。
* その他の社名、製品名などは、一般に各社の表示、商標または登録商標です。

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

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

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

製品および性能に関する情報

¹ 性能は、使用状況、構成、その他の要因によって異なります。詳細については、http://www.intel.com/PerformanceIndex/ (英語) を参照してください。

並列パフォーマンスを向上する最も一般的な手法－インテル® Xeon® プロセッサーおよびインテル® Xeon Phi™ コプロセッサー向けこの記事は、インテル® デベロッパー・ゾーンに掲載されている「Common Best Known Methods for Parallel Performance, from Intel® Xeon® to Intel® Xeon Phi™ Processors」 […]
Parallel Universe マガジン Parallel Universe へようこそ。米国インテル社が四半期に一度オンラインで公開しているオンラインマガジンです。インテルの技術者によるテクノロジーの解説や、最新ツールの紹介など、並列化に関する記事を毎号掲載しています。第1号からのバックナンバーを PDF 形式で用意しました、ぜひご覧ください。 12 […]
インテル® IPP における OpenMP* サポートの変更この記事は、インテル® ソフトウェア・ネットワークに掲載されている「OpenMP support changes in Intel Performance libraries […]
OpenMP* 4.x による新しいレベルの並列化 (全2回) 2015年 12月に東京で開催されたインテル® ソフトウェア・カンファレンス Day 1 で、iSUS 編集長すがわらきよふみが講演した内容を基に、2部にわけて「OpenMP* 4.x による新しいレベルの並列化」についてお届けします。Part 1 では、複数レベルの並列性を活用するための Code […]
データとコードの並べ替え: 最適化とメモリー – パート 2 この記事は、インテル® デベロッパー・ゾーンに公開されている「Putting Your Data and Code in Order: Data and layout - Part 2」の日本語参考訳です。この記事の PDF 版はこちらからご利用になれます。パフォーマンスとメモリーに関するこのシリーズ (全 2 […]