戻る   次へ

2.10. Autovectorization

Intel OpenCL SDKやMax OS XはAutovectorizationという、自動ベクトル化、すなわちSIMD命令コードに自動で翻訳してくれる機能が附属しています。

複雑なロジックをベクトル化するのは、手間がかかりバグの余地も大きくるため、最近はOpenCLの実装ライブラリを提供する大半のベンダーが、自動ベクトル化をサポートしているので、2010〜2012年頃のベクトルだらけのコードベースは可読性をおとします。

自動ベクトル化がサポートされている場合はベクトル型を使わないという選択肢があることを念頭におくようにしてください。

2.10.1. 明示的ベクトル化

明示的ベクトル化は「Explicit Autovectorization」の直訳です。明示的の意味をあえて解説するならば、OpenCLのベクトル型で定義したベクトルにすることです。

これにより各ベクトルの要素がSSE/AVXのSIMDの幅で4または8個の32ビットデータ型の場合は、1実行サイクルで並列に処理します。Intelの内蔵GPUではfloat4やint4/uint4が推奨されています。

2.10.2. 暗黙的ベクトル化

SIMDを活用する場合にAuto-Vectorizationが明示的となると毎回、float8にアルゴリズムを改変するように迫られることは開発者にとって大きな負担となります。

暗黙的ベクトル化は、4または8要素のfloatのベクトル型に修正せずとも、ワークアイテムをSSE/AVX等のSIMDチャンネルに自動マッピングするために、ロックステップでのSIMD処理をOpenCL実装がバックグラウンドで行ないます。

SIMDの演算器のビット数（32bit）に適合したデータ型、特にfloat型、int型がワークアイテムにマップされやすいとされます。

暗黙的ベクトル化は複雑なアルゴリズムになればなるほど重宝する機能であり、むしろベクトルに無理にでも変換しようとする強引な手法を減らす良い契機になるかもしれません。

Mac OS X Autovectorizer

本書の検証環境はMac OS XのためAutovectorizerという機能を使います。下表のように、cl-auto-vectorize-enableが真の値になってる必要があります。既定の設定は「自動ベクトル化有効」となります。

設定

型

既定値

コマンドラインフラグ

Auto-vectorizer

Boolean

YES

 cl-auto-vectorize-enable
cl-auto-vectorize-disable

次にMac OS XのOpenCLフレームワークでは、自動ベクトル化を阻むのでやらない方がいい３つの非推奨設定・コーディングがあります。以下がやるべきでない事です。

  • ワークアイテムのIDに依存した制御フロー
  • 特定のデバイスへの最適化
  • 配列の要素にアクセスするときは、連続するワークアイテムIDが、連続する配列要素にアクセスする（アクセスパターン、例：行列乗算の際に、行列を転置）
  • メモリーへのアクセスを制御フロー外に移動

最後の点は以下のようなコードを指します。

if(条件)
    a[id] = 10;
else
    a[id] = 50;

まずこうした分岐の入る制御フローがあるので、可能な限り避けます。もし制御フローが必要な場合は、if文からメモリーへのアクセスを外に移動させてしまいます。

if(条件)
    tmp = 10;
else
    tmp = 50;

a[id] = tmp;

Mac OS Xではスカラ型の浮動小数点演算をコードした場合、CPUではコアの一部分のみが処理をして、他はアイドル状態となるが、自動ベクトル化によってこの問題を解決（コアの稼働率を上げることが）できるとしています。

Intel Implicit Autovectorization Module

Intelの暗黙的自動ベクトル化モジュールは原則として4バイト（32ビット）なので、floatやintのスカラ型を処理する限りカーネルは自動でベクトル化されます。

32bit長以外のベクトルを使う場合はベクトル型の属性をコンパイラに知らせておくことで、自動ベクトル化をコンパイラがやりやすくするための助けとします。

__attribute__((vec_type_hint(<typen>)))

typenのデフォルトはint型です。例えば、typenをfloat4とする場合は、カーネルで行なう処理の大半が、float4でベクトル化されることをコンパイラに伝えます。Intel CPUのAVXであれば256ビット長のレジスタがあるので、float型が8つおさまりますが、float4が2つ並ぶようにしたい場合には以下のように宣言します。

__kernel __attribute__((vec_type_hint(float4)))
void foo( __global float4 *p ) {
}

基本演算幅が、intの32ビットでなく、128bitということをコンパイラに知らせることにより、自動ベクトル化の際の幅の決定がやりやすくなります。

Intelが推奨する基本ガイドラインには以下のようなものがあります。

  • メモリーアラインメントを調整する。
  • ポインタエイリアシング（画像処理のエイリアシングではなく、メモリ操作の一種）を避けること。restrict修飾子を使用する。
  • 分岐の入る制御フローを避ける。
  • forループの使用を避ける。
  • 配列構造体（SOA、Structure of Array）をやめて、構造体配列(AOS, Array of Structure)にする。

最後の点について補足です。SOAは以下のように、ベクトル型の配列を指します。

__kernel void mul(__global float4* input, __global float* output)
{
    size_t gid  = get_global_id(0);
    output[gid] = input[gid].x * input[gid].y * input[gid].z * input[gid].w;
}

これをAOSにするには、単純にベクトル型をアンパックしてスカラ型に変更します。

__kernel void mul(
    __global float* x,
    __global float* y,
    __global float* z,
    __global float* w,
    __global float* output)
{
    size_t gid  = get_global_id(0);
    output[gid] = x[gid] * y[gid] * z[gid] * w[gid];
}

SOAの方が明らかにエレガントなコードではあるのですが、パフォーマンス上ではペナルティを受ける可能性があります。このように処理するデータによっては最適化によってコードが見苦しくなるケースもあります。

戻る 上に戻る 次へ
 ホーム 

Copyright 2018-2019, by Masaki Komatsu