koturnの日記

転職したい社会人2年生の技術系日記.ブログ上のコードはコピペ自由です.

不思議なHello World

背景

最近になって,x64をちゃんと勉強したので,少しだけひねったHello WorldC言語で書いてみることにした. (※x64のLinux環境でしか動作しない)

const char main[] = "\x48\xc7\xc0\x01\x00\x00\x00\xba\x0d\x00\x00\x00\xbf\x01\x00\x00\x00\xe8\x0c\x00\x00\x00\x48\x81\xc6\x11\x00\x00\x00\x0f\x05\x31\xc0\xc3\x48\x8b\x34\x24\xc3Hello World!\n";

何とワンライナーである. ヘッダのインクルードが無く,main関数ではなく,main配列となっている点が通常のHello Worldと比較すると異質な点であろう.

ワンライナーではわかりづらいので,コメントを入れてみる. 要は機械語である.

const char main[] =
  // mov $0x01,%rax
  "\x48\xc7\xc0\x01\x00\x00\x00"
  // mov $0x0d,%edx
  "\xba\x0d\x00\x00\x00"
  // mov $0x01,%edi
  "\xbf\x01\x00\x00\x00"
  // callq 0x0c
  "\xe8\x0c\x00\x00\x00"
  // add $0x11,%rsi
  "\x48\x81\xc6\x11\x00\x00\x00"
  // syscall
  "\x0f\x05"
  // xor %eax,%eax
  "\x31\xc0"
  // ret
  "\xc3"
  // mov (%rsp),%rsi
  "\x48\x8b\x34\x24"
  // ret
  "\xc3"
  // String data
  "Hello World!\n";

先にタネ明かしをすると,以下を行っているだけにすぎない.

// 第1引数:ファイルディスクリプタ (1: stdout)
// 第2引数:文字列領域へのポインタ
// 第3引数:出力文字数
write(1, "Hello World!\n", 13);

このことを念頭に置き,解説を行っていく.

解説

全体像

mainというのは所詮 crt*.o 中から呼び出しされているに過ぎない. 従って,main という名前の指し示すアドレスの先に関数と同様のデータがあれば問題はない.

基本的に1つの配列に納めたかったので,コードとデータを同一の領域に置いてある. 別々にした場合,データのアドレスを取得するのが面倒になるためだ. (引数でデータのポインタを渡す等の処理が必要になるが,今回はmainを作っているため難しい)

配列 main 中の "Hello World!\n" より前は実行するコードであり,"Hello World!\n" は当然データである. 後は,コード中のいずれかの位置で,その地点のアドレスを取得し, "Hello World!\n" 領域のアドレスを取得すればよい.

システムコール

x64におけるシステムコールでは各レジスタが以下のように役割を果たす. 今回利用するのは第三引数まで.

レジスタ 役割
rax システムコール番号
rdi システムコールの第1引数
rsi システムコールの第2引数
rdx システムコールの第3引数

今回利用する write システムコールC言語APIは以下の通り.

size_t write(int fd, const void *buf, size_t count)

引数 役割
fd ファイルディスクリプタ
buf 出力データのサイズ
count 出力文字数

以上より,レジスタを以下の状態にすることを目指す.

レジスタ 中身
rax 1 (writeシステムコールの番号は1)
rdi 1 (stdoutは1)
rsi Hello Worldへのアドレス
rdx 13 (“Hello World!\n” の文字数は13)

現在位置の取得

Hello World! 文字列アドレスの取得のために,実行しているコードの位置を取得する必要がある. 現在位置はレジスタripに入っているが,直接値を取り出すことはできない. したがって,

  1. call命令を利用し,rspレジスタの管理アドレスにripの値を書き込ませる.
  2. rspが書き込んだ領域からripを取得(mov (%rsp),%rsi

という手段で取得する.

コードまとめ

ここまでの話をまとめると,配列の中身は以下の通りになる.

命令 動作
mov $0x01,%rax raxに1を格納
mov $0x0d,%edx rdxの下位32bitに0x0d(0x14)をセット
mov $0x01,%edi rdiの下位32bitに0x01をセット
callq 0x0c 12bbyte先を関数として呼び出し
add $0x11,%rsi 0x0fをrsiに加算.この命令位置から17byte先が "Hello World!\n"
syscall システムコール
xor %eax,%eax mainの返り値を0に設定(eaxは関数の返り値)
ret main関数のreturn
(%rsp),%rsi 呼出元のアドレスをrsiに格納
ret 呼出位置アドレス取得関数のreturn(add $0x0f,%rsiに復帰)
"Hwllo World!\n" 文字列データ

その他いろいろ

通常,関数の返り値はraxに格納するものだが,面倒なので今回は直接rsiに放り込むという手段を取っている. 自分でコードを書く分には,x64の呼出規約を無視しても構わないだろう.

まとめ

この記事では,少し奇妙な形のC言語Hello Worldを紹介した. 一見,難解に見えるHello Worldのコードであるが,わかってしまえば大したことはない.

今回はx64用のコードを紹介した. なので,Wandboxに貼り付ければ実行できる.

x86については読者の課題としよう(笑)

Javaで簡潔にディープコピーを行う

はじめに

発端は以下の一連のツイート.

Javaのオブジェクトのコピーといえば clone() メソッドを用いるものである. しかし,標準ライブラリのArrayListでは以下のどちらの手段を用いても,シャローコピーとなるようだ.

// list01 の型は ArrayList<Hoge>
ArrayList<Hoge> list02 = (ArrayList<Hoge>) list01.clone();
ArrayList<Hoge> list03 = new ArrayList<>(list01);

これを解消するためには,一々要素のコピーを行う必要があるだろう.

ArrayList<Hoge> list03 = new ArrayList<>(Arrays.asList(list01.stream()
        .map(elm -> {
            Hoge obj = null;
            try {
                obj = (Hoge) elm.clone();
            } catch (CloneNotSupportedException e) {
                throw new RuntimeException(e);
            }
            return obj;
        }).toArray(Hoge[]::new)));

しかし,こういうのはちょっと面倒なように思える上, ArrayList にしか適用できない. そこで,任意のオブジェクトに対して,ディープコピーを行うことを考える

ディープコピー

C# だと,オブジェクトのシリアライズ,デシリアライズを利用してコピーを行うのが一般的となっている. それをJavaで真似てやってみる.

以下のメソッドを定義する.

@SuppressWarnings("unchecked")
public static <T> T deepcopy(T obj) throws IOException, ClassNotFoundException {
    ByteArrayOutputStream baos = new ByteArrayOutputStream();
    new ObjectOutputStream(baos).writeObject(obj);
    return (T) new ObjectInputStream(new ByteArrayInputStream(baos.toByteArray())).readObject();
}

このメソッドは以下のように利用できる. 返り値の型はジェネリクス型推論により,引数と同一の型になるため,呼び出し側でキャストを行う必要はない.

// list01 の型は ArrayList<Hoge>
ArrayList<Hoge> list02 = deepcopy(list01);

このように簡潔にディープコピーが可能である. ただし,シリアライズ,デシリアライズを行うためには,対象のオブジェクトのクラスが Serializable インタフェースを実装しておく必要(メンバも含め)がある点には気をつけておくこと.

また,シリアライズ,デシリアライズのコストはやや高めなので,自作クラスのディープコピーは clone() メソッドを適切に実装するのがよいだろう.

サンプルコード

import java.io.ByteArrayInputStream;
import java.io.ByteArrayOutputStream;
import java.io.IOException;
import java.io.ObjectInputStream;
import java.io.ObjectOutputStream;
import java.io.Serializable;
import java.util.ArrayList;
import java.util.Arrays;

public class Main {
    public static void main(String[] args) throws ClassNotFoundException, IOException{
        ArrayList<Hoge> list01 = new ArrayList<>();
        list01.add(new Hoge(10));
        list01.add(new Hoge(20));
        @SuppressWarnings("unchecked")
        ArrayList<Hoge> list02 = (ArrayList<Hoge>) list01.clone();
        ArrayList<Hoge> list03 = new ArrayList<>(list01);
        ArrayList<Hoge> list04 = new ArrayList<Hoge>(Arrays.asList(list01.stream()
                .map(elm -> {
                    Hoge obj = null;
                    try {
                        obj = (Hoge) elm.clone();
                    } catch (CloneNotSupportedException e) {
                        throw new RuntimeException(e);
                    }
                    return obj;
                }).toArray(Hoge[]::new)));
        ArrayList<Hoge> list05 = deepcopy(list01);

        list01.get(0).setValue(100);

        System.out.println("list01: Original");
        list01.stream().mapToInt(t -> t.getValue()).forEach(System.out::println);
        System.out.println("list02: Shallow copy");
        list02.stream().mapToInt(t -> t.getValue()).forEach(System.out::println);
        System.out.println("list03: Shallow copy");
        list03.stream().mapToInt(t -> t.getValue()).forEach(System.out::println);
        System.out.println("list04: Deep copy");
        list04.stream().mapToInt(t -> t.getValue()).forEach(System.out::println);
        System.out.println("list05: Deep copy");
        list05.stream().mapToInt(t -> t.getValue()).forEach(System.out::println);
    }

    @SuppressWarnings("unchecked")
    public static <T> T deepcopy(T obj) throws IOException, ClassNotFoundException {
        ByteArrayOutputStream baos = new ByteArrayOutputStream();
        new ObjectOutputStream(baos).writeObject(obj);
        return (T) new ObjectInputStream(new ByteArrayInputStream(baos.toByteArray())).readObject();
    }
}

/**
 * Cloneable は 要素を1つずつ追加するdeep copyの例のため
 * Serializable は シリアライズ / デシリアライズによる deep copyのため
 */
class Hoge implements Cloneable, Serializable {
    private static final long serialVersionUID = 1145141919810L;
    private int value;

    Hoge(int value) {
        this.value = value;
    }

    int getValue() {
        return value;
    }

    void setValue(int value) {
        this.value = value;
    }

    @Override
    public Hoge clone() throws CloneNotSupportedException {
        Hoge obj = (Hoge) super.clone();
        obj.value = value;
        return obj;
    }
}

まとめ

Javaの標準ライブラリの中にはディープコピーが楽ではないものもある.

リアリアズ,デシリアライズ用いてdeep copyを行うためには以下のようなメソッドを用意するとよい.

@SuppressWarnings("unchecked")
public static <T> T deepcopy(T obj) throws IOException, ClassNotFoundException {
    ByteArrayOutputStream baos = new ByteArrayOutputStream();
    new ObjectOutputStream(baos).writeObject(obj);
    return (T) new ObjectInputStream(new ByteArrayInputStream(baos.toByteArray())).readObject();
}

Vimのあまり使われない機能について~暗号化と印刷~

はじめに

この記事はVim Advent Calendar 2016の15日目の記事です.

今回はVimであまり知られていない,使われていない機能であると思われる暗号化機能について書きたいと思います.

暗号化機能

Vimテキストエディタですが,単なるテキストエディットだけに留まらず,テキストファイルを暗号化する機能も備わっています. 編集したテキストファイルを保存するとき,ユーザが指定したキーによって暗号化し,再度そのファイルオープンするときにキーの入力を求めます.

f:id:koturn:20161214012205g:plain

そして,オープン時に入力されたキーによって復号化を行うので,正しいキーであれば元のテキストが表示され,異なったキーであればデタラメなテキストが表示されるようになっています.

利用可能な暗号化方式

VIm 8.0の時点では,以下の表に示す3つの暗号化方式が利用可能となっています. 各暗号化方式についての詳細は :h cryptmethod を参照しましょう.

暗号化方式 強度 バージョン
zip 7.2 以前
blowfish 中強度(脆弱性あり) 7.3 以降
blowfish2 中強度 7.4.401 以降

cryptmethod オプションに上記のいずれかの値をセットすることで,暗号化方式を指定することができます.

デフォルトでは zip の指定となっているのですが,最新のVimを利用しており,古いバージョンのVimで開き直すことを考えないのであれば, blowfish2 を指定するべきでしょう. 利用できる最新の暗号化方式を指定していない場合,暗号化するときにVimからワーニングメッセージが出て邪魔になります.

あらゆる環境を想定していて,その環境のVimで利用可能な最新の暗号化方式を利用したいという奇特な方は,以下の設定を.vimrcに記述するとよいでしょう.

if has('cryptv')
  if v:version > 704 || v:version == 704 && has('patch401')
    set cryptmethod=blowfish2
  elseif v:version >= 703
    set cryptmethod=blowfish
  else
    set cryptmethod=zip
  endif
endif

blowfish2blowfish が利用可能かどうかはそれぞれ has('crypt-blowfish2')has('crypt-blowfish') で確認することもできるのですが,これらの条件が利用できるようになったのは Vim7.4.1009 以降でおり,blowfish2 が利用可能になったバージョンよりも後です. したがって,Vimのバージョンによって分岐する方がより古いVimに対応しているといえるでしょう.

暗号化を行う

さて,実際に暗号化を行ってみましょう.

コマンド経由

現在のファイルを暗号化したい場合,ファイル保存前に :X コマンドを実行し,暗号化キーを指定してください. よくあるパスワード入力のように,確認のため2回同じキーを入力する必要があります.

f:id:koturn:20161214012209g:plain

オプション経由

オプション key の値に暗号化に用いるキーを設定し,その後にファイルを保存することでも暗号化できます. ただし, .vimrc に記述すると,暗号化キーが筒抜けなので,通常は記述するべきではないでしょう. (.vimrcに記述した場合は,何もしなくても常に暗号化ができるというメリットはあります) 利用する場合は,保存前に :set key=hogehoge といった感じで,手動でオプションに値を設定しましょう.

オプション key は扱いがやや特殊であり, set keyecho &key としても,設定値を確認することはできません. また,ヒストリにも set key= までしか残らず,暗号化キーが分からないように細工されています.

key オプションに値を設定する場合は,暗号化キーの入力が一度で済みますが,その分typoに気付かなかった場合,ファイルの復号化が難しくなるでしょう. 僕としては, :X コマンドを利用して暗号化キーを設定するとよいと思います.

暗号化を解除する

一度暗号化したファイルは,再度オープンし,保存し直しても,暗号化されたままです. 暗号化を解除したくなった場合は,空のキーを設定し直すとよいです.

:X コマンドを用いる場合は,キーを何も入力せずに <CR> だけを2回,key オプションを利用する場合は, set key= としましょう.

印刷機能

実はVimには印刷機能が備わっています. 印刷は簡単で, :hardcopy コマンドを実行するだけです.

印刷設定

印刷設定に関するオプションは7種類あり,以下の通りです.

オプション 機能
printdevice :hardcopy! のようにbang付きで実行されたときのプリンタを指定
printencoding 印刷に用いるエンコーディングを設定
pritfont Ascii文字の印刷に用いるフォントを設定
printheader 印刷のヘッダを設定. statusline と同じ書式で指定する
printmbcharset マルチバイトの文字セットを指定
printmbfont マルチバイトの印刷に用いるフォントを設定
printoptions 余白,行番号の有無,折り返しの有無等を設定

上記7つの内,基本的に設定しておくべきなのは以下の3つでしょう.

  1. printfont
  2. printmbfont
  3. printoptions

あとは,自動設定に任せる or 用いない等で設定しなくても大丈夫でしょう.

僕はWindowsユーザなのですが,以下のように設定しています.

set printoptions=number:y,header:0,syntax:y,left:5pt,right:5pt,top:10pt,bottom:10pt
set printfont=Consolas:h8
set printmbfont=r:MS_Gothic:h8,a:yes

標準では行番号の印刷がoffになっているので, number:y を入れておくとよいでしょう. また,僕はヘッダ行を必要としないため header:0 を付与しています.

上記の設定にはありませんが,縦向きor横向き,印刷する用紙の大きさ等を設定することもできます.

印刷設定

印刷には,設定中のカラースキームが一部反映されます. そのため,実際に印刷してみると,悲惨な結果になっていることがあります.

f:id:koturn:20161214012216p:plain

f:id:koturn:20161214012223p:plain

仮想プリンタ等を導入し,実際に出力してみて,どうなっているかを確認してから,紙に出力するとよいでしょう.

まとめ

この記事ではVimの暗号化機能と印刷機能について紹介しました. 暗号化機能の利用シーンとしては,ギョームでちょっと秘密にしておくべきファイルを扱うとき,などが考えられるでしょうか. 印刷機能は何かの用事があってコードを印刷したいときに利用できると思います.

なお,neovimでは暗号化機能は削除されているため,利用することはできません. この点には気を付けましょう.

denite.nvimのsourceを作ってみる

はじめに

unite.vim に替わる新しいプラグインとして,denite.nvimが開発されている. まだまだ開発途上ではあるものの,じわじわとunite.vimから移行する人が見受けられる. denite.nvimはneovimだけでなく,Vim 8.0以降であれば利用可能であるのも魅力の1つであるだろう.

そこで,(ゴミ)プラグイン作者の一人として,denite-sourceの作成方法について記したいと思う.

source / kind作成

サンプルとして作るもの

何よりもまず,実例を示すのが早いだろう. サンプルとして,以下のようなsourceを作成してみることにしよう.

  • sourceの引数で指定されたディレクトリ(複数指定可)のファイル一覧を表示
  • sourceの引数が指定されなかった場合はホームディレクトリのファイル一覧を候補として表示
  • 選択したファイルの行数を print() する

ファイル構成

以下の構成でプラグインディレクトリを作成する. 今回は denite-source のみのプラグインであるため, plugin/autoload/ も存在しない. 同じものは GitHub: koturn/vim-denite-sampleにも置いてある.

vim-denite-sample/
|
+-.gitignore
|
+-README.md
|
+-LICENSE
|
+-rplugin/
  |
  +-python3/
    |
    +-denite/
      |
      +-kind/
      | |
      | +-sample.py
      |
      +-source/
        |
        +-sample.py

コード

vim-denite-sample/rplugin/pyhton3/denite/source/sample.py

本体同梱のsourceを真似て,sourceを作成する. まず, from .base import Base のようにBaseクラスをインポートし,それを継承して, Source クラスを作成する. Source クラスには必要なメソッドを実装するとよい. 特に, gather_candidates メソッドは候補取得のためのメソッドなので,必ず実装することになる.

# -*- coding: utf-8 -*-
# FILE: sample.py
# AUTHOR: koturn <jeak.koutan.apple@gmail.com>
# License: MIT License

from .base import Base
import glob
import itertools
import os


class Source(Base):
    def __init__(self, vim):
        #  このvimという引数はVimとPythonで相互にやりとりするためのインタフェース
        # :help pyth を参照すれば,おおよそのことが記述してある
        super().__init__(vim)
        # Denite xxx の xxx に相当する部分
        self.name = 'sample'
        # kind名を指定
        self.kind = 'sample'

    def on_init(self, context):
        '''
        このsourceが指定されて起動されたときに呼び出される.
        元のバッファのファイルタイプの取得などに利用する.
        今回は必要ないので,とりあえずechomsgしておく
        '''
        # print() は :echomsg と同等
        print('on_init')
        # filetypeなどの値は,以下のように context のキーとして生やす
        # selfのメンバにするのはよくない
        # context['__filetype'] = self.vim.eval('&filetype')

    def on_close(self, context):
        '''
        Denite終了時に呼び出される.
        '''
        # コンストラクタ以外では,self.vim を参照して,Vimインタフェースを利用する
        print('on_close')

    def gather_candidates(self, context):
        '''
        候補の取得を行う関数
        '''
        # sourceの引数context['args']はユーザがDenite実行時に以下のように : 区切りで指定する
        #   Denite sample:arg1:arg2:arg3
        dirpaths = ['~'] if len(context['args']) == 0 else context['args']
        print(dirpaths)
        candidates = filter(lambda path: os.path.isfile(path),
                        itertools.chain.from_iterable(
                            map(lambda dirpath: glob.glob(os.path.expanduser(dirpath) + '/*'), dirpaths)))
        # 辞書を返す
        #   word: 候補欄に表示される文字列
        #   アクション側で利用するための情報を増やしたい場合は 'action__xxx'
        #   という名前のキーを利用すること
        return list(map(
            lambda candidate: {
                'word': candidate},
            candidates))

vim-denite-sample/rplugin/pyhton3/denite/kind/sample.py

denite-kindはアクション部分を記述するものである. denite-kindに関しても,sourceと同様に from .base import Base として,Base クラスをインポートする. そして,それを継承した Kind クラスを作成し,必要なメンバ,メソッドを実装する. メンバの中で特に重要なのが default_action であり,これは <CR> を押下したときのアクションに相当する.

# -*- coding: utf-8 -*-
# FILE: sample.py
# AUTHOR: koturn <jeak.koutan.apple@gmail.com>
# License: MIT License

from .base import Base

class Kind(Base):
    def __init__(self, vim):
        super().__init__(vim)
        # kindの名前.この名前をsourceで指定することにより,関連付けられる
        self.name = 'sample'
        # デフォルトのアクション(候補上で<CR>を押下したときに実行するアクション)を指定する.
        # self.default_action = 'xxx' と指定すると,このクラスのメソッド:
        # action_xxx() が呼び出される
        self.default_action = 'sample'

    def action_sample(self, context):
        '''
        コンストラクタで指定したように,デフォルトのアクションを記述する関数
        '''
        for target in context['targets']:
            filepath = target['word']
            print(filepath + ' has ' + str(sum(1 for line in open(filepath, encoding='utf-8'))) + ' lines')

    def action_preview(self, context):
        '''
        候補上にカーソルが移動する度に呼び出される関数
        :Denite -auto-preview ... のように -auto-preview オプションを付加すると,previewモードになる
        付加しない場合でも,以下のキーを押下することでpreview動作を行うことは可能
          insertモード: <C-v>
          normalモード: p
        return Trueとしないと,deniteが終了する
        '''
        # context['targets'] に source側のgether_candidatesの返り値が格納されている
        filepath = context['targets'][0]['word']
        print(filepath + ' has ' + str(sum(1 for line in open(filepath, encoding='utf-8'))) + ' lines')
        return True

サンプルの使い方

以下のように起動すると,ホームディレクトリのファイル一覧が表示される. 選択したファイルの行数が print() に渡されるので, :message 等で確認するとよい.

:Denite sample

次のように,sourceの引数とディレクトリを渡すと,ホームディレクトリではなく,そのディレクトリのファイル一覧が表示される.

:Denite sample:~/Desktop

引数を複数指定することも可能だ.

:Denite sample:~/Desktop:~/Documents/

テンプレート

最低限のdenite-sourceを実装するのであれば,以下のようなテンプレートを用意しておくと楽になるだろう. on_initon_close については,必ずしも利用するわけではないので,コメントアウトしてある.

denite-source のテンプレート

# FILE: xxx.py
# AUTHOR: xxx <xxx.yyy.zzz@gmail.com>
# License: MIT license

from .base import Base


class Source(Base):
    def __init__(self, vim):
        super().__init__(vim)
        self.name = 'xxx'
        self.kind = 'xxx'

    # def on_init(self, context):
        # TODO

    # def on_close(self, context):
        # TODO

    def gather_candidates(self, context):
        # TODO: Following code is a sample
        candidates = ['apple', 'banana', 'cake']
        return list(map(
            lambda candidate: {
                'word': candidate},
            candidates))

denite-kind のテンプレート

# FILE: sample.py
# AUTHOR: xxx <xxx.yyy.zzz@gmail.com>
# License: MIT license

from .base import Base


class Kind(Base):
    def __init__(self, vim):
        super().__init__(vim)
        self.name = 'xxx'
        self.default_action = 'xxx'

    def action_sample(self, context):
        # TODO

    # def action_preview(self, context):
        # TODO

余談

実は,unite.vimにもdeite-sourceが実装され,unite-sourceをdeniteから利用することが可能になった. 以下のようにして,deniteからuniteを利用する.

" xxx がunite-source名
:Denite unite:xxx

このため,既にプラグインにunite-sourceを実装している場合は,改めて同等のdenite-sourceを実装し直す必要はないということだ. しかし,ユーザとしてはdenite-sourceも提供されていた方が気持ちが良いと思われる.

ちなみに,denite-sourceはPythonであるため,当然バイトコードが生成される. そのため, .gitignore に,

*.py[cod]

と追記しておくのがよいだろう.

おまけ

実はちょっとした自作プラグインに取り込んである.

:Denite mplayer

とすることで,デフォルトでは ~/ 以下のファイルを再帰的に表示し,mplayer を用いて,選択したファイルを再生するようになっている.

デフォルトディレクトリは g:mplayer#default_dir を設定することで変更可能であるし,

:Denite mplayer:~/Music/

のように指定してもよい.

また,プレビュー機能にも対応しているため,

:Denite mplayer:~/Music/ -auto-preview

とすることで,Deniteバッファ上で,カーソル移動の度にカーソル下のファイルを再生できる.

このように,ちょっとした自作プラグインに導入することで,格段に便利になる.

まとめ

この記事では,最も単純なdenite-sourceの作成方法を示した. プラグイン作成初心者のdenite-source作成の足掛かりとなれば幸いである.

ただし,あくまで,最も初歩的な作成方法について述べただけであり,触れていない項目の一例として matcherやsorter,converterといった話もある. それらの機能についてはdocやdenite.nvim自体のソースコードを参照して欲しい.

なお,deniteは開発が活発なプラグインであるため,数ヶ月後にはこの記事の内容が古いものとなっている可能性は十分にある.

Osaka.vim #8 に行ってきた

はじめに

昨日,10/29(土)にOsaka.vim #8に行ってきました. 前回は大学の後輩を誘って参加しましたが,今回は都合が悪かったので参加できず...残念.

やったこと

denite.nvimの勉強

最初はShougo/denite.nvimについて勉強するかーと思って,denite.nvimのhelpと本体付属のsourceを読んでました. 少し前に自作プラグイン用のdenite.nvimのsourceを書いたときにちょっと勉強してたのもあって,これ自体はそこそこに終えました. で,以前書いたsourceの作法としてマズかった部分を修正しました.

プラグインの修正

ここ最近手を加えている自作プラグインkoturn/vim-mplayerをいじっていました. このプラグインVimからmplayerを起動,操作するプラグインです. 2年前に作ったプラグインであり,Osaka.vim #2のときに機能作り込みしてた記憶があります. 2年前に基本機能を作って寝かせ,1年前にunite sourceやctrlpの拡張,fzfと連携できるようにして寝かせ,今年はjob対応,みたいな感じで,なんか1年周期でいじっているプラグインですね.

いじったと言っても,機能追加は特になく,内部をスッキリさせたぐらいですね. あらかたの修正はここ数週間で既に終えていたというのもありますね.

Osaka.vim #8までの「ここ数週間での修正」というのは,

って感じです.

以前は,:MPlayer hoge.mp3 みたいなコマンド経由で操作するだけしかできないようにしていましたが,オブジェクト指向チックにするに伴い,複数のmplayerプロセスを操作することができるようになりました.

以下のような感じで,他のプラグインから利用可能になりました. (mplayer利用するゴミプラグインはいくつか作ったので,この変更は自分のためですね) 単純なmplayer操作プラグインにライブラリとしての側面を追加した形,って感じですかね. 以下の形で利用できます.

" mplayerのインスタンス生成
let s:mplayer = mplayer#new()
" mplayerプロセス起動と指定ファイル再生
call s:mplayer.play('~/Music/hoge.mp3')
" 音量レベルを70%に
call s:mplayer.set_volume(70)
" 終了(プロセス終了)
" call s:mplayer.stop()

help書き

koturn/vim-mplayerのhelpを面倒で書いてなかったので,そろそろ書くか~って思って書いていました. helpをゼロから書くのは面倒なので,LeafCage/vimhelpgeneratorを利用して,雛形を生成し,ごにょごにょと書いてました. 上記のhelp雛形生成プラグインは,今回のOsaka.vim #8 参加者の一人である@LeafCageさんのベンリプラグインですね.

英語力無いマンであり,それなりに書くことがあったので,helpはまだ書き終えていません()

出てきた話とか

とりあえず,箇条書きで.

前回と同様,電源タップを2個持っていったのですが,みんな持ってきてて,むしろ必要なかったぐらいでした. fzf関連の話で自分のブログが検索の上の方に出てきたのは面白かったです. 他にも話があった気がするけど,思い出しながら書くと抜けがあるので,あとで追記する可能性あります.

あと書き

自分のもくもく作業としては,そこそこ進捗はあったと思います. 後半に@kozo2さんが「LTというか質問という形で...」と前に出て,対話するスタイルで話を始めたのですが,バンバンと各個人の知識が出てくるのはよかったな~と思いました. また次回も参加したいですね.

Windowsでneovimを使う(改訂版1)

はじめに

前回の記事(約1年前1)Windowsでneovimを使う方法について書いた. しかし,約1年も経過すれば,色々と変化はあり,前回の記事内容そのままというわけにはいかない.

Windows neovim」でググると,僕の記事が一番上に出てくるというのもあり,新しい記事を書くべきだと思い立ったわけだ. (プラグイン制作にあたり,Vimのjobとneovimのjobについて調査していたのもきっかけ)

ダウンロード

公式Wiki 公式でneovim本体のバイナリとQtで実装されたフロントエンドのバイナリが配布されているので,それらをダウンロードしてくる.

本体

公式のneovim本体のバイナリは,以下のAppVeyorの成果物へのリンクから取得できる(Wikiのリンクそのまま).

フロントエンド

公式Wikiにある通り,equalsraf/neovim-qt にてバイナリが配布されているので取ってくる.

ファイル配置

システムのランタイム

上記ダウンロード手順により落とした Neovim.zip と neovim-qt.zip を展開し,どちらの中にもある share/ を以下のディレクトリ下に配置する(無ければ作成). つまり,併合して配置する.

  • C:/Program Files (x86)/nvim/

すなわち,nvimのデフォルトのシステムファイルは,

  • C:/Program Files (x86)/nvim/share/nvim/runtime/

である. このディレクトリに置きたく無い場合は,環境変数 $VIMRUNTIME を書き換えろという話らしい.

最後に,

C:/Program Files (x86)/nvim/share/nvim/runtime/share/nvim-qt/runtime/

C:/Program Files (x86)/nvim/share/nvim/runtime/share/nvim/runtime/

に併合する. この操作により, nvim_gui_shim.vim が起動時に読み込まれるようになり,フォント変更やフルスクリーン切り替えなどのコマンドを利用することが可能となる. 詳細は後述する.

バイナリ

Neovim.zip と neovim-qt.zip を展開した中にある bin/ を併合し,適当な位置に配置するとよい. 公式では C:/Program Files (x86)/nvim/bin としてある.

なお, Neovim.zip と neovim-qt.zip のどちらにも libwinpthread-1.dll が含まれるが,サイズの大きい neovim-qt.zip のものをとりあえず使っておけばよい.

設定

設定ファイル

設定ファイルの位置は XDG directory specification に従う環境変数 %XDG_CONFIG_HOME% が定義されている場合, %XDG_CONFIG_HOME%\init.vim が読み込まれるが,定義されていない場合, %USERPROFILE%\AppData\Local\nvim\init.vim が読み込まれる.

通常の ~/_vimrc~/.vimrc) と ~/vimfiles を利用したい場合,シンボリックリンクを作成するとよい. (%USERPROFILE%Linuxでいう $HOMEWindowsでもHOME環境変数を定義してもよい) 以下の例では, %XDG_CONFIG_HOME% を定義しているものとする.

> mklink /D %XDG_CONFIG_HOME%\nvim %USERPROFILE%\vimfiles\
> mklink %USERPROFILE%\vimfiles\init.vim %USERPROFILE%\_vimrc

なお,シンボリックリンクの作成には管理者権限が必要となるので,気をつけること.

GUI設定

neovim-qt特有の設定は %XDG_CONFIG_HOME%\ginit.vim に書く. 特に必要になると思われるフォント設定の例は以下の通り. これは gvim における set guifont=Consolas:h9 に相当する設定である.

Guifont Consolas:h9

ただし,フォント(特に小さいサイズにした場合に対して)によっては以下のようなワーニングが出ることがある.

Warning: Font "Consolas" reports bad fixed pitch metrics

これを無視して設定したい場合, Guifont コマンドに bang を与えるとよい.

Guifont! Consolas:h9

また,起動時にフルスクリーンにしたい場合は以下の通り.

" 引数が1でフルスクリーン,0でフルスクリーン解除
call GuiWindowFullscreen(1)

起動時のWindowサイズを調整したい場合,以下のように nvim-qt.exe の引数に -qwindowgeometry オプションを指定する.

> nvim-qt.exe -qwindowgeometry 800x600

nyaovim

日本人のVimmerとして有名なrhysdさんがelectronにて作成したneovimのフロントエンドがある.

これを利用してもよい. nyaovimのインストールは以下の通り.

> npm install -g nyaovim

あとはneovim本体のバイナリにパスが通っている状態で以下のコマンドを実行するとよい.

> nyaovim

雑感

依然として,公式配布のバイナリでは,utf-8以外の文字コードを扱うことができない. この記事によると,自前ビルドで頑張れば,utf-8以外の文字コードを扱うことも可能になりそうだ. また,前回の記事で,新規ファイルの作成ができない,と書いたが,その問題は解消されたようだ.

まとめ

公式配布のneovimのバイナリをWIndosで利用する手順を示した. 1年前と比較すると,色々と手順に変化がある.

OpenCLのオフラインコンパイル

はじめに

OpenCLといえば,カーネルのコードに以下の2つのコンパイル方式がある.

オンラインコンパイルは,実行時にOpenCLカーネルコードを文字列として関数に渡し,プログラムオブジェクトを構築する手法である. 反対に,オフラインコンパイルは事前にOpenCLカーネルコードをコンパイルし,コンパイル結果のバイナリを生成しておく. そして,生成したバイナリを実行時に読み込んで,プログラムオブジェクトを構築する手法である. オフラインコンパイルは事前にコンパイルを行う分,実行時のコンパイル時間を削減することができるわけだ.

オフラインコンパイル

オフラインコンパイルには,大別して2つの方針がある.

  1. OpenCLSDK付属のオフラインコンパイラを利用する
  2. OpenCLAPIを用いて,コンパイルプログラムを書く

前者はSDKが必須となるが,後者は不要となる. この記事では,OpenCLAPIを用いて,オフラインコンパイラを書くことにする.

オフラインコンパイルを行うプログラムを作成する

clCreateProgramWithSource()clBuildProgram() により生成した cl_program から clGetProgramInfo() を用いることで,バイナリの情報およびバイナリそのものを取得する形になる. 具体的には以下のようになる. このコードは,0番目のプラットフォームID,0番目のデバイスIDを対象に,コマンドライン引数で指定したカーネルソースコードのオフラインコンパイルを行うものだ.

  • oclc.cpp
// oclc.cpp
// g++ -gnu++11 -O3 oclc.cpp -lOpenCL -o oclc
#include <iostream>
#include <cmath>
#include <cstdlib>
#include <cstring>
#include <ctime>
#include <fstream>
#include <memory>
#include <vector>

#ifdef __APPLE__
#  include <OpenCL/opencl.h>
#else
#  include <CL/cl.h>
#endif


static constexpr cl_uint kNDefaultPlatformEntry = 16;
static constexpr cl_uint kNDefaultDeviceEntry = 16;


/*!
 * @brief プラットフォームIDを取得
 * @param [in] nPlatformEntry  取得するプラットフォームID数の上限
 * @return  プラットフォームIDを格納した std::vector
 */
static inline std::vector<cl_platform_id>
getPlatformIds(cl_uint nPlatformEntry = kNDefaultPlatformEntry)
{
  std::vector<cl_platform_id> platformIds(nPlatformEntry);
  cl_uint nPlatform;
  if (clGetPlatformIDs(nPlatformEntry, platformIds.data(), &nPlatform) != CL_SUCCESS) {
    std::cerr << "clGetPlatformIDs() failed" << std::endl;
    std::exit(EXIT_FAILURE);
  }
  platformIds.resize(nPlatform);
  return platformIds;
}


/*!
 * @brief デバイスIDを取得
 * @param [in] platformId    デバイスIDの取得元のプラットフォームのID
 * @param [in] nDeviceEntry  取得するデバイスID数の上限
 * @param [in] deviceType    取得対象とするデバイス
 * @return デバイスIDを格納した std::vector
 */
static inline std::vector<cl_device_id>
getDeviceIds(const cl_platform_id& platformId, cl_uint nDeviceEntry = kNDefaultDeviceEntry, cl_int deviceType = CL_DEVICE_TYPE_DEFAULT)
{
  std::vector<cl_device_id> deviceIds(nDeviceEntry);
  cl_uint nDevice;
  if (clGetDeviceIDs(platformId, deviceType, nDeviceEntry, deviceIds.data(), &nDevice) != CL_SUCCESS) {
    std::cerr << "clGetDeviceIDs() failed" << std::endl;
    std::exit(EXIT_FAILURE);
  }
  deviceIds.resize(nDevice);
  return deviceIds;
}


/*!
 * @brief 指定されたファイル名のファイルを読み込み,std::stringに格納して返却する
 * @param [in] filename  読み込むファイル名
 * @return  ファイルの内容を格納したstd::string
 */
static inline std::string
readSource(const std::string& filename) noexcept
{
  std::ifstream ifs(filename.c_str());
  if (!ifs.is_open()) {
    std::cerr << "Failed to open " << filename << std::endl;
    std::exit(EXIT_FAILURE);
  }
  return std::string((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());
}


/*!
 * @brief 指定された複数のファイル名のファイルを読み込み,std::stringに格納して返却する
 * @param [in] filenames  読み込み対象のファイル名を格納したstd::vector
 * @return  ファイルの内容を格納したstd:stringを格納したstd::vector
 */
static inline std::vector<std::string>
readSource(const std::vector<std::string>& filenames)
{
  std::vector<std::string> srcs(filenames.size());
  for (decltype(srcs)::size_type i = 0; i < srcs.size(); i++) {
    srcs[i] = readSource(filenames[i]);
  }
  return srcs;
}


/*!
 * @brief 指定されたファイル名から拡張子を除いた文字列を返却する
 * @param [in] filename  拡張子を取り除きたいファイル名
 * @return  拡張子を除いたファイル名
 */
static inline std::string
removeSuffix(const std::string& filename) noexcept
{
  return filename.substr(0, filename.find_last_of("."));
}


/*!
 * @brief このプログラムのエントリポイント
 * @param [in] argc  コマンドライン引数の数
 * @param [in] argv  コマンドライン引数
 * @return 終了ステータス
 */
int
main(int argc, char* argv[])
{
  // OpenCLのコンパイラに渡すオプション文字列
  // 今回はとりあえず空
  static const char kOptStr[] = "";

  std::vector<std::string> args(argc - 1);
  if (args.size() < 1) {
    std::cerr << "Please specify only one or more source file" << std::endl;
    return EXIT_FAILURE;
  }
  for (decltype(args)::size_type i = 0; i < args.size(); i++) {
    args[i] = std::string(argv[i + 1]);
  }

  // プラットフォームを取得
  std::vector<cl_platform_id> platformIds = getPlatformIds(1);

  // デバイスを取得
  std::vector<cl_device_id> deviceIds = getDeviceIds(platformIds[0], 1, CL_DEVICE_TYPE_DEFAULT);

  // コンテキスト生成
  cl_int errCode;
  std::unique_ptr<std::remove_pointer<cl_context>::type, decltype(&clReleaseContext)> context(
      clCreateContext(nullptr, 1, &deviceIds[0], nullptr, nullptr, &errCode), clReleaseContext);
  if (errCode != CL_SUCCESS) {
    std::cerr << "clCreateContext() failed" << std::endl;
    return EXIT_FAILURE;
  }

  // ソースコード読み込み
  std::vector<std::string> kernelSources = readSource(args);
  std::pair<std::vector<const char*>, std::vector<std::string::size_type> > kernelSourcePairs;
  kernelSourcePairs.first.reserve(kernelSources.size());
  kernelSourcePairs.second.reserve(kernelSources.size());
  for (const auto& kernelSource : kernelSources) {
    kernelSourcePairs.first.emplace_back(kernelSource.c_str());
    kernelSourcePairs.second.emplace_back(kernelSource.length());
  }

  // プログラム生成
  // 複数ソースファイルに対応
  std::unique_ptr<std::remove_pointer<cl_program>::type, decltype(&clReleaseProgram)> program(
      clCreateProgramWithSource(
        context.get(),
        static_cast<cl_uint>(kernelSourcePairs.first.size()),
        kernelSourcePairs.first.data(),
        kernelSourcePairs.second.data(),
        &errCode),
      clReleaseProgram);
  if (errCode != CL_SUCCESS) {
    std::cerr << "clCreateProgramWithSource() failed" << std::endl;
    return EXIT_FAILURE;
  }

  // カーネルソースコードのコンパイル
  switch (clBuildProgram(program.get(), 1, &deviceIds[0], kOptStr, nullptr, nullptr)) {
    case CL_SUCCESS:
      break;
    case CL_BUILD_PROGRAM_FAILURE:
      {
        // コンパイルエラーを表示
        std::array<char, 2048> buildLog;
        std::size_t logSize;
        clGetProgramBuildInfo(program.get(), deviceIds[0], CL_PROGRAM_BUILD_LOG, buildLog.size(), buildLog.data(), &logSize);
        std::cerr << "Compile error:\n" << buildLog.data() << std::endl;
      }
      break;
    case CL_INVALID_BUILD_OPTIONS:
      std::cerr << "Invalid option is specified" << std::endl;
      return EXIT_FAILURE;
    default:
      std::cerr << "clBuildProgram() failed" << std::endl;
      return EXIT_FAILURE;
  }

  // デバイス数を取得 (このプログラムでは1が返却されるはず)
  cl_uint nDevice;
  if (clGetProgramInfo(program.get(), CL_PROGRAM_NUM_DEVICES, sizeof(nDevice), &nDevice, nullptr) !=  CL_SUCCESS) {
    std::cerr << "clGetProgramInfo() failed" << std::endl;
  }

  // 各デバイス向けのコンパイル後のバイナリのサイズを取得
  std::unique_ptr<std::size_t[]> binSizes(new std::size_t[nDevice]);
  if (clGetProgramInfo(program.get(), CL_PROGRAM_BINARY_SIZES, sizeof(std::size_t) * nDevice, binSizes.get(), nullptr) != CL_SUCCESS) {
    std::cerr << "clGetProgramInfo() failed" << std::endl;
  }

  // コンパイル後のバイナリをコピー
  std::vector<std::unique_ptr<char> > bins(nDevice);
  for (std::size_t i = 0; i < nDevice; i++) {
    bins[i] = std::unique_ptr<char>(binSizes[i] == 0 ? nullptr : new char[binSizes[i]]);
  }
  if (clGetProgramInfo(program.get(), CL_PROGRAM_BINARIES, sizeof(char*) * nDevice, bins.data(), nullptr) != CL_SUCCESS) {
    std::cerr << "clGetProgramInfo() failed" << std::endl;
    return EXIT_FAILURE;
  }

  // コピーしたバイナリを全てファイルに出力
  std::string basename = removeSuffix(args[0]);
  for (std::size_t i = 0; i < nDevice; i++) {
    if (bins[i] == nullptr) {
      continue;
    }
    std::string filename = basename + ".bin";
    if (nDevice > 1) {
      filename += "." + std::to_string(i);
    }
    std::ofstream ofs(filename, std::ios::binary);
    if (ofs.is_open()) {
      ofs.write(bins[i].get(), binSizes[i]);
    } else {
      std::cerr << "Failed to open: " << filename << std::endl;
    }
  }

  return EXIT_SUCCESS;
}

コンパイルは以下のようにして行う.

$ g++ -gnu++11 -O3 oclc.cpp -lOpenCL -o oclc

次に,以下のようなカーネルのコードを用意する.

  • kernel.cl
// kernel.cl
__kernel void
vecAdd(__global float* z, __global const float* x, __global const float* y, int n)
{
  const int para = 4;
  const int end = (n / para) * para;

  for (int i = 0; i < end; i += para) {
    float4 vtmp = vload4(0, x + i) + vload4(0, y + i);
    vstore4(vtmp, 0, z + i);
  }

  for (int i = end; i < n; i++) {
    z[i] = x[i] + y[i];
  }
}

そして,以下のようにしてカーネルソースコードコンパイルする.

$ ./oclc kernel.cl

これでカレントディレクトリに kernel.bin が生成されていれば,オフラインコンパイルは成功である. コンパイルエラーがある場合は,エラーメッセージを表示するようにしてある.

生成したカーネルバイナリのテストには以下のコードを用いる.

// main.cpp
// g++ -std=gnu++11 -O3 main.cpp -lOpenCL -o main
#include <cstdlib>
#include <fstream>
#include <iostream>
#include <memory>
#include <random>
#include <string>

#ifdef __APPLE__
#  include <OpenCL/opencl.h>
#else
#  include <CL/cl.h>
#endif


static constexpr cl_uint kNDefaultPlatformEntry = 16;
static constexpr cl_uint kNDefaultDeviceEntry = 16;




/*!
 * @brief アラインメントされたメモリを動的確保する関数
 * @param [in] size       確保するメモリサイズ (単位はbyte)
 * @param [in] alignment  アラインメント (2のべき乗を指定すること)
 * @return  アラインメントし,動的確保されたメモリ領域へのポインタ
 */
template<typename T = void*, typename std::enable_if<std::is_pointer<T>::value, std::nullptr_t>::type = nullptr>
static inline T
alignedMalloc(std::size_t size, std::size_t alignment) noexcept
{
#if defined(_MSC_VER) || defined(__MINGW32__)
  return reinterpret_cast<T>(_aligned_malloc(size, alignment));
#else
  void* p;
  return reinterpret_cast<T>(posix_memalign(&p, alignment, size) == 0 ? p : nullptr);
#endif  // defined(_MSC_VER) || defined(__MINGW32__)
}


/*!
 * @brief アラインメントされたメモリを解放する関数
 * @param [in] ptr  解放対象のメモリの先頭番地を指すポインタ
 */
static inline void
alignedFree(void* ptr) noexcept
{
#if defined(_MSC_VER) || defined(__MINGW32__)
  _aligned_free(ptr);
#else
  std::free(ptr);
#endif  // defined(_MSC_VER) || defined(__MINGW32__)
}


/*!
 * @brief std::unique_ptr で利用するアラインされたメモリ用のカスタムデリータ
 */
struct AlignedDeleter
{
  /*!
   * @brief デリート処理を行うオペレータ
   * @param [in,out] p  アラインメントされたメモリ領域へのポインタ
   */
  void
  operator()(void* p) const noexcept
  {
    alignedFree(p);
  }
};


/*!
 * @brief プラットフォームIDを取得
 * @param [in] nPlatformEntry  取得するプラットフォームID数の上限
 * @return  プラットフォームIDを格納した std::vector
 */
static inline std::vector<cl_platform_id>
getPlatformIds(cl_uint nPlatformEntry = kNDefaultPlatformEntry)
{
  std::vector<cl_platform_id> platformIds(nPlatformEntry);
  cl_uint nPlatform;
  if (clGetPlatformIDs(nPlatformEntry, platformIds.data(), &nPlatform) != CL_SUCCESS) {
    std::cerr << "clGetPlatformIDs() failed" << std::endl;
    std::exit(EXIT_FAILURE);
  }
  platformIds.resize(nPlatform);
  return platformIds;
}


/*!
 * @brief デバイスIDを取得
 * @param [in] platformId    デバイスIDの取得元のプラットフォームのID
 * @param [in] nDeviceEntry  取得するデバイスID数の上限
 * @param [in] deviceType    取得対象とするデバイス
 * @return デバイスIDを格納した std::vector
 */
static inline std::vector<cl_device_id>
getDeviceIds(const cl_platform_id& platformId, cl_uint nDeviceEntry = kNDefaultDeviceEntry, cl_int deviceType = CL_DEVICE_TYPE_DEFAULT)
{
  std::vector<cl_device_id> deviceIds(nDeviceEntry);
  cl_uint nDevice;
  if (clGetDeviceIDs(platformId, deviceType, nDeviceEntry, deviceIds.data(), &nDevice) != CL_SUCCESS) {
    std::cerr << "clGetDeviceIDs() failed" << std::endl;
    std::exit(EXIT_FAILURE);
  }
  deviceIds.resize(nDevice);
  return deviceIds;
}



/*!
 * @brief カーネル関数へ引数をまとめてセットする関数の実態
 * @param [in] kernel  OpenCLカーネルオブジェクト
 * @param [in] idx     セットする引数のインデックス
 * @param [in] first   セットする引数.可変パラメータから1つだけ取り出したもの
 * @param [in] rest    残りの引数
 * @return OpenCLのエラーコード.エラーが出た時点でエラーコードを返却する.
 */
template<typename First, typename... Rest>
static inline cl_uint
setKernelArgsImpl(const cl_kernel& kernel, int idx, const First& first, const Rest&... rest) noexcept
{
  cl_uint errCode = clSetKernelArg(kernel, idx, sizeof(first), &first);
  return errCode == CL_SUCCESS ? setKernelArgsImpl(kernel, idx + 1, rest...) : errCode;
}


/*!
 * @brief カーネル関数へ最後の引数をセットする
 * @param [in] kernel  OpenCLカーネルオブジェクト
 * @param [in] idx     引数のインデックス
 * @param [in] last    最後の引数
 * @return  OpenCLのエラーコード
 */
template<typename Last>
static inline cl_uint
setKernelArgsImpl(const cl_kernel& kernel, int idx, const Last& last) noexcept
{
  return clSetKernelArg(kernel, idx, sizeof(last), &last);
}


/*!
 * @brief カーネル関数へ引数をまとめてセットする
 * @param [in] kernel  OpenCLカーネルオブジェクト
 * @param [in] args    セットする引数群
 * @return OpenCLのエラーコード.エラーが出た時点でエラーコードを返却する.
 */
template<typename... Args>
static inline cl_uint
setKernelArgs(const cl_kernel& kernel, const Args&... args) noexcept
{
  return setKernelArgsImpl(kernel, 0, args...);
}




/*!
 * @brief このプログラムのエントリポイント
 * @return 終了ステータス
 */
int
main(int argc, char* argv[])
{
  static constexpr int ALIGN = 4096;
  static constexpr std::size_t N = 65536;

  if (argc < 2) {
    std::cerr << "Please specify only one or more source file" << std::endl;
    return EXIT_FAILURE;
  }

  // ホストのメモリを確保
  std::unique_ptr<float[], AlignedDeleter> hostX(alignedMalloc<float*>(N * sizeof(float), ALIGN));
  std::unique_ptr<float[], AlignedDeleter> hostY(alignedMalloc<float*>(N * sizeof(float), ALIGN));
  std::unique_ptr<float[], AlignedDeleter> hostZ(alignedMalloc<float*>(N * sizeof(float), ALIGN));

  // 初期化
  std::mt19937 mt((std::random_device())());
  for (std::size_t i = 0; i < N; i++) {
    hostX[i] = static_cast<float>(mt());
    hostY[i] = static_cast<float>(mt());
  }
  std::fill_n(hostZ.get(), N, 0.0f);

  // プラットフォームを取得
  std::vector<cl_platform_id> platformIds = getPlatformIds(1);

  // デバイスを取得
  std::vector<cl_device_id> deviceIds = getDeviceIds(platformIds[0], 1, CL_DEVICE_TYPE_DEFAULT);

  // コンテキストを生成
  cl_int errCode;
  std::unique_ptr<std::remove_pointer<cl_context>::type, decltype(&clReleaseContext)> context(
      clCreateContext(nullptr, 1, &deviceIds[0], nullptr, nullptr, &errCode), clReleaseContext);

  // コマンドキューを生成
  std::unique_ptr<std::remove_pointer<cl_command_queue>::type, decltype(&clReleaseCommandQueue)> cmdQueue(
      clCreateCommandQueue(context.get(), deviceIds[0], 0, &errCode), clReleaseCommandQueue);

  // デバイスが用いるメモリオブジェクトの生成
  std::unique_ptr<std::remove_pointer<cl_mem>::type, decltype(&clReleaseMemObject)> deviceX(
      clCreateBuffer(context.get(), CL_MEM_READ_WRITE, N * sizeof(float), nullptr, &errCode), clReleaseMemObject);
  std::unique_ptr<std::remove_pointer<cl_mem>::type, decltype(&clReleaseMemObject)> deviceY(
      clCreateBuffer(context.get(), CL_MEM_READ_WRITE, N * sizeof(float), nullptr, &errCode), clReleaseMemObject);
  std::unique_ptr<std::remove_pointer<cl_mem>::type, decltype(&clReleaseMemObject)> deviceZ(
      clCreateBuffer(context.get(), CL_MEM_READ_WRITE, N * sizeof(float), nullptr, &errCode), clReleaseMemObject);

  // ホストのメモリをデバイスのメモリに転送
  errCode = clEnqueueWriteBuffer(cmdQueue.get(), deviceX.get(), CL_TRUE, 0, N * sizeof(float), hostX.get(), 0, nullptr, nullptr);
  errCode = clEnqueueWriteBuffer(cmdQueue.get(), deviceY.get(), CL_TRUE, 0, N * sizeof(float), hostY.get(), 0, nullptr, nullptr);
  errCode = clEnqueueWriteBuffer(cmdQueue.get(), deviceZ.get(), CL_TRUE, 0, N * sizeof(float), hostZ.get(), 0, nullptr, nullptr);

  // コンパイル後のカーネルのバイナリを読み込み
  std::ifstream ifs(argv[1], std::ios::binary);
  if (!ifs.is_open()) {
    std::cerr << "Failed to kernel binary: " << argv[1] << std::endl;
    std::exit(EXIT_FAILURE);
  }
  std::string kernelBin((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());

  // プログラムオブジェクトの生成
  const unsigned char* kbin = reinterpret_cast<const unsigned char*>(kernelBin.c_str());
  std::size_t kbinSize = kernelBin.size();
  cl_int binStatus;
  std::unique_ptr<std::remove_pointer<cl_program>::type, decltype(&clReleaseProgram)> program(
      clCreateProgramWithBinary(context.get(), 1, &deviceIds[0], &kbinSize, &kbin, &binStatus, &errCode), clReleaseProgram);

  // カーネルソースコードのコンパイル (必要な環境もあるらしい?)
  // errCode = clBuildProgram(program.get(), 1, &deviceIds[0], nullptr, nullptr, nullptr);

  // カーネルオブジェクトの生成
  std::unique_ptr<std::remove_pointer<cl_kernel>::type, decltype(&clReleaseKernel)> kernel(
      clCreateKernel(program.get(), "vecAdd", &errCode), clReleaseKernel);

  // カーネル関数に引数を渡す
  errCode = setKernelArgs(kernel.get(), deviceZ.get(), deviceX.get(), deviceY.get(), static_cast<int>(N));

  // カーネルプログラムの実行
  errCode = clEnqueueTask(cmdQueue.get(), kernel.get(), 0, nullptr, nullptr);

  // 終了待機等
  errCode = clFlush(cmdQueue.get());
  errCode = clFinish(cmdQueue.get());

  // 実行結果をデバイスからホストへコピー
  errCode = clEnqueueReadBuffer(cmdQueue.get(), deviceZ.get(), CL_TRUE, 0, N * sizeof(float), hostZ.get(), 0, nullptr, nullptr);

  // 計算結果の確認
  for (std::size_t i = 0; i < N; i++) {
    if (std::abs(hostX[i] + hostY[i] - hostZ[i]) > 1.0e-5) {
      std::cerr << "Result verification failed at element " << i << "!" << std::endl;
      return EXIT_FAILURE;
    }
  }
  std::cout << "Test PASSED" << std::endl;

  return EXIT_SUCCESS;
}

実行は以下のようにする.

$ ./main kernel.bin

これで問題なく実行できればOKである.

雑感

個別のコンパイルプログラムを準備し,事前にオフラインコンパイルを行い,カーネルバイナリを用意するのは正直なところどうなのかと感じた. プログラムの初回起動時に,オンラインコンパイルの結果を保存し,2回目以降にその結果を再利用する形にするのがよさそうだ.

個別のコンパイルプログラムは,OpenCLのプログラムのコンパイルエラーのチェック等に用いるとよいだろう.

もう少しまともな形として作るならば,以下のような形として作るとよさそうだ.

koturn/oclc

まとめ

OpenCLのオフラインコンパイルには,オンラインコンパイルの結果を保存する手法がある.