FC2ブログ

スポンサーサイト

上記の広告は1ヶ月以上更新のないブログに表示されています。
新しい記事を書く事で広告が消せます。

OpenCL の Image2D メモリを使ってみる

OpenCL では,標準的な1次元の配列の他に,2,3次元のデータを取り扱う Image2D,Image3D というデータ型を利用できます.1次元の配列を用いて2,3次元の配列を表現することもできますが,専用のデータ型を用いるとで座標による直観的なアクセスが可能です.ただし,メモリの確保は1次元の配列に比べると若干面倒になります.本記事では Image2D オブジェクトを確保し,画像データにアクセスする方法を説明します.

Keywords: OpenCL, Image2D


OpenCL を利用するためには,Context を作成したり Program を読み込んでビルドしたりといった,おまじないじみた作業が必要になります.このような作業については本記事では説明しません.私は以下の記事を参考にしましたので紹介させていただきます.
THE KERNEL MAGIC
OpenGL de プログラミング


1.Image2D オブジェクトを作成する

clCreateImage2D 関数を用いて Image2D オブジェクトを作成します.ここでは16x16の1チャンネル Intensity 画像を指定します.
static cl_int CreateImage2D(
cl_context context,
size_t image_width,
size_t image_height,
cl_mem *image
)
{
cl_image_format image_format;
image_format.image_channel_order = CL_INTENSITY;
image_format.image_channel_data_type = CL_UNORM_INT8;

cl_int err;
cl_mem m = clCreateImage2D(
context,
CL_MEM_READ_WRITE,
&image_format,
image_width,
image_height,
0,
NULL,
&err
);
if(err != CL_SUCCESS) {
printf("Failed creating Image2D memory object\n");
return err;
}

*image = m;

return CL_SUCCESS;
}
cl_image_format 構造体で,作成する Image オブジェクトのフォーマットを指定します.上記のように指定すると,ホストからデータを読み書きするときには1ピクセルあたり1byteの1チャンネル画像として,カーネル(デバイス側のプログラム)からは[0,1]に正規化された4チャンネルの画像として取り扱うことができます.OpenCL仕様書の247ページに Channel Order ごとの値の展開のされ方が,264ページに Channel Data Type による値の変換についてが記述されているので参考にしてください.
clCreateImage2D の第2引数には,カーネル側からこのメモリに実行できるオペレーションを指定します.ここでは読み書き両方を許可しています.


2.カーネルの引数と関連付ける

clSetKernelArg 関数を用いて,カーネル関数の引数と確保した Image2D オブジェクトを関連付けます.
cl_kernel kernel;
cl_mem imsrc, imdst;
clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&imsrc);
clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&imdst);
imsrc,imdst は先ほど作成した Image2D オブジェクトになります.


3.ホストからデバイスへデータを転送する

clEnqueueWriteImage 関数を用いて,ホストからデバイスへデータを転送します.Image2D と Image3D を同じ関数で取り扱うため,一部の引数が冗長になっています.
static cl_int WriteImage2D(
cl_command_queue command,
cl_mem image,
size_t image_width,
size_t image_height,
size_t image_row_pitch,
const void *host_ptr
)
{
size_t origin[3] = {0};
size_t region[3] = {image_width, image_height, 1};
cl_int err = clEnqueueWriteImage(
command,
image,
CL_TRUE,
origin,
region,
image_row_pitch,
0,
host_ptr,
0,
NULL,
NULL
);
if(err != CL_SUCCESS) {
printf("Failed writing data to device\n");
return err;
}

return CL_SUCCESS;
}
第2引数に転送先の Image2D オブジェクトを指定します.第3引数は書き込み操作がおわるまでこの関数をブロックするかのフラグで,ここではブロックしています.第8引数が書き込むデータブロックのポインタで,今回はあらかじめ定義した
unsigned char[256] = {0,1,2, ... 254,255};
となるデータを書き込んでいます.


4.カーネルを実行する

clEnqueueNDRangeKernel 関数を呼び出して,カーネルを実行します.global_size を16x16とすることで,カーネル側でのコピー作業を並列で行っています.
static cl_int ExecuteKernel(
cl_command_queue command,
cl_kernel kernel,
const size_t global_size[2],
const size_t local_size[2]
)
{
cl_int err = clEnqueueNDRangeKernel(
command,
kernel,
2,
NULL,
global_size,
local_size,
0,
NULL,
NULL
);
if(err != CL_SUCCESS) {
printf("Failed executing kernel on device\n");
return err;
}

return CL_SUCCESS;
}


5.データをデバイスから読み出す

clEnqueueReadImage 関数を用いて,デバイスからデータを読み出します.書き込み操作のときとほとんど同じです.第8引数には,あらかじめ確保済みのデータブロックを渡してください.
static cl_int ReadImage2D(
cl_command_queue command,
cl_mem image,
size_t image_width,
size_t image_height,
size_t image_row_pitch,
void *host_ptr
)
{
size_t origin[3] = {0};
size_t region[3] = {image_width, image_height, 1};
cl_int err = clEnqueueReadImage(
command,
image,
CL_TRUE,
origin,
region,
image_row_pitch,
0,
host_ptr,
0,
NULL,
NULL
);
if(err != CL_SUCCESS) {
printf("Failed reading data from device\n");
return err;
}

return CL_SUCCESS;
}


6.カーネルプログラム

プログラムを実行すると,次のような結果が出力されます.
OpenCL Image2D and Kernel Execution Sample

[global (16, 16), local (1 1)]
0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31
32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47
48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63
64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79
80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95
96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111
112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127
128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143
144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159
160 161 162 163 164 165 166 167 168 169 170 171 172 173 174 175
176 177 178 179 180 181 182 183 184 185 186 187 188 189 190 191
192 193 194 195 196 197 198 199 200 201 202 203 204 205 206 207
208 209 210 211 212 213 214 215 216 217 218 219 220 221 222 223
224 225 226 227 228 229 230 231 232 233 234 235 236 237 238 239
240 241 242 243 244 245 246 247 248 249 250 251 252 253 254 255
カーネルプログラムで,入力イメージから出力イメージにピクセルをコピーできていることが確認できます.下にカーネルプログラムのコードを示します.

CLCopy.cl
const sampler_t sampler =
CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST;

__kernel void CLCopy(read_only image2d_t a, write_only image2d_t b)
{
int2 pos;

pos.x = get_global_id(0);
pos.y = get_global_id(1);

if(pos.x >= 16 || pos.y >= 16) return;

float4 pix = read_imagef(a, sampler, pos);

write_imagef(b, pos, pix);
}


付録

CLImage2D.cpp
#include <stdio.h>
#include <stdlib.h>

#include <fstream>

#include <cl/cl.h>
#include <cl/cl_ext.h>


using std::string;
using std::ifstream;
using std::istreambuf_iterator;


static cl_platform_id *platform_ids = NULL;
static cl_device_id *device_ids = NULL;

static cl_context context = NULL;
static cl_command_queue command = NULL;
static cl_program program = NULL;
static cl_kernel kernel = NULL;

static cl_mem imsrc = NULL;
static cl_mem imdst = NULL;


static const unsigned char zeros[256] = {0};

static const unsigned char sample[] = {
0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,
16,17,18,19,20,21,22,23,24,25,26,27,28,29,30,31,
32,33,34,35,36,37,38,39,40,41,42,43,44,45,46,47,
48,49,50,51,52,53,54,55,56,57,58,59,60,61,62,63,
64,65,66,67,68,69,70,71,72,73,74,75,76,77,78,79,
80,81,82,83,84,85,86,87,88,89,90,91,92,93,94,95,
96,97,98,99,100,101,102,103,104,105,106,107,108,109,110,111,
112,113,114,115,116,117,118,119,120,121,122,123,124,125,126,127,
128,129,130,131,132,133,134,135,136,137,138,139,140,141,142,143,
144,145,146,147,148,149,150,151,152,153,154,155,156,157,158,159,
160,161,162,163,164,165,166,167,168,169,170,171,172,173,174,175,
176,177,178,179,180,181,182,183,184,185,186,187,188,189,190,191,
192,193,194,195,196,197,198,199,200,201,202,203,204,205,206,207,
208,209,210,211,212,213,214,215,216,217,218,219,220,221,222,223,
224,225,226,227,228,229,230,231,232,233,234,235,236,237,238,239,
240,241,242,243,244,245,246,247,248,249,250,251,252,253,254,255
};


template<typename T>
void SafeRelease(T *pointer, cl_int (CL_API_CALL *func)(T*))
{
if(pointer != NULL) func(pointer);
}

template<typename T>
void SafeFree(T **pointer)
{
free(*pointer);
*pointer = NULL;
}


static cl_uint GetNumPlatforms()
{
cl_uint numPlatforms = 0;
cl_int result = clGetPlatformIDs(0, NULL, &numPlatforms);
if(result != CL_SUCCESS) {
printf("Failed getting number of available platforms\n");
}
return numPlatforms;
}

static cl_platform_id* GetPlatformIDs(cl_uint numPlatforms)
{
size_t allocateSize = sizeof(cl_platform_id) * numPlatforms;

cl_platform_id *ids = (cl_platform_id*)malloc(allocateSize);
if(ids != NULL) {
cl_int result = clGetPlatformIDs(numPlatforms, ids, NULL);
if(result != CL_SUCCESS) {
printf("Failed getting platform IDs\n");
SafeFree(&ids);
}
}
return ids;
}

static cl_platform_id* GetPlatformIDs()
{
cl_platform_id *ids = NULL;
cl_uint numPlatforms = GetNumPlatforms();
if(numPlatforms > 0) {
ids = GetPlatformIDs(numPlatforms);
}
return ids;
}


static cl_uint GetNumDevices(cl_platform_id id)
{
cl_uint numDevices = 0;
cl_int result = clGetDeviceIDs(
id, CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices
);
if(result != CL_SUCCESS) {
printf("Failed getting number of available devices\n");
}
return numDevices;
}

static cl_device_id* GetDeviceIDs(cl_platform_id id, cl_uint numDevices)
{
size_t allocateSize = numDevices * sizeof(cl_device_id);
cl_device_id* ids = (cl_device_id*)malloc(allocateSize);
if(ids != NULL) {
cl_int result = clGetDeviceIDs(
id, CL_DEVICE_TYPE_ALL, numDevices, ids, NULL
);
if(result != CL_SUCCESS) {
printf("Failed getting device IDs\n");
SafeFree(ids);
}
}
return ids;
}

static cl_device_id* GetDeviceIDs(cl_platform_id id)
{
cl_device_id *ids = NULL;
cl_uint numDevices = GetNumDevices(id);
if(numDevices > 0) {
ids = GetDeviceIDs(id, numDevices);
}
return ids;
}


static cl_int BuildProgram(
const char *filename,
cl_context context,
cl_program *program
)
{
cl_int err;

ifstream source(filename);
istreambuf_iterator<char> dataBegin(source);
istreambuf_iterator<char> dataEnd;

string code(dataBegin, dataEnd);

const char *program_code = code.c_str();
const size_t program_length = code.size();
cl_program p = clCreateProgramWithSource(
context,
1,
&program_code,
&program_length,
&err
);
if(err != CL_SUCCESS) {
printf("Failed creating program object from source\n", err);
return err;
}

*program = p;

err = clBuildProgram(p, 0, NULL, NULL, NULL, NULL);
if(err != CL_SUCCESS) {
printf("Failed building program object\n", err);
return err;
}

return CL_SUCCESS;
}


static cl_int CreateImage2D(
cl_context context,
size_t image_width,
size_t image_height,
cl_mem *image
)
{
cl_image_format image_format;
image_format.image_channel_order = CL_INTENSITY;
image_format.image_channel_data_type = CL_UNORM_INT8;

cl_int err;
cl_mem m = clCreateImage2D(
context,
CL_MEM_READ_WRITE,
&image_format,
image_width,
image_height,
0,
NULL,
&err
);
if(err != CL_SUCCESS) {
printf("Failed creating Image2D memory object\n");
return err;
}

*image = m;

return CL_SUCCESS;
}


static cl_int WriteImage2D(
cl_command_queue command,
cl_mem image,
size_t image_width,
size_t image_height,
size_t image_row_pitch,
const void *host_ptr
)
{
size_t origin[3] = {0};
size_t region[3] = {image_width, image_height, 1};
cl_int err = clEnqueueWriteImage(
command,
image,
CL_TRUE,
origin,
region,
image_row_pitch,
0,
host_ptr,
0,
NULL,
NULL
);
if(err != CL_SUCCESS) {
printf("Failed writing data to device\n");
return err;
}

return CL_SUCCESS;
}


static cl_int ReadImage2D(
cl_command_queue command,
cl_mem image,
size_t image_width,
size_t image_height,
size_t image_row_pitch,
void *host_ptr
)
{
size_t origin[3] = {0};
size_t region[3] = {image_width, image_height, 1};
cl_int err = clEnqueueReadImage(
command,
image,
CL_TRUE,
origin,
region,
image_row_pitch,
0,
host_ptr,
0,
NULL,
NULL
);
if(err != CL_SUCCESS) {
printf("Failed reading data from device\n");
return err;
}

return CL_SUCCESS;
}


static cl_int ExecuteKernel(
cl_command_queue command,
cl_kernel kernel,
const size_t global_size[2],
const size_t local_size[2]
)
{
cl_int err = clEnqueueNDRangeKernel(
command,
kernel,
2,
NULL,
global_size,
local_size,
0,
NULL,
NULL
);
if(err != CL_SUCCESS) {
printf("Failed executing kernel on device\n");
return err;
}

return CL_SUCCESS;
}


static void Shutdown()
{
SafeRelease(imsrc, clReleaseMemObject);
SafeRelease(imdst, clReleaseMemObject);
SafeRelease(kernel, clReleaseKernel);
SafeRelease(program, clReleaseProgram);
SafeRelease(command, clReleaseCommandQueue);
SafeRelease(context, clReleaseContext);
SafeFree(&device_ids);
SafeFree(&platform_ids);
exit(0);
}


static void CheckError(const cl_int error)
{
if(error != CL_SUCCESS) Shutdown();
}


static void Print(unsigned char image[256])
{
int i, j;
for(j=0; j<16; j++) {
for(i=0; i<16; i++) {
printf("%3d ", image[i + j * 16]);
}
printf("\n");
}
}


int main(int, char**) {
cl_int err;

printf("OpenCL Image2D and Kernel Execution Sample\n");
printf("\n");

platform_ids = GetPlatformIDs();
if(platform_ids == NULL) {
Shutdown();
}

device_ids = GetDeviceIDs(platform_ids[0]);
if(device_ids == NULL) {
Shutdown();
}

// OpenCL Context
context = clCreateContext(NULL, 1, device_ids, NULL, NULL, &err);
CheckError(err);

// OpenCL Command Queue
command = clCreateCommandQueue(context, device_ids[0], 0, &err);
CheckError(err);

// OpenCL Program Build
err = BuildProgram("CLCopy.cl", context, &program);
CheckError(err);

// OpenCL Kernel Object
kernel = clCreateKernel(program, "CLCopy", &err);
CheckError(err);

// Memory Allocation for Image
err = CreateImage2D(context, 16, 16, &imsrc);
err |= CreateImage2D(context, 16, 16, &imdst);
CheckError(err);

err = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&imsrc);
err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&imdst);
CheckError(err);

// Write Image Data to Device
err = WriteImage2D(command, imsrc, 16, 16, 16, (void*)sample);
CheckError(err);

// Execute
const size_t size_global[] = {16, 16};
const size_t size_local[] = { 1, 1};
err = ExecuteKernel(command, kernel, size_global, size_local);
CheckError(err);

// Read Result Image
unsigned char result[256];
err = ReadImage2D(command, imdst, 16, 16, 16, (void*)result);
if(err == CL_SUCCESS) {
printf("[global (%d, %d), local (%d %d)]\n",
size_global[0], size_global[1],
size_local[0], size_local[1]);
Print(result);
printf("\n");
}

Shutdown();
}
スポンサーサイト

テーマ : プログラミング
ジャンル : コンピュータ

コメントの投稿

非公開コメント

プロフィール

Ishida Akihiko

Author:Ishida Akihiko
FC2ブログへようこそ!

免責事項
当サイトに掲載する記事内容は,必ずしも正確性,信頼性,妥当性,有用性,完成度などを保証しません.記事の利用はすべて自己責任でお願いします.当サイトに掲載された内容によって発生したいかなる損害に対しても,管理人は一切の責任を負いかねます.
最新記事
最新コメント
最新トラックバック
月別アーカイブ
カテゴリ
アクセスカウンター
検索フォーム
RSSリンクの表示
リンク
ブロとも申請フォーム

この人とブロともになる

QRコード
QR
上記広告は1ヶ月以上更新のないブログに表示されています。新しい記事を書くことで広告を消せます。