GEN アセンブリーの概要

この記事は、インテル® デベロッパー・ゾーンに公開されている「Introduction to GEN Assembly」の日本語参考訳です。


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

目次

はじめに

OpenCL* カーネルの最適化とデバッグには、アセンブリー・コードが非常に役立つことがあります。この記事では、インテル® SDK for OpenCL* Applications で利用可能なツールについて説明します。個々のカーネル用にオフライン・コンパイラーにより生成されたアセンブリーの表示、OpenCL* C コードに対応するアセンブリー・コードの領域のハイライト、生成されたアセンブリーの異なる部分の高レベルの説明などを行います。レジスター領域の構文の概要と意味を説明した後、異なる種類のレジスター、利用可能なアセンブリー命令とこれらの命令で操作できるデータ型を紹介します。この記事をお読みになったら、あとは作業を開始するだけです。この記事の続編では、インテル® VTune™ Amplifier を使用したアセンブリーのプロファイルおよびアセンブリーのデバッグを説明する予定です。

単純な OpenCL* カーネルのアセンブリー

まず単純なカーネルで始めましょう。

kernel void empty() { 
}

図 1. 単純な OpenCL* カーネル

これは最も単純なカーネルと言えます。このカーネルを Code Builder Session Explorer でビルドしましょう。CODE-BUILDER/OpenCL Kernel Development/New Session を選択して新しいセッションを作成し、上記のカーネルを空の program.cl ファイルにコピーして、このファイルをビルドします。第 5 世代インテル® プロセッサー (Broadwell) または第 6 世代インテル® プロセッサー (Skylake) を利用していれば、program_empty.gen ファイルが生成されているはずです。このファイルをダブルクリックします。次のような画面が表示されます。

図 2. empty カーネルと対応するアセンブリー

カーネルのアセンブリーは右に表示されています。いくつか注釈を追加します。

// Start of Thread
LABEL0
(W)      and      (1|M0)        r2.6<1>:ud    r0.5<0;1,0>:ud    0x1FF:ud         // id:

// End of thread
(W)      mov      (8|M0)        r127.0<1>:ud  r0.0<8;8,1>:ud   {Compacted}                 // id:
         send     (8|M0)        null          r127              0x27      0x2000010 {EOT}  // id:

図 3. empty カーネルの注釈付きアセンブリー

ほとんど何も含まれていませんが、これはあくまでもスタートです。

では、少しだけ複雑にしてみましょう。次のコードを program.cl にコピーしてください。

kernel void meaning_of_life(global uchar* out) 
{ 
 out[31] = 42;
}

図 4. meaning_of_life カーネル

ファイルをリビルドすると、program_meaning_of_life.gen ファイルが生成されます。このファイルをダブルクリックすると、少しだけ複雑になったカーネルが表示されます。

図 5. meaning_of_life カーネルと対応するアセンブリー

左のカーネルの異なる部分をクリックすると、アセンブリーの異なる部分がハイライトされることを確認します。

次はカーネルの最初と対応する命令です。

図 6. meaning_of_life カーネルの最初

カーネルの本体と対応する命令。

図 7. meaning_of_life カーネルの本体

カーネルの最後と対応する命令。

図 8. meaning_of_life カーネルの最後

分かりやすくなるように、アセンブリーを再配置しましょう。

// Start of Thread
LABEL0
(W)      and      (1|M0)        r2.6<1>:ud    r0.5<0;1,0>:ud    0x1FF:ud         // id:
// r3 and r4 contain the address of out variable (8 unsigned quadwords – uq)
// we are going to place them in r1 and r2
(W)      mov      (8|M0)        r1.0<1>:uq    r3.0<0;1,0>:uq                   // id:


// Move 42 (0x2A:ud – ud is unsigned dword) into 32 slots (our kernel is compiled SIMD32)
// We are going to use registers r7, r10, r13 and r16, each register fitting 8 values
         mov      (8|M0)        r7.0<1>:ud    0x2A:ud          {Compacted}                 // id:
         mov      (8|M8)        r10.0<1>:ud   0x2A:ud          {Compacted}                 // id:
         mov      (8|M16)       r13.0<1>:ud   0x2A:ud                          // id:
         mov      (8|M24)       r16.0<1>:ud   0x2A:ud                          // id:

// Add 31 (0x1F:ud) to eight quadwords in r1 and r2 and place the results in r3 and r4
// Essentially, we get &out[31]
 (W)      add      (8|M0)        r3.0<1>:q     r1.0<0;1,0>:q     0x1F:ud          // id:

// Now we spread &out[31] into r5,r6, r8,r9, r11, r12, and r14, r15 – 32 values in all.
         mov      (8|M0)        r5.0<1>:uq    r3.0<0;1,0>:uq                   // id:
         mov      (8|M8)        r8.0<1>:uq    r3.0<0;1,0>:uq                   // id:1
         mov      (8|M16)       r11.0<1>:uq   r3.0<0;1,0>:uq                   // id:1
         mov      (8|M24)       r14.0<1>:uq   r3.0<0;1,0>:uq                   // id:1

// Write to values in r7 into addresses in r5, r6, etc.
         send     (8|M0)        null          r5                0xC       0x60680FF                 // id:1
         send     (8|M8)        null          r8                0xC       0x60680FF                 // id:1
         send     (8|M16)       null          r11               0xC       0x60680FF                 // id:1
         send     (8|M24)       null          r14               0xC       0x60680FF                 // id:1

// End of thread
(W)      mov      (8|M0)        r127.0<1>:ud  r0.0<8;8,1>:ud   {Compacted}                 // id:
         send     (8|M0)        null          r127              0x27      0x2000010 {EOT}                 // id:1

図 9. meaning_of_life カーネルの注釈付きアセンブリー

固定インデックスの代わりに get_global_id(0) を使用して、カーネルをもう少し複雑にしてみましょう。

kernel void meaning_of_life2(global uchar* out) 
{ 
 int i = get_global_id(0);
 out[i] = 42;
}

図 10. meaning_of_life2 カーネル

get_global_id(0) を追加したことでカーネルのサイズが増加している (アセンブリー命令が 9 つ増えている) ことに注意してください。スレッドの各ワークアイテムの増加アドレスを計算する必要があります (32 のワークアイテムがあります)。

// Start of Thread
LABEL0
(W)      and      (1|M0)        r7.6<1>:ud    r0.5<0;1,0>:ud    0x1FF:ud         // id:

// Move 42 (0x2A:ud – ud is unsigned dword) into 32 slots (our kernel is compiled SIMD32)
// We are going to use registers r17, r20, r23 and r26, each register fitting 8 values
         mov      (8|M0)        r17.0<1>:ud   0x2A:ud          {Compacted}                 // id:
         mov      (8|M8)        r20.0<1>:ud   0x2A:ud          {Compacted}                 // id:
         mov      (8|M16)       r23.0<1>:ud   0x2A:ud                          // id:
         mov      (8|M24)       r26.0<1>:ud   0x2A:ud                          // id:
// get_global_id(0) calculation, r0.1, r7.0 and r7.3 will contain the necessary starting values
(W)      mul      (1|M0)        r3.0<1>:ud    r0.1<0;1,0>:ud    r7.3<0;1,0>:ud   // id:
(W)      mul      (1|M0)        r5.0<1>:ud    r0.1<0;1,0>:ud    r7.3<0;1,0>:ud   // id:
(W)      add      (1|M0)        r3.0<1>:ud    r3.0<0;1,0>:ud    r7.0<0;1,0>:ud   {Compacted} // id:
(W)      add      (1|M0)        r5.0<1>:ud    r5.0<0;1,0>:ud    r7.0<0;1,0>:ud   {Compacted} // id:1
// r3 thru r6 will contain the get_global_id(0) offsets; r1 and r2 contain 32 increasing values
         add      (16|M0)       r3.0<1>:ud    r3.0<0;1,0>:ud    r1.0<8;8,1>:uw   // id:1
         add      (16|M16)      r5.0<1>:ud    r5.0<0;1,0>:ud    r2.0<8;8,1>:uw   // id:1
// r8 and r9 contain the address of out variable (8 unsigned quadwords – uq)
// we are going to place these addresses in r1 and r2
 (W)      mov      (8|M0)        r1.0<1>:uq    r8.0<0;1,0>:uq                   // id:1

// Move the offsets in r3 thru r6 to r7, r8, r9, r10, r11, r12, r13, r14
         mov      (8|M0)        r7.0<1>:q     r3.0<8;8,1>:d                    // id:1
         mov      (8|M8)        r9.0<1>:q     r4.0<8;8,1>:d                    // id:1
         mov      (8|M16)       r11.0<1>:q    r5.0<8;8,1>:d                    // id:1
         mov      (8|M24)       r13.0<1>:q    r6.0<8;8,1>:d                    // id:1

// Add the offsets to address of out in r1 and place them in r15, r16, r18, r19, r21, r22, r24, r25
         add      (8|M0)        r15.0<1>:q    r1.0<0;1,0>:q     r7.0<4;4,1>:q    // id:1
         add      (8|M8)        r18.0<1>:q    r1.0<0;1,0>:q     r9.0<4;4,1>:q    // id:1
         add      (8|M16)       r21.0<1>:q    r1.0<0;1,0>:q     r11.0<4;4,1>:q   // id:2
         add      (8|M24)       r24.0<1>:q    r1.0<0;1,0>:q     r13.0<4;4,1>:q   // id:2

// write into addresses in r15, r16, values in r17, etc.
         send     (8|M0)        null          r15               0xC       0x60680FF                 // id:2
         send     (8|M8)        null          r18               0xC       0x60680FF                 // id:2
         send     (8|M16)       null          r21               0xC       0x60680FF                 // id:2
         send     (8|M24)       null          r24               0xC       0x60680FF                 // id:2

// End of thread
(W)      mov      (8|M0)        r127.0<1>:ud  r0.0<8;8,1>:ud   {Compacted}                 // id:
         send     (8|M0)        null          r127              0x27      0x2000010 {EOT}                 // id:2

図 11. meaning_of_life2 カーネルの注釈付きアセンブリー

最後に、読み書きと計算を行うカーネルを見てみましょう。

kernel void modulate(global float* in, global float* out) { 
 int i = get_global_id(0);

 float f = in[i];
 float temp = 0.5f * f;
 out[i] = temp;
}

図 12. 浮動小数点演算を行う単純なカーネル

コードは次のようになります (分かりやすくなるように一部のアセンブリー命令を再配置していることに注意してください)。

// Start of Thread
LABEL0
(W)      and      (1|M0)        r7.6<1>:ud    r0.5<0;1,0>:ud    0x1FF:ud         // id:

// r3 and r4 will contain the address of out buffer
(W)      mov      (8|M0)        r3.0<1>:uq    r8.1<0;1,0>:uq                     // id:
// int i = get_global_id(0);
(W)      mul      (1|M0)        r5.0<1>:ud    r0.1<0;1,0>:ud    r7.3<0;1,0>:ud   // id:
(W)      mul      (1|M0)        r9.0<1>:ud    r0.1<0;1,0>:ud    r7.3<0;1,0>:ud   // id:
(W)      add      (1|M0)        r5.0<1>:ud    r5.0<0;1,0>:ud    r7.0<0;1,0>:ud   {Compacted} // id:
(W)      add      (1|M0)        r9.0<1>:ud    r9.0<0;1,0>:ud    r7.0<0;1,0>:ud   {Compacted} // id:
         add      (16|M0)       r5.0<1>:ud    r5.0<0;1,0>:ud    r1.0<8;8,1>:uw   // id:
         add      (16|M16)      r9.0<1>:ud    r9.0<0;1,0>:ud    r2.0<8;8,1>:uw   // id:

// r1 and r2 will contain the address of in buffer
(W)      mov      (8|M0)        r1.0<1>:uq    r8.0<0;1,0>:uq                   // id:1
// r11, r12, r13, r14, r15, r16, r17 and r18 will contain 32 qword offsets 
         mov      (8|M0)        r11.0<1>:q    r5.0<8;8,1>:d                    // id:1
         mov      (8|M8)        r13.0<1>:q    r6.0<8;8,1>:d                    // id:1
         mov      (8|M16)       r15.0<1>:q    r9.0<8;8,1>:d                    // id:1
         mov      (8|M24)       r17.0<1>:q    r10.0<8;8,1>:d                   // id:1

//  float f = in[i];
         shl      (8|M0)        r31.0<1>:uq   r11.0<4;4,1>:uq   0x2:ud           // id:1
         shl      (8|M8)        r33.0<1>:uq   r13.0<4;4,1>:uq   0x2:ud           // id:1
         shl      (8|M16)       r35.0<1>:uq   r15.0<4;4,1>:uq   0x2:ud           // id:1
         shl      (8|M24)       r37.0<1>:uq   r17.0<4;4,1>:uq   0x2:ud           // id:1
         add      (8|M0)        r19.0<1>:q    r1.0<0;1,0>:q     r31.0<4;4,1>:q   // id:1
         add      (8|M8)        r21.0<1>:q    r1.0<0;1,0>:q     r33.0<4;4,1>:q   // id:2
         add      (8|M16)       r23.0<1>:q    r1.0<0;1,0>:q     r35.0<4;4,1>:q   // id:2
         add      (8|M24)       r25.0<1>:q    r1.0<0;1,0>:q     r37.0<4;4,1>:q   // id:2
// read in f values at addresses in r19, r20, r21, r22, r23, r24, r25, r26 into r27, r28, r29, r30 
         send     (8|M0)        r27           r19               0xC       0x4146EFF                 // id:2
         send     (8|M8)        r28           r21               0xC       0x4146EFF                 // id:2
         send     (8|M16)       r29           r23               0xC       0x4146EFF                 // id:2
         send     (8|M24)       r30           r25               0xC       0x4146EFF                 // id:2

// float temp = 0.5f * f; - 0.5f is 0x3F000000:f
//     We multiply 16 values in r27, r28 by 0.5f and place them in r39, r40
//     We multiple 16 values in r29, r30 by 0.5f and place them in r47, r48
         mul      (16|M0)       r39.0<1>:f    r27.0<8;8,1>:f    0x3F000000:f     // id:3
         mul      (16|M16)      r47.0<1>:f    r29.0<8;8,1>:f    0x3F000000:f     // id:3

//     out[i] = temp;
         add      (8|M0)        r41.0<1>:q    r3.0<0;1,0>:q     r31.0<4;4,1>:q   // id:2
         add      (8|M8)        r44.0<1>:q    r3.0<0;1,0>:q     r33.0<4;4,1>:q   // id:2
         add      (8|M16)       r49.0<1>:q    r3.0<0;1,0>:q     r35.0<4;4,1>:q   // id:2
         add      (8|M24)       r52.0<1>:q    r3.0<0;1,0>:q     r37.0<4;4,1>:q   // id:3

         mov      (8|M0)        r43.0<1>:ud   r39.0<8;8,1>:ud  {Compacted}                 // id:3
         mov      (8|M8)        r46.0<1>:ud   r40.0<8;8,1>:ud  {Compacted}                 // id:3
         mov      (8|M16)       r51.0<1>:ud   r47.0<8;8,1>:ud                  // id:3
         mov      (8|M24)       r54.0<1>:ud   r48.0<8;8,1>:ud                  // id:3

// write into addresses r41, r42 the values in r43, etc.
         send     (8|M0)        null          r41               0xC       0x6066EFF                 // id:3
         send     (8|M8)        null          r44               0xC       0x6066EFF                 // id:3
         send     (8|M16)       null          r49               0xC       0x6066EFF                 // id:3
         send     (8|M24)       null          r52               0xC       0x6066EFF                 // id:4

// End of thread
(W)      mov      (8|M0)        r127.0<1>:ud  r0.0<8;8,1>:ud   {Compacted}                 // id:
         send     (8|M0)        null          r127              0x27      0x2000010 {EOT}                 // id:4

図 13. 単純な浮動小数点演算カーネルの注釈付きアセンブリー

アセンブリー命令の読み方

すべての命令は次のような形式になります。

[(pred)] opcode (exec-size|exec-offset) dst src0 [src1] [src2]

(pred) は、オプションのプレディケートです。ここでは説明しません。

opcode は、add や mov のような命令のシンボルです (opcode のリストは最後の表を参照)。

exec-size は、命令の SIMD 幅です。アーキテクチャーに応じて 1、2、4、8 または 16 になります。SIMD32 コンパイルでは、通常、実行サイズ 8 または 16 の 2 つの命令が 1 つのグループになります。

exec-offset は、ARF レジスターのどの部分を読み書きするか、実行ユニットに知らせる部分です。例えば、(8|M24) は実行マスクのビット 24 から 31 を参照します。SIMD16 または SIMD32 のコードは次のようになります。

         mov  (8|M0)   r11.0<1>:q   r5.0<8;8,1>:d   // id:1
         mov  (8|M8)   r13.0<1>:q   r6.0<8;8,1>:d   // id:1
         mov  (8|M16)  r15.0<1>:q   r9.0<8;8,1>:d   // id:1
         mov  (8|M24)  r17.0<1>:q   r10.0<8;8,1>:d  // id:1

図 14. SIMD32 アセンブリーの mov 命令

GRF のオペランドごとにアクセスできるバイト数の制限により、コンパイラーは 4 つの 8 要素幅の演算を行う必要があります。

dst は、デスティネーション・レジスターです。

src0 は、ソースレジスターです。

src1 は、オプションのソースレジスターです。0x3F000000:f (0.5) や 0x2A:ud (42) のように、即値の場合もあることに注意してください。

src2 は、オプションのソースレジスターです。

ジェネラル・レジスター・ファイル (GRF) レジスター

各スレッドには 128 個のレジスターの専用スペース (r0 から r127) があります。各レジスターは 256 ビットまたは 32 バイトです。

アーキテクチャー・レジスター・ファイル (ARF) レジスター

上記のアセンブリー・コードには、これらの特殊なレジスターの 1 つである null レジスターのみ含まれています。null レジスターは、send 命令のデスティネーションとして使用され、スレッドの最後を記述および示すために使用されます。その他のアーキテクチャー・レジスターを次に示します。

図 15. アーキテクチャー・レジスター・ファイル (ARF) レジスター

レジスターは 32 バイト幅でバイトアドレス指定可能であるため、アセンブリーにはこれらのレジスターに格納された値にアクセスできるレジスター領域構文が用意されています。

この後、一連の図でレジスター領域構文について説明します。

例として、レジスター領域 r4.1<16;8,2>:w を見てみましょう。領域の最後の w はワード (または 2 バイト) 値であることを示しています。利用可能な整数データ型と浮動小数点データ型の一覧は、後の表を参照してください。起点は r4.1 で、レジスター r4 の 2 つ目のワードから始まることを意味します。縦ストライドは 16 で、2 つ目の行から始まるため 16 の要素をスキップする必要があることを意味します。幅パラメーターは 8 で、行の要素数を指します。横ストライドは 2 で、2 つおきに要素を処理することを意味します。ここでは r4 と r5 の両方の内容を参照していることに注意してください。次の図は結果をまとめたものです。

図 16. 16 要素のレジスター領域 (r4.1<16;8,2>:w) の例

この例の、レジスター領域 r5.0<1;8,2>:w について考えてみましょう。領域は r5 の最初の要素から始まります。1 行に 8 つの要素があり、2 つおきの要素を含みます。つまり、最初の行は {0, 1, 2, 3, 4, 5, 6, 7} です。2 つ目の行は 1 ワードのオフセットまたは r5.2 で始まり、{8, 9, 10, 11, 12, 13, 14, 15} を含みます。次の図は結果をまとめたものです。

図 17. 行がインターリーブされた 16 要素のレジスター領域 (r5.0<1;8,2>:w)

次のアセンブリー命令について考えてみましょう。

add(16|M0) r6.0<1>:w r1.7<16;8,1>:b r2.1<16;8,1>:b

src0 は、r1.7 から始まり 8 つの連続するバイトの最初の行、r1.23 から始まる 2 つ目の行が続きます。

src1 は、r2.1 から始まり 8 つの連続するバイトの最初の行、r2.17 から始まる 2 つ目の行が続きます。

dst は、r6.0 から始まり、値をワードで格納します。命令 Add(16) は 16 の値を操作するため、16 の連続するワードを r6 に格納します。

図 18. 直接レジスター・アドレッシングの領域説明の例

次のアセンブリー命令について考えてみましょう。

add(16|M0) r6.0<1>:w r1.14<16;8,0>:b r2.17<16;8,1>:b

src0 は r1.14<16;8,0>:b です。最初のバイトサイズの値は r1.14 から始まります。ストライド値 0 は領域の幅 (8) で値を繰り返します。次に、r1.30 から始まる領域が続き、そこに 8 回格納された値を繰り返します。つまり、値 {1,1,1,1,1,1,1,1, 4, 4, 4, 4, 4, 4, 4, 4} を操作します。

src1 は r2.17<16;8,1>:b で、最初の行は r2.17 から始まる 8 バイト、2 つ目の行は r3.1 から始まる 8 バイトです。

図 19. src0 を複製スカラーのベクトルとして含む直接レジスター・アドレッシングの領域説明の例

レジスター領域の : の後の文字は、格納されるデータ型を示します。利用可能な整数データ型と浮動小数点データ型の 2 つの表を次に示します。

図 20. 実行ユニット (整数データ型)

図 21. 実行ユニット (浮動小数点データ型)

次の表は、利用可能なアセンブリー命令をまとめたものです。

図 22. 利用可能な GEN アセンブリー命令

参考文献 (英語)

インテル® グラフィックス・ドキュメントの第 7 巻:

インテル® グラフィックス・ドキュメントのセット:
https://01.org/linuxgraphics/documentation/hardware-specification-prms

著者紹介

Robert Ioffe は、インテル コーポレーションのソフトウェア & ソリューション・グループのテクニカル・コンサルティング・エンジニアです。OpenCL* プログラミングとインテル® Iris™ グラフィックスまたはインテル® Iris™ Pro グラフィックスにおける OpenCL* ワークロードの最適化のエキスパートで、インテル® グラフィックス・ハードウェアを熟知しています。Khronos* の標準化作業に深くかかわっており、これまでに最新機能のプロトタイプの作成やインテル® アーキテクチャーでの動作検証を行ってきました。最近では、OpenCL* 2.0 の入れ子構造の並列処理 (enqueue_kernel functions) 機能のプロトタイプの作成に取り組み、OpenCL* 2.0 用の GPU クイックソートを含む、いくつかの入れ子構造の並列処理サンプルコードを作成しました。また、以下の単純な OpenCL* カーネルの最適化に関する動画 2 つを公開しています。現在、入れ子の並列処理に関する 3 つ目の動画を作成中です。

以下の記事も参照してください。
GPU-Quicksort in OpenCL 2.0: Nested Parallelism and Work-Group Scan Functions
Sierpiński Carpet in OpenCL 2.0
Optimizing Simple OpenCL Kernels: Modulate Kernel Optimization
Optimizing Simple OpenCL Kernels: Sobel Kernel Optimization

コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください。

関連記事

  • インテル® SDK for OpenCL* Applications 入門インテル® SDK for OpenCL* Applications 入門 この記事はインテル® デベロッパー・ゾーンに公開されている「Getting Started with Intel® SDK for OpenCL™ Applications」の参考訳です。 このガイドは、以下で入手可能なインテル® SDK for OpenCL* Applications […]
  • Intel® OpenCL SDK 1.1 の導入Intel® OpenCL SDK 1.1 の導入 インテルより初めての製品版となる Intel® OpenCL SDK 1.1 が 2011年 6月にリリースされました。この記事では、すでに OpenCL プログラムを開発している開発者、あるいは OpenCL について多少の知識を持つユーザーを対象に、Intel® OpenCL SDK 1.1 […]
  • opencl_program と引数のバインド – パート 4opencl_program と引数のバインド – パート 4 この記事は、インテル® デベロッパー・ゾーンに公開されている「opencl_program and argument binding」の日本語参考訳です。 この記事は、インテル® スレッディング・ビルディング・ブロック (インテル® TBB) 4.4 Update 2 以降で利用可能な新しいノード opencl_node […]
  • インテル® Xeon Phi™ コプロセッサー x100 が搭載されたシステムでインテル® Code Builder for OpenCL* API を実行するインテル® Xeon Phi™ コプロセッサー x100 が搭載されたシステムでインテル® Code Builder for OpenCL* API を実行する この記事は、インテル® デベロッパー・ゾーンに公開されている「Working with Intel® Code Builder for OpenCL™ API on Systems Equipped with Intel® Xeon Phi™ x100 Coprocessor」の日本語参考訳です。 OpenCL* […]
  • インテル® SDK for OpenCL™ Applications 2014 ベータ版インテル® SDK for OpenCL™ Applications 2014 ベータ版 インテル® SDK for OpenCL™ Applications 2014 は、OpenCL API 向けの総合開発プラットフォームです。 Windows* と新たに Android* OS 向けに OpenCL をサポート OpenCL 1.2 の完全なプロファイルをサポート インテル® Iris™ […]