opencl_node テンプレート・クラス

概要

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) を使用します。

opencl_factory Device Selector

エンティティー

使用モデル

Device Filter

カーネルの実行に利用可能なデバイスのセットで opencl_factory を初期化するのに使用します。フィルターされたすべてのデバイスは同じ OpenCL* プラットフォームに属していなければなりません。デフォルトの Device Filter には、最初に検出された OpenCL* プラットフォームで利用可能なすべてのデバイスが含まれます。

Device Selector

特定の opencl_node インスタンスのカーネルを実行するデバイスを選択するのに使用します。デフォルトの Device Selector は、Device Filter によって生成された利用可能なデバイスリストの最初のデバイスを選択します。選択は、カーネルの実行ごとに行われます。

さらに、opencl_node は、OpenCL* エンティティー (デバイス、メモリーバッファー、カーネル、ND 範囲) の処理を容易にするヘルパークラスを提供します。

opencl_node ヘルパークラス

エンティティー

説明

opencl_buffer/opencl_subbuffer

厳密な型の線形配列を抽象化したテンプレート・クラス。OpenCL* カーネルは、ホスト上に割り当てられた特別なメモリー・オブジェクトである opencl_bufferopencl_subbuffer を使用して、ホストとターゲット間のメモリー・トランザクションのロジックをカプセル化します。

opencl_device

OpenCL* デバイスを抽象化したクラス。主に Device Selector と Device Filter の内部で使用され、デバイスに関する情報を照会するパブリック API があります。

opencl_device_list

opencl_device インスタンスのコンテナー。

opencl_program

カーネルの種類の指定、ファイルからのカーネルの読み取り、カーネルのビルドに使用できるクラス。

opencl_program_type

列挙子 opencl_program_type は、opencl_node でサポートされるカーネルの種類のセットを提供します。

  • SOURCE は、OpenCL* C ソースコードです。SOURCE で初期化された opencl_program インスタンスは、実行時にソースファイルを読み取り、カーネルをビルドします。SOURCE は、opencl_program のデフォルト値です。
  • PRECOMPILED は、ターゲットで実行可能な機械可読形式です。この形式のカーネルは、実行時にコンパイルを必要としません。
  • SPIR は、SPIR* (Standard Portable Intermediate Representation) です。

opencl_range

範囲をサポートするため opencl_node を有効にします (詳細は、OpenCL* の clEnqueueNDRangeKernel() メソッドを参照)。

引数バインド

デフォルトでは、opencl_node は、1 つ目の入力ポートを 1 つ目のカーネル引数に、2 つ目の入力ポートを 2 つ目のカーネル引数に、というようにバインドします。opencl_nodeset_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 を参照してください。

関連情報