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

2020年7月7日

この記事は、インテル® デベロッパー・ゾーンに公開されている「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 向けの nextgen コンパイラーでサポートされる 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 向けの nextgen コンパイラーでサポートされる 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 (PVC ターゲットを含む) サポート

OpenMP* オフロードランタイムは、Pixel Visual Core* (PVC) 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 以上のスレッドを実行できます。

チーム内のスレッド間で PVC ハードウェア・バリアの使用が許可されます。
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 レガシー・アプリケーションの PVC をターゲットにするオプションのサポート (以前の調査に基づきます)

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

-fiopenmp
-fopenmp-targets=spir64

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

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

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 <stdio.h>

#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<N; k++) {
    a[k] = k;
    b[k] = k + 1;
    result = result + a[k] + b[k];  
  }
  printf("GPU version was called. ");
  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<N; k++) {
    a[k] = k;
    b[k] = k + 1;
    result = result + a[k] + b[k];
  }
  printf("CPU version was called. ");
  return result;
}

int main() {
  float result=0.0;

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

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

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

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

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

再帰関数呼び出し (コンパイル時の定数式を除く)
プレースメント以外の new と delete
goto 文の制限
register と thread_local ストレージ修飾子
仮想関数修飾子
関数ポインター (コンパイル時の定数式を除く)
仮想関数
例外処理
C+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 オフロードを起動するため、icc -qnextgen、icpc -qnextgen、または ifort -qnextgen コンパイル・オプションを使用してソースコードをコンパイルします。

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

または

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

または

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

実行: 環境変数を設定します: 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/ATS/PVC でサポートされる数学関数のリストです。

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 を設定します。結果は次のようになります。

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

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

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

さらに詳しく

ドキュメント	説明とリンク
OpenMP* 5.0 仕様 (PDF)	The 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 またはその子会社の商標です。

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

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

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

注意事項の改訂 #20110804

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