読者です 読者をやめる 読者になる 読者になる

koturnの日記

転職したい社会人1年生の技術系日記

denite.nvimのsourceを作ってみる

Vim neovim

はじめに

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 に行ってきた

Vim neovim

はじめに

昨日,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)

Vim neovim

はじめに

前回の記事(約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のオフラインコンパイル

C++ 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のオフラインコンパイルには,オンラインコンパイルの結果を保存する手法がある.

Osaka.vim #7に行ってきた

Vim

はじめに

参加感想記事なので,今回は「ですます調」で書きます.

2ヶ月遅れの参加記事執筆となりますが,7月2日にOsaka.vimに行ってきました. 7回目のOsaka.vimだったわけですが,6回目の参加でした. ちなみに,Osaka.vimの参加記事を書くのは今回が初めてです(怠惰).

合流

大阪駅に11時頃に到着するとツイートしたら,ryunixさんに「お昼ご飯一緒にどうですか」とお誘いを受けたので,ryunixさん,thincaさん,つじけんさんと合流. Osaka.vim #1 のとき,thincaさんが参加されてたので,何とか合流できました. 会ったことが無い人が多いと,合流に少し苦労しますね.

お昼ご飯

適当に駅周辺をブラつきながら,ご飯食べるところを探してました. 駅前にいたミュージシャン(?)の演奏がうまかった!

で,しばらく歩いてお好み焼き屋に. 店員が怖かったけど,美味しかったです.

本編

今回は大学のときに所属していた研究室の後輩を誘っていたので,隣に座ってVimの知識を叩き込んでいました. 僕は最初からプラグインを導入するより,基本設定である程度やっていた方がよいと考えている派なので,最初は基本的な設定を教えました

とりあえず,カーソルキーを無効化する設定を書かせました.

noremap <Up> <Nop>
noremap <Down> <Nop>
noremap <Right> <Nop>
noremap <Left> <Nop>

最近,僕はVimに行番号は必要ない派になったので,「 set number は不要やで」と言ったのですが,それは受け入れられなかったようですw あとは,ステータスラインを表示する設定とか.

set statusline=%F%m%r%h%w\ [Format=%{&ff}]\ [TYPE=%Y]\ [ASCII=\%03.3b]\ [%p%%]\ {LEN=%L]
set laststatus=2

で,Vim7.3.xxxみたいな化石バージョンのVimを使っていたため,「Vim7.3は古い」という話になり,homebrewで新しいVimを導入させました. それに伴い,プラグインマネージャを導入し,2, 3程度のプラグインを入れるか,という話になり,結局プラグインを導入させる設定を書きました. 日本におけるプラグインマネージャといえば,dein.vimを導入さ,以下のような雛形を渡しました.

if has('vim_starting')
  let s:deindir = expand('~/.cache/dein')
  let s:deinlocal = s:deindir . '/repos/github.com/Shougo/dein.vim'
  let &runtimepath = s:deinlocal . ',' . &runtimepath
endif

if dein#load_state(s:deindir)
  call dein#begin(s:deindir)
  call dein#add('Shougo/dein.vim')
  call dein#add('itchyny/lightline.vim')

  " ここに
  " call dein#add(...)
  " でプラグインを追加する


  call dein#end()
  call dein#save_state()
endif

filetype plugin indent on
syntax enable

(しれっと雛形にサンプルとしてlightline.vimを入れておいたので,後輩のステータスラインがカッコよくなりました)

基本的にはこれでキャッシュ読み込みありでdein.vimを利用できるようになります. TOMLでプラグイン管理も可能なのですが,それは一旦置いておきました.

他には日本語のヘルプを導入させたり.... そんな感じで後輩のVim設定を増やしました.

僕自身の進捗は微妙でしたw やってたことは,job機能をいじるってことで,stdin/stdoutをゴニョゴニョとやってました. ある程度jobの使い方がわかってきたので,これからのプラグイン作成に活かしたいですね.

最後に

やっぱり,こういうVim会に参加すると,Vimに対する意欲が高まるのでいいですね.

Dein.vim で NeoBundleList と同様のことをする

Vim

はじめに

発端はこの会話.

この質問を受け,dein.vimにおいて, :NeoBundleList に相当する関数を探したが,直接的に該当するものはなかった. しかし,dein.vimの関数を組み合わせることで実現できそうだったので,実装を試みた.

実装方針

:NeoBundleList では,管理している全プラグインがそれぞれ

  1. インストールされているかどうか
  2. sourceされている(runtimepath に追加されている)かどうか

echomsg で表示する. これを目指す.

dein#get() を引数を与えずに呼び出すことで,管理プラグインの辞書を得ることができる. この辞書から必要な情報を抜き出すとよい.

インストールされているかどうか

dein.vimが想定するインストール先のパスを確認するとよい. イメージとしては以下の通り.

for pair in items(dein#get())
  echomsg pair[0] 'is' (isdirectory(pair[1].path) ? 'installed' : 'not installed')
endfor

sourceされているかどうか

それぞれのプラグインの辞書に sourced というキーがあるので,それを利用する. イメージとしては以下の通り.

for pair in items(dein#get())
  echomsg pair[0] 'is' (pair[1].sourced ? 'sourced' : 'not sourced')
endfor

実装

以上のことをまとめ,以下のコードにするとよい. ついでに, DeinList というコマンドを定義しておく.

function! s:dein_list() abort
  echomsg '[dein] #: not sourced, X: not installed'
  for pair in items(dein#get())
    echomsg (!isdirectory(pair[1].path) ? 'X'
          \ : pair[1].sourced ? ' '
          \ : '#') pair[0]
  endfor
endfunction
command! DeinList call s:dein_list()

ちなみに,空引数の dein#get()dein#_plugins の浅いコピーを返却するだけなので,以下のようにした方が無駄がない.

function! s:dein_list() abort
  echomsg '[dein] #: not sourced, X: not installed'
  for pair in items(dein#_plugins)
    echomsg (!isdirectory(pair[1].path) ? 'X'
          \ : pair[1].sourced ? ' '
          \ : '#') pair[0]
  endfor
endfunction
command! DeinList call s:dein_list()

ただ,privateを意図したような変数を直接いじるのは,やや気乗りしないものがある.

まとめ

dein.vimの関数を組み合わせることで,:NeoBundleList と同等のものを実現することができる.

MSYS2でOpenCLを使う

C++ OpenCL

はじめに

MSYS2でOpenCLのプログラムをコンパイル&実行したかった. pacmanOpenCL関連のヘッダを導入することはできるが,OpenCLのインポートライブラリは導入することはできない. ここでは,MSYS2でOpenCLの環境を導入する一連の手順を紹介する.

OpenCLのヘッダを導入する

pacmanを用いると,OpenCLのヘッダを導入できる. x64ならば,

$ pacman -S mingw-w64-x86_64-opencl-headers

x86ならば,

$ pacman -S mingw-w64-i686-opencl-headers

とするとよい.

OpenCLのインポートライブラリを作成する

上記の手順では,ヘッダファイルしか導入できない. インポートライブラリは自分で作成する必要がある. この手順は,ここを参考にした.

  • x64環境
$ mkdir lib64 && gendef - /c/Windows/system32/OpenCL.dll > lib64/OpenCL.def
$ dlltool -l lib64/libOpenCL.a -d lib64/OpenCL.def -A -k
$ mkdir lib32 && gendef - /c/Windows/SysWOW64/OpenCL.dll > lib32/OpenCL.def
$ dlltool -l lib32/libOpenCL.a -d lib32/OpenCL.def -A -k

あとは必要に応じて, /usr/local/lib 下に置くなどするとよい.

コンパイルの確認

適当に以下のようなソースコードを準備する. (面倒なので, 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


/*!
 * @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 OpenCLのデバイス情報を表示する
 * @param [in] device  OpenCLのデバイスID
 */
static inline void
showDeviceInfo(cl_device_id device) noexcept
{
  char info[2048];
  std::size_t size;

  int err = clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(info), info, &size);
  std::cout << "  Device: " << info << std::endl;
  err = clGetDeviceInfo(device, CL_DEVICE_VERSION, sizeof(info), info, &size);
  std::cout << "  CL_DEVICE_VERSION: " << info << std::endl;
}


/*!
 * @brief OpenCLのプラットフォーム情報を表示する
 * @param [in] platformId  OpenCLのプラットフォーム情報
 */
static inline void
showPlatformInfo(cl_platform_id platformId) noexcept
{
  static constexpr int MAX_DEVICE_IDS = 8;
  char info[2048];

  std::cout << "========== Platform Information ==========" << std::endl;
  std::size_t size;
  int err = clGetPlatformInfo(platformId, CL_PLATFORM_NAME, sizeof(info), info, &size);
  std::cout << "Platform: " << info << std::endl;
  err = clGetPlatformInfo(platformId, CL_PLATFORM_VERSION, sizeof(info), info, &size);
  std::cout << "CL_PLATFORM_VERSION: " << info << std::endl;

  cl_device_id devices[MAX_DEVICE_IDS]; cl_uint nDevice;
  err = clGetDeviceIDs(platformId, CL_DEVICE_TYPE_ALL, MAX_DEVICE_IDS, devices, &nDevice);
  if (err != CL_SUCCESS) {
    std::cerr << "Error (clGetDeviceIDs): " << err << std::endl;
    std::exit(EXIT_FAILURE);
  }
  for (cl_uint j = 0; j < nDevice; j++) {
    showDeviceInfo(devices[j]);
  }
  std::cout << "==========================================" << std::endl;
}


/*!
 * @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()
{
  static constexpr int ALIGN = 4096;
  static constexpr std::size_t N = 65536;
  static const char KERNEL_FILENAME[] = "kernel.cl";

  // ホストのメモリを確保
  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);

  // プラットフォームを取得
  cl_platform_id platformId;
  cl_uint nPlatform;
  cl_int errCode = clGetPlatformIDs(1, &platformId, &nPlatform);
  showPlatformInfo(platformId);

  // デバイス情報を取得
  cl_device_id deviceId;
  cl_uint nDevice;
  errCode = clGetDeviceIDs(platformId, CL_DEVICE_TYPE_DEFAULT, 1, &deviceId, &nDevice);

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

  // コマンドキューを生成
  std::unique_ptr<std::remove_pointer<cl_command_queue>::type, decltype(&clReleaseCommandQueue)> cmdQueue(
      clCreateCommandQueue(context.get(), deviceId, 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(KERNEL_FILENAME);
  if (!ifs.is_open()) {
    std::cerr << "Failed to read kernel program: " << KERNEL_FILENAME << std::endl;
    return EXIT_FAILURE;
  }
  std::string kernelSource((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());

  // プログラムオブジェクトの生成
  const char* ksrc = kernelSource.c_str();
  std::string::size_type srcSize = kernelSource.length();
  std::unique_ptr<std::remove_pointer<cl_program>::type, decltype(&clReleaseProgram)> program(
      clCreateProgramWithSource(context.get(), 1, &ksrc, &srcSize, &errCode), clReleaseProgram);

  // カーネルソースコードのコンパイル
  errCode = clBuildProgram(program.get(), 1, &deviceId, 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;
}

そして,カーネルソースコードとして以下のものを準備する. ファイル名は kernel.cl とする.

__kernel void
vecAdd(__global float* z, __global float* x, __global 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];
  }
}

コンパイルと実行は以下の通り. これで問題なくコンパイル&実行できればOKである.

$ g++ -std=gnu++11 main.cpp -Llib64 -lOpenCL -o main

余談

カーネル関数へ引数を渡す関数: clSetKernelArg() は,引数のインデックスと引数へのポインタを渡す形となっており,非常にダサい. また,カーネル関数の引数の数だけ呼び出さなければならないのも大変だ.

int n = static_cast<int>(N);
clSetKernelArg(kernel.get(), 0, sizeof(deviceZ.get()), &deviceZ.get());
clSetKernelArg(kernel.get(), 1, sizeof(deviceX.get()), &deviceX.get());
clSetKernelArg(kernel.get(), 2, sizeof(deviceY.get()), &deviceY.get());
clSetKernelArg(kernel.get(), 3, sizeof(n), &n);

うまく可変引数テンプレートを用いれば,このダサさを解消できると考え,以下のような関数の実装を行った.

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;
}


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);
}


template<typename... Args>
static inline cl_uint
setKernelArgs(const cl_kernel& kernel, const Args&... args) noexcept
{
  return setKernelArgsImpl(kernel, 0, args...);
}

実装としては,ユーザが呼び出しを行う setKernelArgs() と実際の処理を担当する setKernelArgsImpl() の2つに分けている. setKernelArgsImpl() は,可変テンプレートのアンパックとインデックスのインクリメントを行う. 引数のサイズは sizeof 演算子で取得するので,引数の型はちゃんと意識する必要がある.

呼び出し側は以下のようになる. 呼び出し側の見た目としては,

  • インデックスが必要ない
  • カーネル関数の引数をポインタとして渡す必要がない
  • 引数のサイズを渡す必要がない

ため,非常にスッキリしていると思う.

setKernelArgs(kernel.get(), deviceZ.get(), deviceX.get(), deviceY.get(), static_cast<int>(N));

C++11様々である.

まとめ

MSYS2でOpenCLを使用するには,pacmanでOepnCL関連のヘッダのインストールを行い,WindowsOpenCLのDLLからインポートライブラリを作成する必要がある.

参考文献