OpenCL Cについて書く前に知っておくべきこと

書かれたとおり	float4 val = (0, 0, 0, 0);
著者が書きたかったこと	float4 val = (float4)(0, 0, 0, 0);
書き方	float4 val = 0;

OpenCLに出くわしたり、実行する予定で、最初と2番目のオプションの違いがわからない場合、3番目のオプションで「コンパイルできますか？」-catを歓迎します。言語のニュアンスはたくさんあり、APIとパフォーマンスの最適化についてはまったくありません。

最も高価なコンピューターメモリはプログラマの頭の中にあります。 これがおそらく、2つの最も人気のあるGPUプログラミングテクノロジーであるCUDAとOpenCLが、特定のアーキテクチャ用のアセンブラーではなく、根本的に新しい並列言語に基づいているのではなく、適応C ++とCに基づいている理由です.OpenCLの場合、最も一般的なカーネル記述言語はOpenCL C 1.2です- ISO C99に基づくC方言。 標準ライブラリが削除され、それらに対するベクトル型と操作が追加されました。他のフロー間での位置をロックおよび決定するためのいくつかの関数、4つのアドレス空間。 C11と互換性のない単純なアトミック操作が追加されました（C11からのアトミック操作とロックは、まだ普及していないOpenCL C 2.0に追加されました）。 標準化された組み込み関数のような、Cにはない便利な関数をいくつか追加しました。

言語とAPIに関する多くの例がありますが、そのほとんどは2つのベクトルの追加です。 乾燥しているとはいえ、優れた公式仕様 、いくつかの英語の書籍、およびデバイスメーカーからの最適化のヒントがあります。 プログラマーは自分のタスクの書き方を理解するとすぐに、 get_global_id()ような新しい関数を使用して使い慣れたC99に書き込み、言語に関するすべてが明確で原始的なように見えます。 このような馴染みのあるC99で、簡単にトラップに陥り、当面は気付かないことがあります。 はい、OpenCL CはCに非常に似ていますが、C99には類似物がないため忘れられないほど忘れがちな非常に有用な違いがあります。

私は多くのOpenCL Cコードを見ましたが、書き始めたばかりの人も同じ間違いを犯します。

• ベクトルリテラルと型キャストを混同します。
• すばらしい型変換メカニズムを使用しないでください
• ベクトル型の変換の微妙な違いを忘れてください。

リストからわかるように、それはすべて型変換に関するものです。 OpenCL 1.2仕様では、これらはセクション6.2。* Conversions and Type Castingです。 さらに、次のセクション6.3演算子は cで、 誰も読みません 。 経験が示すように、仕様の多くは明確に記述されておらず、あまりにも退屈です。これらのトピックに関するアクセス可能なロシア語のドキュメントのギャップをこの記事で埋めようとします。

ベクトルリテラルまたは明示的なキャスト

OpenCL Cの新しいデザインは、ベクターの値を設定できるベクターリテラルです。 残念ながら、その構文は明示的な型キャストに非常に似ています。

 ( )(   ) 

例えば

 (int2)(1,2); 

または

 //      int2 a = (int2)(1, 2); // a=[1,2] //       int2 b = (int2)(3, 4); // b=[3,4] //          int4 c = (int4)(a, b); // c=[1,2,3,4] //         int3 d = (int3)(1, c.xy); // d=[1,1,2] //       ?! float2 e = (float2)(1); // e=[1.0f,1.0f] 

ただし、 (float2)(1)および上記の他の例は型キャストではなく、新しいデザインです（ OpenCL 1.2仕様の 6.1.6 Vector Literalsを参照）。

2番目の括弧内には、最初の括弧内のベクトル型と同じ数のベクトルのスカラーまたはコンポーネントの合計があるはずです。 例外が1つあります-右側に括弧で囲まれたスカラー値が1つしかない場合、それ自体が必要な数のベクトルコンポーネントに「乗算」されます。

Cスタイルでのベクトル型の明示的なキャストは、単に言語ではありません。 ベクトルリテラルではなく、「目を閉じた」状態で「なじみのある」型変換を見ると、致命的な間違いを犯す可能性があります。 次に、先頭の括弧内の型を削除できます。「結局、既にコンパイルされています。なぜ不要な型変換なのでしょうか？ すでに暗黙的に導かれています。」

実際の例：

 int2 coords = (get_global_id(0), get_global_id(1)); 

coordsはベクトルリテラルによって設定されません。ベクトルリテラルの場合は、ベクトルタイプを追加する必要がありました。

 int2 coords = (int2)(get_global_id(0), get_global_id(1)); 

しかし、次のようになりました： (get_global_id(0), get_global_id(1))そしてこれは通常のCからの構築です-括弧内で、演算子「」（コンマ）を使用して2つの関数を呼び出します。つまり、両方の関数が実行され、式が結果を返します2番目の関数、私たちが書いたかのように：

 get_global_id(0); int2 coords = get_global_id(1); 

スカラーのベクトルへの暗黙的な変換（さらに少し）が機能し[get_global_id(1), get_global_id(1)] [get_global_id(0), get_global_id(1)]ではなく、ベクトル[get_global_id(1), get_global_id(1)]があります。

幸いなことに、単純なケースでは、コンパイラは「 warning: expression result unused 」などの警告を発行する場合がありますが、これに頼るべきではありません。

このようなコードは、正しく機能しないため、すぐに見つけることができます。 ただし、次の例は、色が灰色である限り機能します。 色を変更したいときは、何らかの理由で、まだ灰色になります。

 //  ,    float3 color = (0.5f, 0.5f, 0.5f); // color=[0.5f, 0.5f, 0.5f] 

コードは機能しており、プロジェクトが送信されました。 そして、突然、小さな変更が必要になりました-グレーの色を濃紺にするためです。

 //  ,    float3 color = (0.1f, 0.1f, 0.5f); // color=[0.5f, 0.5f, 0.5f] 

ベクトルリテラルを使用する必要がありました。

 //   float3 color = (float3)(0.1f,0.1f, 0.5f); // color=[0.1f,0.1f, 0.5f] 

ブール値からベクトルへの変換

 int val = true; int2 val2 = true; 

valの意味は何ですか？ val2はどれがありますか？

スカラーの場合、 bool値を変換するときにISO C99ルールが適用され（ bool型と定数trueおよびfalseはC99およびOpenCL Cにあります）、 falseはゼロになり、 trueは1になります。 これらはスカラーのルールです。 したがって、 valは「1」になります。 これは必ずしも便利ではありませんが、そのような動作はプログラマの脳に組み込まれていますx+=(a>b)ような構造x+=(a>b)もはや驚くべきことでx+=(a>b)ません。

ただし、OpenCL Cでは、ベクトル整数型に変換する場合、 bool型の値は、すべてのビットがゼロであるか、すべてのビットが1である整数を返します。これは(int)-1対応します。 このトピックに関する仕様の内容は次のとおりです（セクション6.2.2明示的なキャスト ）。

boolをベクトル整数データ型にキャストすると、bool値がtrueの場合、ベクトルコンポーネントは-1（つまり、すべてのビットが設定されます）に設定され、そうでない場合は0に設定されます。

したがって、 val2はベクトル[-1, -1]ます。 これは、式が最初にベクトルコンポーネントの型に変換され、次に乗算される型変換のコンテキストでは少し予想外です-他の型と同様に、この動作はboolに対して宣言されます。 適切に使用すれば、条件式をビットごとの演算に置き換えることができます。

「コンパイルするかどうか？ 変数の値は何ですか？」と書いて、 opencl-sandboxプロジェクトをgithubにアップロードしました 。 私のマシンでこの記事のすべての例をチェックしました。 これを含む：

 __kernel void bool_to_int_vec() { int val = true; int2 val2 = true; printf("int val = true; // val=%d\n", val); printf("int2 val2 = true; // val2=%v2d\n", val2); if(val2.x == -1 && val2.y == -1) { printf("Compiler follows specification for bool->intn conversion, OK\n"); } else { printf("Compiler does not follow specification for bool->intn conversion, FAILED\n"); } } 

ご存知のように、コンパイラ開発者も人間であり、仕様を暗記していません。
2つのOpenCLプラットフォーム-AMDおよびIntel。 各プラットフォームには、GPUとCPUの2つのデバイスがあります。 そして、GPU用のAMDコンパイラ（最も成熟した）のみが仕様に従っており、他の3つは-1ではなくval2にユニットのベクトルを記述しています。


システムには5つのデバイスがあります。 AMDコンパイラも同じように動作します。 Intelの最新のコンパイラは「修正」され、現在は標準に従って動作します。 NVidiaコンパイラは、ベクトル型に変換するだけでなく、2番目のprintf()行を表示することもできませんでした。

これには2つの結論があります。

1. 仕様の知識がなければ、移植可能なコードを書くことはできません。
2. 各プラットフォームは独自の方法で仕様を理解するため、OpenCLカーネルをテストでカバーする必要があります。

ベクトルの論理演算子と比較演算子

boolをintにキャストする場合、対応する演算子はスカラーとベクトルに対して異なる動作をします。 演算子の実行結果の値> 、 < 、 >= 、 <= 、 == != 、 && 、 || ! 、これはintです。 スカラーの場合-0または1。ベクトルの場合-値0または-1を持つintの対応する長さのベクトル（すべてのビットが1に設定されます）。

 int a = 1 > 0; // a=1 int4 b = (int4)(1) > (int4)(0); // b=[-1,-1,-1,-1] 

4つのコンパイラをチェックすると、今回はすべてが正しい結果を出しました。

ベクトルの三項演算子

「 exp1 ? expr2 : expr3 」という形式の三項演算子 exp1 ? expr2 : expr3 」は、スカラーとベクトルに対しても同様に異なる動作をしexp1 ? expr2 : expr3 。 C99のようなスカラーの場合、 expr1ゼロでない場合、式の結果はexpr2 、 expr1ゼロの場合、 exp3 expr1ます。

ベクトルの場合、最初に、 expr1型は整数のみになります。 第二に、 expr1条件をチェックするときexpr1チェックはゼロと等しくないため、最初のビットでなく、 最上位ビットのためです 。 同時に、演算子はコンポーネントごとに機能します。 式expr2とexpr3がベクトルで、もう一方がスカラーである場合、スカラーは対応するコンポーネントを使用して暗黙的にベクトル型に変換されます。

 int a = 1 ? 1 : 0; // a=1 int4 b = (int4)(1, 0, 1, 0) ? (int4)(1) : 0; // b=[0,0,0,0] int4 c = (int4)(-1, 0, -1, 0) ? 1 : (int4)(0); // c=[1,0,1,0] int4 d = (uint4)(0x80000000u, 0, 0, 0) ? (int4)(1) : (int4)(0); // d=[1,0,0,0] //  C99  ,     OpenCL     float e = 0.0f ? 1 : 2; //  , expr1     float4 f = (float4)(0) ? (float4)(1) : (float4)(2); // expr2  expr3     float4 g = (int4)(-1, 1, -1, 1) ? (float4)(1) : (float4)(0); // g=[1.0f, 0.0f, 1.0f, 0.0f] 

ご覧のとおり、ここで再び類似性のtrapに陥ることがあります。 同じコードを正確にベクトル化と比較します。

 int a = 1 ? 1 : 0; // a=1 int4 b = (int4)(1) ? (int4)(1) : (int4)(0); // b=[0, 0, 0, 0] 

ベクトルb 、仕様とCプログラマーの戸惑いに完全に一致して、ゼロで満たされています。

実数型と整数型をOpenCL Cに変換する

スカラー型の場合、整数型から実数型、実数型から整数型への変換は、C99と同じ規則に従って実行されます。つまり、実数型から整数型への変換では小数部は拒否されます。元の整数と同じ値。 変換が行われる型の範囲に数値が収まらない場合、結果は実装によって異なります。

あるタイプのデータを別のタイプのデータとして解釈する必要がある場合、C99でこれを行う唯一の方法は、 memcpy関数を使用することだけです。 OpenCLにはmemcpyがありませんが、C99とは異なり、ユニオンを使用してデータを異なるタイプのデータとして解釈することは完全に合法です。

OpenCL言語は、ユニオンを拡張して、プログラムが異なるタイプのメンバーを使用してユニオンオブジェクトのメンバーにアクセスできるようにします。

さらに、飽和演算用のベクトル型とハードウェア機能がサポートされています-これにより、OpenCLでの型変換の機能が決まります。

次のタイプのタイプ変換がサポートされています。

1. 暗黙的な変換
2. 明示的なキャスト
3. 明示的な変換
4. データを別のタイプのデータとして解釈する（データを別のタイプとして再解釈する）。

OpenCLでは、アイテム1と2はC99に似ており、アイテム3と4は、ベクタータイプでの作業の利便性と明確さのための革新です。

暗黙的な変換と明示的なCスタイルのキャスト

C99のように、さまざまなタイプのオペランドが式で見つかった場合、それらは1つの共通タイプに変換されます。 違いは、これがベクトルに対してどのように機能するかです。 スカラー型の場合、C99と同じ方法で暗黙的な型変換と明示的な型変換がサポートされます。

 float a = 5.1f; int b = a; // b = 5 float c = 1.6f; int d = (int)c; // d = 1 

明示的または暗黙的にスカラー型からベクトル型に変換する場合、スカラーは最初にC99と同様の規則に従ってベクトル要素の型に変換され、次に乗算されます
ベクタータイプのサイズ：

 float a = 4.7f; float4 b = 5; // int 5 → float 5.0f → (float4)(5.0f, 5.0f, 5.0f, 5.0f) int4 c = 4.7f; // float 4.7f → int 4 → (int4)(4, 4, 4, 4) int4 d = 1; // int 1 → (int4)(1, 1, 1, 1) int4 e = (int4) a;// float 4.7 → int 4 → (int4)(4, 4, 4, 4)      float4 f = a; // float 4.7f → (float4)(4.7f, 4.7f, 4.7f, 4.7f) 

あるベクトル型から別のベクトル型への暗黙的な変換および明示的なCスタイルのキャストは禁止されています。 同じ数のコンポーネントがある場合でも。

 float4 a = (float4)(5.0f, 5.0f, 5.0f, 5.0f); //   ,     int4 b = a; // ,        float4 c = 0; int4 d = (int4)c; // ,        int4 e = (int4)(c); // ,      —      float4  int4 int4 f = (int4)(c.xy, c.zw); // ,      —      float2  int2 int4 g = (int4)(cx, cy, cz, cw); //      ,       float  int 

ベクトル型の明示的なキャストはありませんが、スカラーをベクトル型にキャストできます。 これにより、ベクトルリテラルがさらに混乱します。 同じコンポーネントを持つベクトルを指定する3つの方法を比較します。

 float2 a = (float2)(1); //   float2 b = (float2)1; //       float2 c = 1; //       

コンポーネントが異なるベクターの場合、同じコードは機能しません。ベクターリテラルを使用するだけです。 最悪なのは、以下のすべてのコードが正常にコンパイルされることです。結果だけが適切です。

 float2 a, b, c, d; //    a = (float2)(1, 2); // a=[1, 2] // 1       , 2 —  b = (float2)1, 2; // b=[1, 1] // 1       , 2 —  c = 1, 2; // c=[1, 1] // 1 , 2        d = (1, 2); // d=[2, 2] 

実数型と整数型の明示的な変換

Cスタイルのキャストに加えて、OpenCLには、オーバーフロー状態を処理し、ベクトルを処理する型キャストメカニズムがあります。 これは一連の機能です。

 convert_() 

より一般的な機能

 convert_<_sat><_>() 

さらに、オーバーフローモードと丸めモードを使用します。 スカラーとベクトルの場合、関数は同じように機能します。 ソースのベクトルの要素数と結果の型は一致する必要があります。

 float a = 5.5f; int b = convert_int(a); // b = 5 float4 c = a; // c=[5.5, 5.5, 5.5, 5.5] float2 d = convert_float2(c); // ,           //         //        int4 e = convert_int4(c); // e=[5,5,5,5] 

整数型にキャストする場合、オーバーフローの動作はオプションで決定されます
修飾子_sat 。 これがないと、C99で通常どおり整数型のオーバーフローが発生し、飽和が機能するため、型で許可される範囲外の値は、変換後の型で表現できる最も近い値に削減されます。

 int a = 257; uchar b = convert_uchar(a); // b = 1,   b = convert_uchar_sat(a); // b = 255,   

実際の型にキャストする場合、 _sat許可されません。 これは必要ありません。実際の型がオーバーフローすると、それらはすでに±INFになるためです。

丸めを制御するには、修飾子_rte （最も近い偶数に_rtz ）、 _rtz （ゼロに_rtp ）、 _rtp （正の無限大に_rtn ）、および_rtn （負の無限大に丸める）が使用され、最も近い整数への丸め、ゼロへの丸め、丸めるプラス無限大およびマイナス無限大への丸め。 丸め修飾子がない場合は、 _rtz使用して実数から整数に変換し、整数から実数に変換するときに_rtz使用します。 _rteは、通常の数学ではなく、最も近い整数への丸めのいわゆる「 バンキング 」バージョンを使用します。 小数部が正確に0.5である場合、最も近い整数は1つではなく、2つの最も近い整数から偶数が選択されます。

 int a = convert_int_rtp(4.2f); // a = 5 a = convert_int(4.2f); // a = 4 int4 b = convert_int4_rte((float4)M_PI_F); // b = [3, 3, 3, 3] 

さまざまな丸めモードでfloatをint変換します（ ここでチェックします ）：

	0.5	-0.5	1.1	-1.1	1.5	-1.5	1.7	-1.7
最も近い整数への丸め （最も近い偶数に丸め、rte）	0	0	1	-1	2	-2	2	-2
ゼロへの丸め （ゼロに向かって丸め、rtz）	0	0	1	-1	1	-1	1	-1
正の無限大への丸め （正の無限大に向かって丸め、rtp）	1	0	2	-1	2	-1	2	-1
負の無限大への丸め （負の無限大に向かって丸め、rtn）	0	-1	1	-2	1	-2	1	-2

Wikipediaの丸めに関する英語の記事にはすばらしい説明があります。 その上のrteモードは、「偶数」、 rtz 「切り上げ→ゼロ」、 rtz 「 rtz 」、 rtz 「 rtz 」に対応します。

データを別のタイプのデータとして解釈する

unionに加えて、OpenCLで1つのタイプのデータを別のタイプのデータとして解釈するために、スカラーとベクトル用の関数as_()ファミリーがあります。

 float a = 25.0f; int b = as_int(a); // b=0x41C80000,     25.0f 

元の型と新しい型のバイト単位のサイズが一致しない場合、 as_はコンパイルエラーを引き起こします。

 int a = 0; char b = as_char(a); //, sizeof(int)!=sizeof(char) float2 c = 0; float8 d = as_float8(c); //, sizeof(float2)!=sizeof(float8) 

元の型と新しい型の要素の数が一致しない場合（ただし、型のサイズは同じです）、結果がOpenCL（実装定義）の実装に依存します。ただし、オペランドが4コンポーネントベクトルであり、結果が3コンポーネントベクトルである場合は例外です。 したがって、32ビットワードのバイトをベクトル要素として取得すると便利です。

 uint word = 0x01020304; uchar4 bytes = as_uchar4(word); 

ただし、この場合の結果は、特定のOpenCL実装の裁量で、 [1, 2, 3, 4] [4, 3, 2, 1] 、および[1, 2, 3, 4] 、またはその他のものになります。 ただし、OpenCLのいずれかのバージョンを最適化して作業する場合、 as_このような使用は非常に受け入れられます。

オペランドが4成分のベクトルであり、結果が3成分のベクトルである場合、
as_ 、元の型のビットを変更せずに返すために必要です-標準
要素のサイズが同じ場合、3つのコンポーネントのベクトルのサイズは4つのコンポーネントのベクトルのサイズに等しくなります。

 float4 a = 1.0f; int3 b = as_int3(a); // ,   sizeof(int3)==sizeof(float4) // b=[0x3f800000, 0x3f800000, 0x3f800000] char3 c = as_char3(a); // , sizeof(char3)!=sizeof(float4) 

おわりに

OpenCL Cは、通常のC99との類似点で潜行的です。 この記事を読んだ後、

• ベクトルリテラルと明示的なキャストを混同しないでください。
• ベクトル型の論理演算子に該当しないでください。
• convert_ *およびas_ *関数を兵器庫に追加します。