| On this page | |
| Since | 16.0 |
OpenCL SOPには、ジオメトリに対してOpenCLカーネルを作成して実行するための一般的なインターフェースが備わっています。 このSOPは、そのカーネルのOpenCLパラメータに定数、アトリビュート、ボリュームデータをバインドすることができます。
Warning
このノードでは、あなたがOpenCLを理解していることが求められます。 このノードを使用して不正なコードを書いてしまうことが非常に起こりやすいです。
構文 ¶
利用可能な構文に関する基本情報は、VEXユーザ向けのOpenCLを参照してください。
パラメータ ¶
Kernel ¶
Kernel Name
読み込んだプログラムを使って実行するOpenCLカーネルの名前。
Use Code Snippet
外部ディスクファイルではなくKernel Codeパラメータで指定されたコードを使用します。これは、OpenCL SOPを簡単に編集/作成することができます。
Kernel File
コンパイルするOpenCLプログラムファイルのパス。ここには、ディスクファイルまたはアセットのパスを含めることができます。
Kernel Code
Code Snippetモードの時に実行するOpenCLカーネル。
Enable @-Binding
Code Snippetモードでは、カーネルの生成とそのカーネルのジオメトリ操作を簡単にできるようにした@接頭辞マクロの利用が有効になります。
Kernel Options
カーネルで必要なコンパイルフラグを指定します。 よくある使い方は、-Dフラグを使ってプリプロセッサ用の#defineディレクティブを定義することです。
Note
Apple OSX OpenCLコンパイラは、カーネルオプション間に1個のスペースだけを必要とします!
Houdiniは、カーネルをコンパイルする時にOpenCLデバイスに応じて追加フラグを定義します。
__H_GPU__または__H_CPU__のフラグは、GPUとCPUのデバイスを区別し、__H_NVIDIA__、__H_AMD__、__H_INTEL__、__H_APPLE_のフラグはハードウェアベンダーを意味します。
Houdiniを起動する前にHOUDINI_OCL_REPORT_BUILD_LOGS環境変数を1に設定することで、プリプロセッサフラグと一緒にコンパイルされたすべてのカーネルのダンプを取得することができます。
Option Attribute
カーネルのコンパイルフラグとして追加される文字列Detailアトリビュートを指定します。 これは、入力ジオメトリ内からこの名前のDetailアトリビュートを取得します。 そのDetailアトリビュートが文字列だった場合、その文字列がカーネルオプションとして注入されます。 例えばdefineを指定する場合は、この文字列に-Dオプションを入れてください。
Note
文字列の値を頻繁に変えないでください。 そうするとカーネルが再コンパイルを続けることになって非常に負荷が大きくなってしまいます。
Use Write Back Kernel
カーネルを実行した後に、それにバインドされた同じパラメータのセットで2番目のカーネルを即座に実行することができます。 複数のスレッドが同じデータに書き込みたい時に、それを2回のパス処理に分けることで、競合状態を回避することができます。
Write Back Kernel Name
書き戻し処理に使用するカーネルの名前。これは、メイン処理と同じ関数シグネチャを持ちます。
Recompile Kernel
ディスクからカーネルを読み込んだ時、そのカーネルは、計算の度に再生成されないようにキャッシュ化されます。 これを有効にすると、そのカーネルが強制的に再読み込みされて再コンパイルされます。 これは、#includeファイルが変更されたコードを参照していたり、カーネルファイルが外部テキストエディタで変更されている場合に役立ちます。
プロトタイプが完成したら、これを常に無効にしてください。
Options ¶
Run Over
指定したOpenCLカーネルを一度だけ呼び出します。とはいえ、グローバルidの数はこの設定で制御されます。 First Writeableアトリビュートは、それを書き込み可能とマークされた最初にバインドされたアトリビュートのサイズに設定します。 同様に、First Writeable Volumeは、それを書き込み可能とマークされた最初にバインドされたボリュームの合計のボクセルに設定します。
worksetモードでは、整数配列のDetailアトリビュートを使って、worksetの数とそのサイズを決定します。
Note
このグローバルidは、GPU上で効率的に処理できるように上手くまとめ上げられるので、常にget_global_id(0)とバインドされたアトリビュートの実際の長さを比較してください。
Iterations
カーネルは、可変回数で再実行することができます。これによって、より多くのノードを使ってFor Loopを作成する必要がなくなり、一連の評価の間にすべてのデータをビデオカード上に残せるようになります。
Worksets Begin Attr.
各worksetのstart値を指定した整数配列のDetailアトリビュート。
Worksets Length Attr.
各worksetのlength値を指定した整数配列のDetailアトリビュート。 カーネルは、これらのサイズ毎に呼び出されます。 ゼロのサイズはスキップされます。
Use Single Workgroup If Possible
GPU上でworksetを実行する時、別々のworkset毎にカーネルを実行するよりも、1回のカーネルコール内でGPU上でたくさんの小さいworksetを実行して、各worksetの後にそのカーネル内で同期をさせた方が高速です。
このオプションを有効にすると、一番大きいworksetがOpenCL GPUデバイス上の1ワークグループ内に収まれば、SINGLE_WORKGROUPプリプロセッサフラグが定義され、 Worksets Begin と Worksets Length の配列全体がカーネルに渡されます。
各worksetの後での同期はカーネルに依存し、通常ではbarrier(CLK_MEM_GLOBAL_FENCE)を使用します。
Generate Kernel ボタンで生成されるコードには、この同期を制御する方法が表示されます。
Finish Kernels
Finish Kernelsが無効な時、次のソルバを続行する前にOpenCLカーネルが完了するのを待つ必要はありません。 これによって、OpenCLカーネルの実行結果が実際に必要になるまではバックグラウンドでOpenCLカーネルを走らせることができます。 ただ、デバッグを簡単にするには、必ず正しい箇所でエラーが検出できるようにFinish Kernelsを有効にした方が役立ちます。
Include Time
現行時間をパラメータとして含めます。
Include Timestep
現行タイムステップをパラメータとして含めます。
Time Scale
タイムステップの一定乗数。
Time Method
いくつかの処理では、タイムステップの累乗を知りたいことがあります。
カーネルで再計算するのではなく、これをe^Timestepに設定することで、その指数を事前計算させることができます。
Include Simplex Noise Data
OpenCLカーネルからSimplexノイズやCurlノイズを生成するために、<xnoise.h>のsimplex noise関数に渡すことができるOpaqueポインタを含めます。
Precision
このノードの精度を制御します。
fprealタイプとexintタイプは、ここで指定した精度に合うように、生成されるコード内で定義されます。
ベクトルバリアント(例えば、fpreal3、fpreal4など)も定義されます。
さらにFPREAL_PRECシンボルは、halfには16、floatには32、doubleには64として定義されます。
Autoは、入力ジオメトリの優先精度を使用します。これは、
Attribute Cast SOPで設定することができます。
Note
たいていのドライバでは、計算に16ビットを使用することはできません。
Bindings ¶
OpenCL Parameters ¶
各パラメータでは、カーネルを呼び出す前に評価される固定定数値を指定したり、ボリュームやジオメトリアトリビュートから読み/書きすることができます。
Parameter Name
パラメータの名前。これは、Generate Kernelボタンで使用されますが、そうでない場合はコメントとして存在するだけです。
OpenCLカーネルの実際のバインドは、名前ではなくパラメータ順で行なわれます。
Options ¶
Parameter Type
作成とバインドをするパラメータのタイプ。
Integer
定数の整数値。これは、チャンネル参照や事前に計算されるエクスプレッションをバインドすることができます。
Float
定数の浮動小数点値。オプションで、これをtimestepでスケールすることができます。
Float Vec4
float4 OpenCLパラメータにバインドされた4つの浮動小数点の定数タプル。
Ramp
スカラーランプ。OpenCLカーネル内のスプラインベースのランプの評価は複雑なので、代わりに、このランプが均一な浮動小数点の配列にサンプリングされます。
Ramp Sizeパラメータは、使用するサンプルの数を制御します。
Attribute
ジオメトリアトリビュートをバインドします。
Volume
ボリュームをバインドします。
VDB
VDBをバインドします。
Integer
整数モードを使用する整数値。
Float
浮動小数点モードを使用する浮動小数点値。
Float 4
float 4モードを使用するFloat 4値。
Time Scale
指定した浮動小数点値をtimestepでスケールする方法。 パラメータ評価の時点でtimeincを知ることができないので、カーネルを評価する前にtimeincを定数として計算して、それを浮動小数点値に適用することができます。
Ramp
浮動小数点値のリストとして用意するランプデータ。
Ramp Size
ランプを評価する浮動小数点値の数。
Volume
バインドするボリュームまたはVDBのプリミティブの名前または番号。
Force Alignment
カーネルを単純化するために、たいていの場合、すべてのボリュームが解像度とトランスフォームが揃っているものと仮定することができます。 Force Alignmentを設定すると、これが強制されて、ボリュームがずれているとエラーを生成します。
Voxel Resolution
ボリュームの解像度をパラメータとして追加します。
Voxel Size
ボリュームのサイズをSOP空間でパラメータとして追加します。
Volume Transform to World
ボリュームのボクセル座標からSOP座標に変換する行列トランスフォームを追加します。
Volume Transform to Voxel
SOP座標からボリュームのボクセル座標に変換する行列トランスフォームを追加します。
Attribute
バインドするアトリビュート。オプションのフラグを設定しない限りは、このアトリビュートが見つからなければエラーになります。
Attributes用に存在します。
Class
アトリビュートのタイプ。First Writableアトリビュートは反復順を決定することができるので、これはOpenCL Solverで処理されるグローバルidの番号を決定することができます。
すべてのバインドアトリビュートが同じタイプである必要はなく、同じジオメトリデータから派生している必要もありません。
Attributes用に存在します。
Type
バインドするアトリビュートのタイプ。浮動小数点と整数のアトリビュートは、順番通りにすべてのエレメント値を含んだ単一配列としてバインドされます。
タプルは、交互に配置されます。つまり、Pはxyzxyzxyzとしてバインドされます。
配列アトリビュートは、2つの配列としてバインドされます。1つ目の配列には、各エレメントの配列データが含まれます。そのため、オフセットのペアはエレメント配列の長さが異なります。 2つ目の配列は、単一配列に連結されたすべてのエレメントの配列のデータです。
Attributes用に存在します。
Size
バインドするアトリビュートのタプルサイズ。 このサイズが1以上であれば、そのタプルサイズのアトリビュートを用意しなければなりません。 0であれば、自動的にバインドされ、そのタプルサイズに格納するための追加パラメータが生成されます。
Attributes用に存在します。
Precision
このパラメータのデータと紐付ける精度を制御します。
Nodeオプションは、このノードの精度を使用するので、その設定に応じて変わり、それに相当するカーネルコードは、
fprealまたはexintの定義を使用します。
これは、ビデオカード上に格納されるデータの精度なので、低い精度を使用することでGPUメモリを節約することができます。
しかし、halfに相当する16ビットは、たいていの場合では計算に使用することができません。
vload_halfを使用することで、それを計算用にfloatにプロモートさせることができます。
同じアトリビュートが異なる精度にバインドされると、そのバインドが失敗します。
現在のところ、ボリュームは32ビットデータ精度にのみバインドすることができます。
Readable
OpenCLカーネルがこのアトリビュートまたはボリュームから読み込むかどうかを決定します。設定しなかった場合、アトリビュート値がGPU上にコピーされなくなります。 これは、書き込み専用アトリビュートが無駄にコピーされるのを回避するのに役立ちますが、初期化されていないデータが存在してしまうので注意が必要です。
Writeable
OpenCLカーネルがこのアトリビュートまたはボリュームに書き戻すかどうかを決定します。CPUバージョンのアトリビュートまたはボリュームを古いとマークさせるので、次回にそれが必要になった時は、GPUからコピーされるようになります。
Optional
アトリビュートまたはボリュームを不要とマークします。ジオメトリ内にアトリビュートまたはボリュームが存在しなかった場合、エラーを起こさずに、カーネルオプションに#defineを設定して、そのアトリビュートを無効にします。 これは、パラメータシグネチャも変更するので、 Generate Code ボタンを使ってその構文を検証するべきです。
Note
パラメータ名は#defineで使用されるので、そのパラメータ名を変更すると、そのコードを変更する必要があります。
Default Value
オプションのアトリビュートやボリュームが見つからない場合でも、パラメータ値をカーネルにバインドするべきであることをマークします。
#defineがカーネルオプションで設定され、そのアトリビュートが無効になり、単一値に切り替わります。
Note
これはパラメータのシグネチャまでも変更してしまうので、 Generate Code ボタンを使用して構文を確認してください。
バインドされたパラメータの値は、このパラメータの整数値または浮動小数点値から取得されます。
Generated Code ¶
Generate Kernel
Enable @-Binding が有効になっている場合、実際のコンパイラーに送信されるコードには、完全に展開されたコードが生成されます。
これによって、コンパイラが#lineディレクティブを評価しなくてもエラーの行番号を解決することができ、さらに、@マクロの挙動を理解するのにも役立ちます。
Note
この@マクロの実展開に依存するべきではありません。
Enable @-Binding が無効になっている場合、必須のカーネル関数のプロトタイプを作成して、現在選択されているパラメータすべてが考慮されます。 これは、関数を書き始める時のスタートポイントとして使用したり、新しいパラメータを追加/削除した時にインターフェースを更新することができます。
Generated Code
すべての@バインドが展開されたコードスニペット。
Note
このパラメータは使用されず、単に情報提供が目的です。
| See also |
Gas OpenCL
