opencl_node は、統合型および外付け型グラフィックス・プロセシング・ユニット (GPU) や CPU のような、OpenCL* をサポートするデバイスで OpenCL* カーネルを実行できるようにします。opencl_node を利用して、インテル® TBB フローグラフを使用するプログラムに OpenCL* カーネルを簡単に統合できます。また、opencl_node はそのようなプログラムでメモリーバッファー管理も行います。
template < typename JP, typename Factory, typename... Ports > class opencl_node < tuple < Ports... >, JP, Factory > // Factory - default_opencl_factory を使用 template < typename JP, typename... Ports > class opencl_node < tuple < Ports... >, JP > // Factory - default_opencl_factory を使用 // JP - queueing policy を使用 template < typename... Ports > class opencl_node < tuple < Ports... > >
#define TBB_PREVIEW_FLOW_GRAPH_NODES 1 #define TBB_PREVIEW_FLOW_GRAPH_FEATURES 1 #include "tbb/flow_graph_opencl_node.h"
"flow_graph_opencl_node.h" ヘッダーは、"tbb/tbb.h" と "tbb/flow_graph.h" にインクルードされていません。 このため、"flow_graph_opencl_node.h" を直接インクルードします。
opencl_node は、OpenCL* をサポートするデバイスを使用するための簡単なインターフェイスを提供します。opencl_node は、opencl_factory を使用する streaming_node として実装されます。opencl_factory は、カーネルの実行、デバイスとのメモリーの同期、結果の取得に関連するすべての複雑さを隠蔽します。opencl_factory は、Factory コンセプトを実装して、streaming_node 実行ワークフローに従います。
opencl_factory は、デバイスの初期化と実行管理に 2 つの異なるエンティティー (Device Filter と Device Selector) を使用します。
エンティティー |
使用モデル |
---|---|
Device Filter |
カーネルの実行に利用可能なデバイスのセットで opencl_factory を初期化するのに使用します。フィルターされたすべてのデバイスは同じ OpenCL* プラットフォームに属していなければなりません。デフォルトの Device Filter には、最初に検出された OpenCL* プラットフォームで利用可能なすべてのデバイスが含まれます。 |
Device Selector |
特定の opencl_node インスタンスのカーネルを実行するデバイスを選択するのに使用します。デフォルトの Device Selector は、Device Filter によって生成された利用可能なデバイスリストの最初のデバイスを選択します。選択は、カーネルの実行ごとに行われます。 |
さらに、opencl_node は、OpenCL* エンティティー (デバイス、メモリーバッファー、カーネル、ND 範囲) の処理を容易にするヘルパークラスを提供します。
エンティティー |
説明 |
---|---|
opencl_buffer/opencl_subbuffer |
厳密な型の線形配列を抽象化したテンプレート・クラス。OpenCL* カーネルは、ホスト上に割り当てられた特別なメモリー・オブジェクトである opencl_buffer と opencl_subbuffer を使用して、ホストとターゲット間のメモリー・トランザクションのロジックをカプセル化します。 |
opencl_device |
OpenCL* デバイスを抽象化したクラス。主に Device Selector と Device Filter の内部で使用され、デバイスに関する情報を照会するパブリック API があります。 |
opencl_device_list |
opencl_device インスタンスのコンテナー。 |
opencl_program |
カーネルの種類の指定、ファイルからのカーネルの読み取り、カーネルのビルドに使用できるクラス。 |
opencl_program_type |
列挙子 opencl_program_type は、opencl_node でサポートされるカーネルの種類のセットを提供します。
|
opencl_range |
範囲をサポートするため opencl_node を有効にします (詳細は、OpenCL* の clEnqueueNDRangeKernel() メソッドを参照)。 |
引数バインド
デフォルトでは、opencl_node は、1 つ目の入力ポートを 1 つ目のカーネル引数に、2 つ目の入力ポートを 2 つ目のカーネル引数に、というようにバインドします。opencl_node の set_args メソッドは、実行ごとに変更がなければ、各カーネル引数に特定の値をバインドできます。代わりに値を入力ポートから取得するように指定するには、port_ref ヘルパー関数を使用します。このドキュメントの「streaming_node」セクションにある説明と使用例を参照してください。
次のコードは、一部の計算を GPU へオフロードするように立方体と四角形のサンプルを変更したものです。
カーネル__kernel void cuber( __global int* buf ){ int idx = get_global_id(0); buf[idx] = buf[idx] * buf[idx] * buf[idx]; } __kernel void squarer( __global int* buf ){ int idx = get_global_id(0); buf[idx] = buf[idx] * buf[idx]; }サンプル
#define TBB_PREVIEW_FLOW_GRAPH_NODES 1 #define TBB_PREVIEW_FLOW_GRAPH_FEATURES 1 #include "tbb/flow_graph_opencl_node.h" typedef tbb::flow::opencl_buffer<cl_int> opencl_vector; class gpu_device_selector { public: template <typename DeviceFilter> tbb::flow::opencl_device operator()(tbb::flow::opencl_factory<DeviceFilter>& f) { // Set your GPU device if available to execute kernel on const tbb::flow::opencl_device_list &devices = f.devices(); tbb::flow::opencl_device_list::const_iterator it = std::find_if( devices.cbegin(), devices.cend(), [](const tbb::flow::opencl_device &d) { cl_device_type type; d.info(CL_DEVICE_TYPE, type); return CL_DEVICE_TYPE_GPU == type; }); if (it == devices.cend()) { std::cout << "Info: could not find any GPU devices. Choosing the first available device (default behaviour)." << std::endl; return *(f.devices().begin()); } else { // Return GPU device from factory return *it; } } }; int main() { const int vector_size = 10; int sum = 0; tbb::flow::graph g; tbb::flow::broadcast_node< opencl_vector > broadcast(g); // GPU computation part gpu_device_selector gpu_selector; tbb::flow::opencl_program<> program("vector_operations.cl"); tbb::flow::opencl_node< std::tuple< opencl_vector > > vector_cuber(g, program.get_kernel("cuber"), gpu_selector); tbb::flow::opencl_node< std::tuple< opencl_vector > > vector_squarer(g, program.get_kernel("squarer"), gpu_selector); // Define kernel argument and problem size vector_cuber.set_args(tbb::flow::port_ref<0>); vector_cuber.set_range({{ vector_size }}); vector_squarer.set_args(tbb::flow::port_ref<0>); vector_squarer.set_range({{ vector_size }}); // Computation results join tbb::flow::join_node< std::tuple< opencl_vector, opencl_vector > > join_data(g); // CPU computation part tbb::flow::function_node< std::tuple< opencl_vector, opencl_vector > > summer( g, 1, [&](const std::tuple< opencl_vector, opencl_vector > &res ) { opencl_vector vect_cubed = std::get<0>(res); opencl_vector vect_squared = std::get<1>(res); for (int i = 0; i < vector_size; i++) { sum += vect_cubed[i] + vect_squared[i]; } }); // Graph topology make_edge( broadcast, vector_cuber ); make_edge( broadcast, vector_squarer ); make_edge( vector_cuber, tbb::flow::input_port<0>(join_data) ); make_edge( vector_squarer, tbb::flow::input_port<1>(join_data) ); make_edge( join_data, summer ); // Data initialization opencl_vector vec(vector_size); std::fill( vec.begin(), vec.end(), 2 ); // Run the graph broadcast.try_put(vec); g.wait_for_all(); std::cout << "Sum is " << sum << "\n"; return 0; }
namespace tbb { namespace flow { template < typename... Args > class opencl_node; template < typename JP, typename Factory, typename... Ports > class opencl_node < tuple < Ports... >, JP, Factory > : public streaming_node < tuple < Ports... >, JP, Factory > { public: typedef implementation-dependent kernel_type; opencl_node( graph &g, const kernel_type& kernel ); opencl_node( graph &g, const kernel_type& kernel, Factory &f ); template < typename DeviceSelector > opencl_node( graph &g, const kernel_type& kernel, DeviceSelector d, Factory &f); }; template < typename JP, typename... Ports > class opencl_node < tuple < Ports... >, JP > : public opencl_node < tuple < Ports... >, JP, opencl_info::default_opencl_factory > { public: typedef implementation-dependent kernel_type; opencl_node( graph &g, const kernel_type& kernel ); template < typename DeviceSelector > opencl_node( graph &g, const kernel_type& kernel, DeviceSelector d ); }; template < typename... Ports > class opencl_node < tuple < Ports... > > : public opencl_node < tuple < Ports... >, queueing, opencl_info::default_opencl_factory > { public: typedef implementation-dependent kernel_type; opencl_node( graph &g, const kernel_type& kernel ); template < typename DeviceSelector > opencl_node( graph &g, const kernel_type& kernel, DeviceSelector d ); }; } }
次の表は、このテンプレート・クラスのメンバーの詳細な情報を提供します。ほかのすべてのパブリックメンバーは、基本 streaming_node クラスに説明があります。
メンバー |
説明 |
---|---|
typename... Ports |
ノードの入力/出力データ型。 |
typename JP |
結合ポリシー。詳細は、join_node クラスの説明を参照してください。 |
typename Factory |
デバイス固有の Factory 型。このノードでは、デフォルトで opencl_factory が使用されます。 |
opencl_node( graph &g, const kernel_type& kernel ) opencl_node( graph &g, const kernel_type& kernel, Factory &f ) template < typename DeviceSelector > opencl_node( graph &g, const kernel_type& kernel, DeviceSelector d ) template < typename DeviceSelector > opencl_node( graph &g, const kernel_type& kernel, DeviceSelector d, Factory &f) |
メイン・コンストラクター。詳細は、streaming_node を参照してください。 |