opencl_node の基本インターフェイスと opencl_buffer – パート 2

2016.03.28
インテル® oneTBB
C/C++OpenCL

この記事は、インテル® デベロッパー・ゾーンに公開されている「opencl_node basic interfaces and opencl_buffer」(https://software.intel.com/en-us/blogs/2015/12/14/opencl-node-basic-interfaces-and-opencl-buffer/) の日本語参考訳です。

この記事は、インテル® スレッディング・ビルディング・ブロック (インテル® TBB) 4.4 Update 2 以降で利用可能な新しいノード opencl_node について説明するシリーズのパート 2 です。このノードは、インテル® TBB のフローグラフで OpenCL* デバイスの利用と連携を可能にします。このシリーズのパート 1 はこちらからご覧いただけます。

opencl_node の基本インターフェイス

以下は、opencl_node のジェネリック宣言です。

template< typename... PortTypes, 
    typename Policy = queueing, 
    typename Factory = default_opencl_factory > 
class opencl_node< tuple<PortTypes...>, Policy, Factory >;

opencl_node には入力ポートと出力ポートがあり、tuple<PortTypes…> でそれぞれの種類のポートを定義します。必須のテンプレート引数はこれだけです。

opencl_node は、join_node と同様のバッファーポリシーをサポートします。opencl_node で有効なポリシーは、queueing (デフォルト) と key_matching です。

第 3 引数の Factory は OpenCL* コンテキストを抽象化し、カーネル実行デバイスを指定します。デフォルトでは、opencl_node は default_opencl_factory で初期化されます。現在、default_opencl_factory はシステムで利用可能な最初の OpenCL* プラットフォームの最初のデバイスを選択します (将来、この動作は変更される可能性があります)。

入力ポートにメッセージが送られると、メッセージの関連データがデバイスで利用できるようになります。OpenCL* カーネルは、すべての入力ポートがメッセージを受け取ると実行されます。メッセージは異なるレートで入力ポートに送られますが、各ポートの合計メッセージ数は同じでなければなりません。カーネルに送られた後、メッセージはサクセサーへフォワードされます。

opencl_node のインスタンスを作成するには、少なくとも opencl_graph、opencl_program、およびカーネル名を作成時に指定する必要があります。

template <typename DeviceSelector> 
opencl_node( opencl_graph &g, 
    const opencl_program<Factory> &p, 
    const std::string &kernel_name, 
    DeviceSelector d = /* デフォルト・デバイス・セレクター */, 
    Factory &f = /* デフォルトの OpenCL* factory */ );

opencl_node には 1 つのカーネル名のみ指定できます。つまり、1 つの OpenCL* カーネルのフローグラフ抽象化のみ利用できます。アプリケーションで複数のカーネルを使用する場合は、複数の opencl_node オブジェクトを作成し、make_edge インターフェイスを用いてオブジェクト間の依存性を設定します。次に例を示します。

#define TBB_PREVIEW_FLOW_GRAPH_NODES 1
#include "tbb/flow_graph_opencl_node.h"

#include <numeric>

int main() {
    using namespace tbb::flow;
    
    opencl_graph g;
    opencl_node<tuple<cl_int>> cl1( g, "simple_dependency.cl", "k1" );
    opencl_node<tuple<cl_int>> cl2( g, "simple_dependency.cl", "k2" );
    opencl_node<tuple<cl_int>> cl3( g, "simple_dependency.cl", "k3" );
        
    make_edge( output_port<0>(cl1), input_port<0>(cl2) );
    make_edge( output_port<0>(cl1), input_port<0>(cl3) );
 
    cl1.set_ndranges( { 1 } );
    cl2.set_ndranges( { 1 } );
    cl3.set_ndranges( { 1 } );
    input_port<0>(cl1).try_put( 0 );

    g.wait_for_all();
    
    return 0;
}

simple_dependency.cl:

kernel void k1( int b ) { printf("kernel #1\n"); }
kernel void k2( int b ) { printf("kernel #2\n"); }
kernel void k3( int b ) { printf("kernel #3\n"); }

上記のコードは、

kernel #1 
kernel #2 
kernel #3

または

kernel #1 
kernel #3 
kernel #2

を出力します。2 つ目と 3 つ目のカーネルは、1 つ目のカーネルに依存しますが、互いには依存しません。

make_edge( output_port<0>(cl1), input_port<0>(cl2) ); 
make_edge( output_port<0>(cl1), input_port<0>(cl3) );

そのため、カーネル k1 は常に最初に実行されますが、2 つ目と 3 つ目のカーネルの実行順序は非決定的です。

opencl_buffer

OpenCL* カーネルは、ホスト上に割り当てられた特殊なメモリー・オブジェクトを操作できます。フローグラフは、厳密な型の線形配列の抽象化である opencl_buffer テンプレート・クラスを提供します。

template <typename T, typename Factory = default_opencl_factory>
class opencl_buffer {
public:
    typedef /* 実装定義 */ iterator;
    
    // データアクセサー
    T* data();
    iterator begin();
    iterator end();
    T& operator[] ( ptrdiff_t k );
    size_t size();
    
    // コンストラクター
    opencl_buffer( Factory &f, size_t size );
    opencl_buffer( opencl_graph &g, size_t size );
};

バッファーの作成には、opencl_factory または opencl_graph が必要です。opencl_graph の場合、バッファーは (opencl_graph オブジェクトの) default_opencl_factory により作成されます。さらに、配列の要素数も指定する必要があります。次に例を示します。

opencl_graph g;
const int N = 1000;
opencl_buffer<cl_int> buf( g, N );

上記の例では、1000 要素の整数配列が作成されます。バッファーにデータを格納するには、次のようにホスト側でデータアクセサーを利用します。

std::iota( buf.begin(), buf.end(), 0 );

上記のコードは、0 から 999 の値をバッファーに格納します。OpenCL* デバイスでバッファーを並列に処理するには、同じサイズの ndrange を使用します。

std::vector<int> ndrange( 1, N );
clNode.set_ndranges( ndrange );

前述の例とは異なり、この例では ndrange に std::vector を使用しています。ベクトルには N と等しい 1 つの要素が含まれており、 set_ndrange に渡されると、1 次元の範囲 [0,N) として扱われます。

set_ndranges 関数には、begin() および end() メソッドを提供する任意のコンテナーを渡すことができます。std::initializer_list、std::vector, std::array を含む標準の C++ 型の多くは、これらのメソッドを提供しています。範囲の次元はコンテナーに含まれる要素数と同じで、各次元のサイズは対応する要素の値になります。

以下に、サンプルプログラム全体を示します。

#define TBB_PREVIEW_FLOW_GRAPH_NODES 1
#include "tbb/flow_graph_opencl_node.h"

#include <numeric>
#include <cassert>

int main() {
    using namespace tbb::flow;
    
    opencl_graph g;
    const int N = 1000;
    opencl_buffer<cl_int> buf( g, N );
    std::iota( buf.begin(), buf.end(), 0 );
    
    opencl_node<tuple<opencl_buffer<cl_int>>> 
        clNode( g, "opencl_buffer.cl", "increment" );
            
    std::vector<int> ndrange( 1, N );
    clNode.set_ndranges( ndrange );
    input_port<0>(clNode).try_put( buf );
    
    g.wait_for_all();
    
    assert( std::accumulate( buf.begin(), buf.end(), 0 ) == N*(N+1)/2 );
    
    return 0;
}

opencl_buffer.cl:

kernel void increment ( global int* arr ) { 
    const int i = get_global_id(0);
    arr[i] += 1;

このサンプルは、整数配列に 0 から 999 を格納し、opencl_node を用いて各要素を並列に 1 インクリメントして、assert で結果を確認します。

このシリーズの次の記事では、カーネルの実行に使用するデバイスの選択について述べます。

コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください

タイトルとURLをコピーしました