ここ数日でOpenCLについて学んだこと。
- 各ベンダはOpenCLのPlatformを提供する。
- Platformはひとつ以上のDeviceを持つ。
- 特定のDeviceで実行する環境をContextという。
- Deviceに与える(OpenCL Cで書かれた)Programはひとつ以上のエントリ(関数の入り口)を持つ。
- ひとつのエントリとそれに与える引数のセットをKernelという。
- Device内のBuffer(メモリ領域)への読み書き、
そしてKernelによる実行などをCommandQueueを介して行う。
これらを踏まえてこしらえた習作、お題は「並列処理が向かないかもしれない例」の奇遇転置ソート:
1 7 6 5 8 2 4 3 があったとき、
[1] ふたつずつのペアをつくる。1 7 6 5 8 2 4 3
[2] それぞれのペアについて必要なら入れ替え 1 7 5 6 2 8 3 4
[3] ひとつずらしてペアをつくる 1 7 5 6 2 8 3 4
[4] それぞれのペアについて必要なら入れ替え1 5 7 2 6 3 8 4
[5] [1]～[4] をくりかえし。
[2],[4] の際、ペアの入れ替えは互いに独立だから並行処理できんぢゃん、と。
なんたってGPUですからね、コアはたっぷりあります。
目下愛用のRADEON HD5670、シェーダ・プロセッサを400基載せてます。
#include <iostream> // cout, endl
#include <array> // array
#include <vector> // vector
#include <map> // map
#include <string> // string
#include <numeric> // iota
#include <algorithm> // copy, for_each
#include <iterator> // back_inserter
#include <utility> // make_pair
#include <cassert> // assert
#include <Windows.h> // GetTickCount
// error時は例外処理に一任する
#define __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp>
using namespace std;
int main() {
try {
// Platformを列挙し、ひとつ目を使う
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
assert( !platforms.empty() );
cl::Platform platform = platforms[0];
cout << "Platform: " << platform.getInfo<CL_PLATFORM_NAME>() << '\t';
// 最初に見つかったPlatformからデバイスを探し、GPU-Contextを生成
cl_context_properties cprops[3] = { CL_CONTEXT_PLATFORM,(cl_context_properties)platform(), 0};
cl::Context context(CL_DEVICE_TYPE_GPU,cprops);
// Context中のDeviceを列挙し、ひとつ目を使う
vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
assert( !devices.empty() );
cl::Device device = devices[0];
cout << "Device: " << device.getInfo<CL_DEVICE_NAME>() << endl;
// DeviceにCommandQueueを用意
cl::CommandQueue queue(context,device);
// Programを準備 : 奇遇ソート
static string source_strs[] = {
"__kernel void swap_even( \n"
" __global float* data, \n"
" __global int* swapped) \n"
"{ \n"
" int i = get_global_id(0); \n"
" int j = i + 1; \n"
" if ( !(i & 1) ) { \n"
" if ( data[i] > data[j] ) { \n"
" float tmp = data[i]; \n"
" data[i] = data[j]; \n"
" data[j] = tmp; \n"
" *swapped = 1; \n"
" } \n"
" } \n"
"} \n"
,
"__kernel void swap_odd( \n"
" __global float* data, \n"
" __global int* swapped) \n"
"{ \n"
" int i = get_global_id(0); \n"
" int j = i + 1; \n"
" if ( (i & 1) ) { \n"
" if ( data[i] > data[j] ) { \n"
" float tmp = data[i]; \n"
" data[i] = data[j]; \n"
" data[j] = tmp; \n"
" *swapped = 1; \n"
" } \n"
" } \n"
"} \n"
,
""
};
cl::Program::Sources sources;
for ( string* str = source_strs; !str->empty(); ++str ) {
sources.push_back(make_pair(str->c_str(), str->size()));
}
// Programのビルド(失敗時にはログを出力)
cl::Program program(context, sources);
try {
program.build(devices);
} catch ( cl::Error& ) {
cerr << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device) << endl;
return -1;
}
// ProgramからKernelを生成
map<string,cl::Kernel> kernels;
{
vector<cl::Kernel> knls;
program.createKernels(&knls);
// 関数名とKernelとの対応表をつくる
for_each(knls.begin(),knls.end(),
[&](cl::Kernel k) { kernels[k.getInfo<CL_KERNEL_FUNCTION_NAME>()] = k; });
}
// ホストに要素数Nの配列を用意
const int N = 10000;
array<float,N> data;
int swapped;
// data[] = { N-1, N-2, ... 2, 1 }
iota(data.begin(), data.end(), 1.0f);
reverse(data.begin(),data.end());
// ふつーにやればこんなの。
DWORD t = GetTickCount();
do {
swapped = 0;
for ( int i = 0; i < N-1; i += 2 ) {
if ( data[i] > data[i+1] ) {
float tmp = data[i]; data[i] = data[i+1]; data[i+1] = tmp;
swapped = 1;
}
}
for ( int i = 1; i < N-1; i += 2 ) {
if ( data[i] > data[i+1] ) {
float tmp = data[i]; data[i] = data[i+1]; data[i+1] = tmp;
swapped = 1;
}
}
} while ( swapped );
t = GetTickCount() - t;
assert( is_sorted(data.begin(), data.end()) );
if ( is_sorted(data.begin(),data.end()) ) {
cout << "CPU DONE. " << t << "[ms]" << endl;
}
reverse(data.begin(),data.end());
// Device側にバッファを確保
cl::Buffer buffer_data(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
sizeof(data) , &data[0]);
cl::Buffer buffer_swapped(context, CL_MEM_WRITE_ONLY,
sizeof(swapped), nullptr);
// 2つのKernelに引数を設定
cl::Kernel kernel_0 = kernels["swap_even"];
kernel_0.setArg(0,buffer_data);
kernel_0.setArg(1,buffer_swapped);
cl::Kernel kernel_1 = kernels["swap_odd"];
kernel_1.setArg(0,buffer_data);
kernel_1.setArg(1,buffer_swapped);
t = GetTickCount();
do {
// swappedを0(false)に
swapped = 0;
queue.enqueueWriteBuffer(buffer_swapped,CL_TRUE, 0, sizeof(swapped), &swapped);
// 偶数
queue.enqueueNDRangeKernel(kernel_0, cl::NullRange, cl::NDRange(N-1), cl::NullRange);
queue.enqueueBarrier();
// 奇数
queue.enqueueNDRangeKernel(kernel_1, cl::NullRange, cl::NDRange(N-1), cl::NullRange);
queue.enqueueBarrier();
// swappedを取り出し
queue.enqueueReadBuffer(buffer_swapped, CL_TRUE, 0, sizeof(swapped), &swapped);
} while ( swapped );
t = GetTickCount() - t;
// 結果を取り出して検証
queue.enqueueReadBuffer(buffer_data, true, 0, sizeof(data), &data[0]);
assert( is_sorted(data.begin(), data.end()) );
if ( is_sorted(data.begin(), data.end()) ) {
cout << "GPU DONE. " << t << "[ms] " << endl;
}
} catch ( cl::Error& ex ) {
cerr << ex.err() << ':' << ex.what() << endl;
return -1;
}
}
実行結果:
Platform: ATI Stream Device: Redwood
CPU DONE. 156[ms]
GPU DONE. 2886[ms]
あ、あるえぇぇ...? ぜんっぜん速くねぇ。
赤いトコがGPU内でチョー並列に動いてはくれるんだけども、
それを呼んでる青いトコはホスト側にあります。
青いトコは並列でもなんでもなく、赤いトコがいっくら速くても
青いトコがそれを相殺してるように思えます。
GPGPUでかっとびアプリを書くにはまだまだ修行が足らんなーと。
まぁ二、三日でここまで書けるよになったんだから僕の感覚もそんなに鈍っちゃいないぜ、と。