図9.2 SAMPLE.png

図9.3 SAMPLE3.png

この章でのプロジェクトの構造は以下のようになります。ファイルの出力はtarget/classesのパッケージ内とします。サンプルコードでは、「image.png」というファイルを出力しますが、target/classes下に出力されています。

図：ディレクトリ構造. 

ImgProject
├── pom.xml
├── resources
├── src
│   ├── main
│   │   └── java
│   │       └── com
│   │           └── book
│   │               └── jocl
│   │                   └── image
│   │                       ├── ImageTest.java
│   │                       ├── SAMPLE.png
│   │                       ├── SAMPLE3.png
│   │                       └── image.cl
│   └── test
│       └── java
└── target
    ├── classes
    │   └── com
    │       └── book
    │           └── jocl
    │               └── image
    │                   ├── ImageReadWriteTest.class
    │                   ├── ImageTest.class
    │                   ├── SAMPLE.png
    │                   ├── SAMPLE3.png
    │                   ├── image.cl
    │                   └── image.png
    └── test-classes

clCreateImage関数は1Dイメージ、1Dイメージバッファ、1Dイメージ配列、2Dイメージ、2Dイメージ配列、3Dイメージオブジェクトを生成します。

関数の定義は以下のようになります。

cl_mem org.jocl.CL.clCreateImage(
    cl_context context, //(1)
    long flags, //(2)
    cl_image_format image_format, //(3)
    cl_image_desc image_desc, //(4)
    Pointer host_ptr, //(5)
    int[] errcode_ret) //(6)

引数にはcl_image_foramt、cl_image_descという構造体があります。

では例を見てみましょう。

cl_image_desc desc = new cl_image_desc();
desc.image_type = CL_MEM_OBJECT_IMAGE2D; //(1)
desc.image_width = img.getIconWidth(); //(2)
desc.image_height = img.getIconHeight(); //(3)
desc.num_mip_levels = 0;
desc.num_samples = 0;

cl_image_format format = new cl_image_format();
format.image_channel_order = CL_RGBA; //(4)
format.image_channel_data_type = CL_UNORM_INT8; //(5)


cl_mem img_mem = clCreateImage(context,
        CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
        format, //(6)
        desc, //(7)
        Pointer.to(dbb.getData()), //(8)
        null);

この例では2次元のイメージを既に読み込んだ事を前提とします。また幅と高さについても既知の情報とします。つまりイメージオブジェクトをclCreateImage関数を用いて生成する際のは、イメージを読み込み後となります。

チャネル順序とチャネルデータ型についても既知の情報とします。チャネル順序についてはJava画像処理パッケージが取得することができます。

チャネルデータ型についてはread_imagefで取り扱える正規化符号無し整数（CL_UNORM_INT8）を指定します。最後の8はビット数を示します。イメージのビット数（libpngではIHDRのbit_depthで取得可能）に合わせることと、浮動小数点としてアクセスしたいか、整数としてアクセスしたいかによってデータ型を使い分けることになります。

それでは実装をしてみたいと思います。以下のソースコードでは、SAMPLE3.pngというファイルをJavaの標準パッケージを使い読みこみます。読み込んだデータをclCreateImage関数でイメージオブジェクトを生成します。

さらにイメージオブジェクトをカーネルで処理後に、clEnqueueReadImage関数を使って、イメージオブジェクトからホストのメモリ領域にデータを読み込みます。

最後にImageIO.writeという関数を使って、ホストメモリ領域にあるイメージデータをimage.pngファイルに書き込みます。

package com.book.jocl.image;

import static org.jocl.CL.*;

import org.jocl.CL;
import org.jocl.Pointer;
import org.jocl.Sizeof;
import org.jocl.cl_command_queue;
import org.jocl.cl_context;
import org.jocl.cl_context_properties;
import org.jocl.cl_device_id;
import org.jocl.cl_image_desc;
import org.jocl.cl_image_format;
import org.jocl.cl_kernel;
import org.jocl.cl_mem;
import org.jocl.cl_platform_id;
import org.jocl.cl_program;

import java.awt.image.BufferedImage;
import java.awt.image.DataBufferInt;
import java.awt.*;
import javax.swing.*;
import javax.imageio.*;
import java.io.*;
import java.net.URL;
import java.nio.file.Paths;
import java.util.Scanner;


public class ImageTest {
        private static final String IMAGE_PATH = "SAMPLE.png";
        private static final String KERNEL_PATH = "image.cl";
        private static final String KERNEL_FUNC = "image_filter";

    private static cl_context context;
    private static cl_command_queue queue;
    private static cl_program program;
    private static cl_kernel kernel;

    private static ImageIcon img;
    private static int width;
    private static int height;

        public static void main(String[] args) throws Exception {

                CL.setExceptionsEnabled(true);

                cl_platform_id platform[] = new cl_platform_id[1];
                clGetPlatformIDs(1,  platform, null);
                cl_device_id device[] = new cl_device_id[1];
                int num_devices[] = new int[1];
                clGetDeviceIDs(platform[0], CL_DEVICE_TYPE_GPU, 1, device, num_devices);
                cl_context_properties props = new cl_context_properties();
                props.addProperty(CL_CONTEXT_PLATFORM, platform[0]);
                context = clCreateContext(props, 1, device, null, null, null);
                queue = clCreateCommandQueue(context, device[0], 0, null);

                StringBuffer sb = new StringBuffer();
                URL resource = ImageTest.class.getResource(KERNEL_PATH);
                String path = Paths.get(resource.toURI()).toFile().getAbsolutePath();
                System.out.println(path);
                Scanner sc = new Scanner(new File(path));
                while(sc.hasNext()) {
                        sb.append(sc.nextLine() + "\n");
                }
        program = clCreateProgramWithSource(context,
                1, new String[]{sb.toString()}, null, null);

        clBuildProgram(program, 0, null, null, null, null);

        kernel = clCreateKernel(program, KERNEL_FUNC, null);

                URL imgResource = ImageTest.class.getResource(IMAGE_PATH);
                String imgPath = Paths.get(imgResource.toURI()).toFile().getAbsolutePath();
                System.out.println(imgPath);
        img = new ImageIcon(imgResource);

        width = img.getIconWidth();
        height = img.getIconHeight();

        BufferedImage bimg = new BufferedImage(width, height, BufferedImage.TYPE_INT_RGB);
        Graphics g = bimg.createGraphics();
        img.paintIcon(null, g, 0,0);
        g.dispose();
        DataBufferInt dbb = (DataBufferInt)bimg.getRaster().getDataBuffer();
            int[] intArray = dbb.getData();
            System.out.println(intArray.length);


        cl_image_format format = new cl_image_format();
        format.image_channel_data_type = CL_UNORM_INT8;
        format.image_channel_order = CL_RGBA;
        cl_image_desc desc = new cl_image_desc();
        desc.image_height = img.getIconHeight();
        desc.image_width = img.getIconWidth();
        desc.num_mip_levels = 0;
        desc.num_samples = 0;
        desc.image_type = CL_MEM_OBJECT_IMAGE2D;

        cl_mem img_mem = clCreateImage(context,
                        CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
                        format,
                        desc,
                        Pointer.to(dbb.getData()),
                        null);

        cl_mem dst_img_mem = clCreateImage(context,
                        CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR,
                        format,
                        desc,
                        null,
                        null);

        clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(img_mem));
        clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(dst_img_mem));

        long[] global_work_size = new long[] {300,300,0};
        long[] local_work_size = new long[] {1,1,0};
        clEnqueueNDRangeKernel(queue, kernel, 2, null, global_work_size, local_work_size, 0, null, null);

        int[] dstArray = new int[90000];

        clEnqueueReadImage(queue, dst_img_mem, CL_TRUE, new long[]{0,0,0}, new long[]{300,300,1},0, 0, Pointer.to(dstArray), 0, null, null);

                BufferedImage image = new BufferedImage(width, height, BufferedImage.TYPE_INT_RGB);
                image.setRGB(0, 0, width, height, dstArray, 0, width);

                System.out.println(width+"-"+height);

                resource = ImageTest.class.getResource("./");
                path = Paths.get(resource.toURI()).toFile().getAbsolutePath();
                System.out.println(path+"/image.png");
        ImageIO.write(image, "png", new File(path+"/image.png"));

        clReleaseMemObject(img_mem);
        clReleaseMemObject(dst_img_mem);
                clReleaseDevice(device[0]);
                clReleaseContext(context);
                clReleaseCommandQueue(queue);
                clReleaseProgram(program);

        }

}

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
             CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; //(1)

__kernel void helloworld(read_only image2d_t src_image, //(2)
                        write_only image2d_t dst_image) { //(3)

        int2 coord = (int2)(get_global_id(0), get_global_id(1)); //(4)
        float4 pixel = read_imagef( //(5)
                                        src_image, //(6)
                                        sampler, //(7)
                                        (int2)(coord.x,coord.y)); //(8)

        write_imagef(dst_image, coord, pixel); //(9)
}

host_ptr 引数に指定するイメージデータに割り当てるメモリ領域は以下の表のようにイメージ型によって異なります。

イメージの形式を記述する構造体は次の通り定義されています。

 cl_image_format

     typedef struct _cl_image_format {
          cl_channel_order    image_channel_order;
          cl_channel_type     image_channel_data_type;
     } cl_image_format;

image_channel_orderは、チャネル数とチャネルの配置、つまりイメージ内に保存されるチャネルのメモリ配置を指定します。

image_channel_data_type は、チャネルのデータ型を指定します。 image_channel_data_type と  image_channel_order で定められる要素ごとのビット数は、2のべきでなければなりません。

image_channel_data_type の値では、 CL_UNORM_SHORT_565 、 CL_UNORM_SHORT_555 、 CL_UNORM_INT_101010 は特別なイメージ形式となります。

各チャネル要素は unsigned short 型または unsigned int 型に詰められています。データの詰め方としては、最初のチャネルがビットフィールド内の最上位ビットになるように詰められます。次のチャネルは占有するビット領域が連続になるようにすぐ下のビット領域に詰められます。

CL_UNORM_SHORT_565 の場合、Rが15から11ビット、Gが10から5ビット、Bが4から0ビットに詰められます。

CL_UNORM_SHORT_555 の場合、15ビット目は未定義となり、Rが14から10ビット、Gが9から5ビット、Bが4から0ビットに詰められます。

CL_UNORM_INT_101010 の場合、31から30ビットは未定義となり、R が29から20ビット、G が19から10ビット、B が9から0ビットに詰められます。

/* 入力イメージ形式を設定します。*/
png_format.image_channel_order = CL_RGBA;
png_format.image_channel_data_type = CL_UNORM_INT8;

OpenCLでイメージを処理する場合、イメージ記述子を設定しなくてはいけません。記述子で最低限設定が必要なものは、 cl_mem_object_type と、高さ、幅、奥行きに関する情報です。

cl_mem_object_type については以下のものを設定できます。

幅、高さ、奥行きについてはピクセル数などを指定できます。本書では 300 x 300 画素の画像を使うことが多いです。

イメージ記述子はC構造体のため、以下のように設定することができます。

/*イメージ属性を設定します。*/
cl_image_desc image_desc;
image_desc.image_type = CL_MEM_OBJECT_IMAGE2D;
image_desc.image_width = width;
image_desc.image_height = height;
image_desc.num_mip_levels = 0;
image_desc.num_samples = 0;
image_desc.buffer = NULL;

clGetSupportedImageFormats関数はコンテキスト内のデバイスがサポートするイメージ型式の和集合を戻します。イメージメモリオブジェクトに関する以下の情報を指定したときにOpenCL実装がサポートするイメージ形式の一覧取得に使うことができます。

関数の定義は以下のようになります。

int org.jocl.CL.clGetSupportedImageFormats(
    cl_context context, //(1)
    long flags, //(2)
    int image_type, //(3)
    int num_entries, //(4)
    cl_image_format[] image_formats, //(5)
    int[] num_image_formats) //(6)