この記事は、インテル® デベロッパー・ゾーンに公開されている「Device selection」(https://software.intel.com/en-us/blogs/2015/12/16/device-selection/) の日本語参考訳です。
この記事は、インテル® スレッディング・ビルディング・ブロック (インテル® TBB) 4.4 Update 2 以降で利用可能な新しいノード opencl_node について説明するシリーズのパート 3 です。このノードは、インテル® TBB のフローグラフで OpenCL* デバイスの利用と連携を可能にします。このシリーズのパート 1 はこちらからご覧いただけます。
パート 2 では、基本インターフェイスを説明しました。パート 3 では、カーネルの実行に使用するデバイスの選択について述べます。
デバイスの選択
opencl_node では、実行デバイスの選択方法がいくつかあります。
- 適切なデバイスのリストで default_opencl_factory を初期化する
- opencl_node の作成時にデバイスセレクターを指定する
- カスタム OpenCL* factory を作成する
次の例は、利用可能な OpenCL* デバイスのリストを取得し、標準出力へ出力します。
#define TBB_PREVIEW_FLOW_GRAPH_NODES 1
#include "tbb/flow_graph_opencl_node.h"
#include <iostream>
int main() {
using namespace tbb::flow;
opencl_graph g;
opencl_device_list dl = g.available_devices();
for ( opencl_device d : dl ) {
std::cout << "Device: " << d.name() << std::endl;
std::cout << " Platform: " << d.platform_name() << std::endl;
}
return 0;
}
私のテストマシンでは、次の出力が得られました。
Device: Intel(R) HD Graphics 4400 Platform: Intel(R) OpenCL Device: Intel(R) Core(TM) i5-4300U CPU @ 1.90GHz Platform: Intel(R) OpenCL Device: Intel(R) Core(TM) i5-4300U CPU @ 1.90GHz Platform: Experimental OpenCL 2.0 CPU Only Platform
opencl_device クラスには、OpenCL* デバイスに関する情報を照会するためのいくつかのメソッドがあります。
class opencl_device {
public:
std::string platform_profile(); // CL_PLATFORM_PROFILE
std::string platform_version(); // CL_PLATFORM_VERSION
std::string platform_name(); // CL_PLATFORM_NAME
std::string platform_vendor(); // CL_PLATFORM_VENDOR
std::string platform_extensions(); // CL_PLATFORM_EXTENSIONS
template <typename T>
void info( cl_device_info i, T &t );
std::string version(); // CL_DEVICE_VERSION
std::string extensions(); // CL_DEVICE_EXTENSIONS
// extensions() の出力を解析
bool extension_available( const std::string &ext );
std::string name(); // CL_DEVICE_NAME
cl_device_type type(); // CL_DEVICE_TYPE
};
最も一般的なメソッドを次に示します。
template <typename T> void info( cl_device_info i, T &t );
このメソッドを使用して OpenCL* デバイスに関する情報を照会できます。
opencl_graph g; opencl_device d = *g.available_devices().begin(); std::string profile; d.info( CL_DEVICE_PROFILE, profile );
cl_device_info の値については、clGetDeviceInfo のドキュメントを参照してください。
特定のデバイスでデフォルトの factory を初期化するには、init メソッドを使用します。
g.opencl_factory().init( {*dl.begin()} );
上記の例は、最初に利用可能なデバイスで factory を初期化します。これはデフォルトの動作でもあります。上記の例から、特定のデバイスで OpenCL* factory を初期化する方法が分かります。
factory の初期化に複数のデバイスを指定することもできますが、同じプラットフォームのデバイスでなければなりません。次に例を示します。
opencl_device_list dl = g.available_devices();
opencl_device_list chosenDevices;
for ( opencl_device d : dl ) {
if ( d.platform_name() == "Intel(R) OpenCL" )
chosenDevices.add( d );
}
g.opencl_factory().init( chosenDevices );
1 つの factory に複数のデバイスが指定されている場合、デフォルトでは opencl_node は最初のデバイスのみ使用します。opencl_node の作成時にユーザー定義のデバイスセレクターを指定することで、各カーネル呼び出しごとに個別のデバイスを選択することもできます。デバイスセレクターは、(各メッセージごとではなく) 1 パックの入力メッセージにつき 1 回のみ呼び出されます。次に、ラウンドロビン方式で複数のデバイスから選択する方法を示します。
#define TBB_PREVIEW_FLOW_GRAPH_NODES 1
#include "tbb/flow_graph_opencl_node.h"
#include <atomic>
#include <iostream>
int main() {
using namespace tbb::flow;
opencl_graph g;
opencl_device_list dl = g.available_devices();
opencl_device_list chosenDevices;
for ( opencl_device d : dl ) {
if ( d.platform_name() == "Intel(R) OpenCL" )
chosenDevices.add( d );
}
g.opencl_factory().init( chosenDevices );
std::atomic<int> deviceNum = 0;
opencl_node<tuple<cl_int>> clNode( g, "multiple_devices.cl", "foo",
[&deviceNum](const opencl_device_list &dl) {
return *(dl.begin() + deviceNum++ % dl.size());
}
);
clNode.set_ndranges( { 10 } );
for ( int i = 0; i < 10; ++i )
input_port<0>(clNode).try_put( 0 );
g.wait_for_all();
std::cout << std::endl;
return 0;
}
multiple_devices.cl:
kernel void foo ( int c ) {
if ( get_global_id(0) == 0 ) {
int num = get_num_groups(0);
printf("%d ", num);
}
}
カーネルは、(デバイスの特性に応じてワークの分割状況が異なるため) デバイスを区別できるようにワークグループの数を出力します。ただし、過度な出力を避けるため、最初のワークのみこの情報を出力します。
次のような出力が得られます。
5 1 5 1 5 1 5 1 5 1
異なる数値が出力されていることから、複数のデバイスで実行されていることが分かります。
この default_opencl_factory の初期化アプローチには、複数のプラットフォームに同じグラフを使用できないという制限があります。この制限に対応するには、カスタム factory を使用します。
フローグラフには、カスタム factory としてインスタンス化できる opencl_factory テンプレート・クラスがあります。
template <typename DeviceFilter> class opencl_factory;
opencl_factory テンプレート・クラスには、ユーザー定義の DeviceFilter のみ必要です。opencl_factory の作成時に、必要なデバイスを選択するためこのフィルターが呼び出されます。default_opencl_factory とは異なり、必要に応じて複数の opencl_factory インスタンスを作成できます。次の例は、2 つのインスタンスを作成します。1 つ目は OpenCL* 1.2 をサポートするプラットフォーム用で、2 つ目は OpenCL* 2.0 をサポートするプラットフォーム用です。
#define TBB_PREVIEW_FLOW_GRAPH_NODES 1
#include "tbb/flow_graph_opencl_node.h"
#include <algorithm>
#include <string>
using namespace tbb::flow;
struct ChooseDevice_1_2 {
opencl_device_list operator()( opencl_device_list dl ) const {
return { *std::find_if( dl.begin(), dl.end(),
[]( opencl_device d ) -> bool {
const std::string ver = "OpenCL 1.2";
return d.version().compare(0, ver.length(), ver) == 0;
} )
};
}
};
int main() {
typedef opencl_factory<ChooseDevice_1_2> F12;
auto ChooseDevice_2_0 = [](opencl_device_list dl) -> opencl_device_list {
return { *std::find_if( dl.begin(), dl.end(),
[]( opencl_device d ) -> bool {
const std::string ver = "OpenCL 2.0";
return d.version().compare(0, ver.length(), ver) == 0;
} )
};
};
typedef opencl_factory<decltype(ChooseDevice_2_0)> F20;
opencl_graph g;
F12 f12(g);
F20 f20(g);
opencl_node<tuple<cl_int>, queueing, F12> clNode12( g, "multiple_factories.cl", "version", f12 );
opencl_node<tuple<cl_int>, queueing, F20> clNode20( g, "multiple_factories.cl", "version", f20 );
clNode12.set_ndranges( { 1 } );
clNode20.set_ndranges( { 1 } );
input_port<0>(clNode12).try_put( 0 );
input_port<0>(clNode20).try_put( 0 );
g.wait_for_all();
return 0;
}
multiple_factories.cl:
kernel void version ( int c ) {
const int v = __OPENCL_VERSION__;
const int major = v / 100;
const int minor = v % 100 / 10;
printf("OpenCL version is %d.%d\n", major, minor);
}
カーネルは、次のように OpenCL* のバージョンを出力します。
OpenCL version is 2.0 OpenCL version is 1.2
ユーザー定義のデバイスセレクターを custom factory と併用することもできます。opencl_node コンストラクターの引数リストで factory object よりも先に指定します。次に例を示します。
opencl_node<tuple<cl_int>, queueing, F12> clNode12( g,
"multiple_factories.cl", "version",
[](const opencl_device_list &dl) {
return *dl.begin();
}, f12 );
異なる factory で opencl_nodes をインスタンス化した場合、それぞれの factory は個別の OpenCL* コンテキストを作成し、異なる OpenCL* コンテキストのリソースを直接共有することは不可能なため、インスタンス化した opencl_nodes は make_edge インターフェイスで接続できないことに注意してください。さらに、対応するカスタム factory 型とオブジェクトで、すべてのメモリー・オブジェクトをインスタンス化する必要があります。
opencl_buffer<cl_float, F12> b12( f12, size ); opencl_buffer<cl_float, F20> b20( f20, size );
異なる factory で作成された opencl_node にバッファーを渡すことはできません。そのため、異なる factory でインスタンス化された 2 つの opencl_node を make_edge インターフェイスで接続することはできません。あるバッファーから別のバッファーへデータをコピー/転送する場合は、中間の function_node を使用すべきです。ただし、オーバーヘッドが増える可能性があります。
このシリーズの次の記事では、OpenCL* プログラムを指定し、プログラムからカーネルを選択して、カーネル呼び出しに引数をバインドする方法を説明します。
コンパイラーの最適化に関する詳細は、最適化に関する注意事項を参照してください
