いんちょき通信

GnuplotグラフをWEB用に出力する

新年度が始まってから、妙にGnuplotについて書いた記事のアクセスが多いので、ここは読者に迎合してGnuplotに関する記事を書いてみようと思う。

現在日常的に使用しているPCにGnuplotが入っていなかったので、ダウンロードのため久しぶりにGnuplotの公式ページを訪れてみた。

そこのサポートフォーマットについて見てみると、「mouseable web display formats: HTML5, svg」とあり、どうも知らない間にSVGとHTML5(canvas要素)に対応しているらしい。早速、詳細を確かめるべく公式マニュアルをざっくり読んでみると、次のような記載があった。

CANVAS

set terminal canvas {size , } {background }
{font {}{,}} | {fsize }
{{no}enhanced} {linewidth }
{rounded | butt}
{solid | dashed {dashlength 
}}
{standalone {mousing} | name ''}
{jsdir 'URL/for/javascripts'}
{title ''}

SVG

set terminal svg {size , {|fixed|dynamic}}
{{no}enhanced}
{fname ""} {fsize }
{mouse} {standalone | jsdir }
{name }
{font "{,}"}
{fontfile }
{rounded|butt} {solid|dashed}

要するに、JPEGやPNG、EPSなんかの画像出力と同じようにset terminalでsvgまたはcanvasと指定してあげるとよいということらしい。案外あっけなかった。

せっかくなので、Gnuplotのでもページから拝借した次のスクリプトを出力してみる。

surface1.4.gnu

# set terminal pngcairo  transparent enhanced font "arial,10" fontscale 1.0 size 500, 350 
# set output 'surface1.4.png'
set grid nopolar
set grid xtics nomxtics ytics nomytics noztics nomztics \
 nox2tics nomx2tics noy2tics nomy2tics nocbtics nomcbtics
set grid layerdefault   linetype -1 linecolor rgb "gray"  linewidth 0.200,  linetype -1 linecolor rgb "gray"  linewidth 0.200
set samples 21, 21
set isosamples 11, 11
set title "3D gnuplot demo" 
set xlabel "X axis" 
set xlabel  offset character -3, -2, 0 font "" textcolor lt -1 norotate
set xrange [ -10.0000 : 10.0000 ] noreverse nowriteback
set ylabel "Y axis" 
set ylabel  offset character 3, -2, 0 font "" textcolor lt -1 rotate by -270
set yrange [ -10.0000 : 10.0000 ] noreverse nowriteback
set zlabel "Z axis" 
set zlabel  offset character -5, 0, 0 font "" textcolor lt -1 norotate
splot x**2+y**2, x**2-y**2, x*y, (x**3+y**3)/10

このスクリプトのsplotを実行する前に次のコードを挿入する。

set terminal svg
set output "sample_svg.svg

これで出力したSVGが次のようになる。

つぎにHTML5で出力するときも同様に

set terminal canvas
set output "sample_canvas.html"

この場合はcanvas要素とjavascriptを含んだhtmlファイルが生成される。javascriptファイルだけ生成したい場合はset terminal canvs name "sample_canvas" とする必要がある。また、「canvastext.js」、「gnuplot_common.js」、「gnuplot_mouse.css」が必要となるので、gnuplotインストールディレクトリ以下の「share/js」からコピーし、html内の参照パスを書き換える必要がある。

ちなみに出力はリンク先のようになる。

グラフを公開するときには便利な機能かなと思う。

dlopen等を利用した共有クラスの動的ロード

Linux上で共有ライブラリないのクラスを動的ロードしようと調べてみたが、時間が経つと忘れてしまいそうなので、記録しておく。

まずは、基本的なおさらいとして、Linux上で共有ライブラリ内の関数を動的ロードする手順を書いておく。

1. dlopenでライブラリを開き、ハンドルを取得する。
2. dlsymでシンボルのアドレスを取得する。
3. 2.で取得したものを利用する。
4. dlcloseでハンドルを閉じる。

続いて共有ライブラリ内のクラスを動的にロードする方法を記載する。サンプルとして使用する共有ライブラリは以下のソースの様になっている。

• libtest.hpp

#ifndef __LIBTEST_HPP__
#define __LIBTEST_HPP__

#include <iostream>
#include <string>
#include <cstdlib>

using namespace std;

namespace otter {

  class Message {
  private:
    string message;
  protected:
  public:
    static char* DEFAULT_MESSAGE;

    Message();
    Message(string message);
    ~Message();

    bool say();
  };

  typedef Message* CreateMessageDefault();
  typedef Message* CreateMessage(string);
  typedef void DestroyMessage(Message*);
}

#endif //__LIBTEST_HPP__

• libtest.cpp

#include <libtest.hpp>

using namespace std;

char* otter::Message::DEFAULT_MESSAGE="Hello World!";

extern "C" otter::Message* createDefaultObject()
{
  return new otter::Message;
}

extern "C" otter::Message* createObject(string message)
{
  return new otter::Message(message);
}

extern "C" void destroyObject(otter::Message* object)
{
  delete object;
}

otter::Message::Message()
{
  this->message = otter::Message::DEFAULT_MESSAGE;
}

otter::Message::Message(string message)
{
  this->message = message;
}

otter::Message::~Message()
{
}

bool otter::Message::say()
{
  cout << this->message << endl;
  return EXIT_SUCCESS;
}

クラスの内容としては単純で、インスタンス生成時に引数に指定した文字列もしくは引数が内場合はデフォルトの文字列をメンバ関数say()を使って出力するというものになっている。

ソースファイルないにはオブジェクト生成用の関数を定義している。dlsymでクラス自体をシンボルとして取得できないようなので、オブジェクトを生成する関数のシンボルを取得する必要があるようだ。また、シンボル取得時に関数ポインタを扱うため、型名が長くてよく分からなくなる。そのため、ヘッダーファイル内にtypedefを記載している。

これを共有ライブラリとしてビルドして、利用可能な状態にしておく。そこら辺の処理は話の本筋から外れてしまうためここでは記載しない。

続いて共有クラスを利用する方法を記載する。以下のようにmain関数を作成する。

• main.cpp

#include <libtest.hpp>
#include <dlfcn.h>

using namespace std;
using namespace otter;

int main(int argc, char* argv[], char* envp[])
{ 
  void* handle = dlopen("/home/inchoki/dev/lib/libtest.so", RTLD_LAZY);

  CreateMessageDefault* createDefault =
    (CreateMessageDefault*)dlsym(handle, "createDefaultObject");
  CreateMessage* create = (CreateMessage*)dlsym(handle, "createObject");
  DestroyMessage* destroy = (DestroyMessage*)dlsym(handle, "destroyObject");

  otter::Message* p_msg = (otter::Message*)create("Hello from Dynamic Load Lib!");
  p_msg->say();
  destroy(p_msg);

  otter::Message* p_dmsg = (otter::Message*)createDefault();
  p_dmsg->say();
  destroy(p_dmsg);

  dlclose(handle);

  return EXIT_SUCCESS;
}

基本的に共有ライブラリから関数を呼び出す流れと同じようになる。dlsymを利用してオブジェクトせいせい関数を準備し、後はその関数からオブジェクトのポインタを取得して利用するという流れになる。

当然dl*関数を利用するのでdlfcn.hが必要になり、ビルド時にはlibdlをリンクする必要がある。

試してはみたものの、プラグインの仕組みなんかが必要にならない限り、ビルド時にリンクすれば十分な気もする。

内容については次のサイトを参考にした。

• C++ dlopen mini HOWTO

PythonでシンプルなSession管理クラス

Pythonで自分用の簡単なWebアプリを作ろうと思ったが、ぱっと探した感じでは標準っぽいシンプルなセッション管理モジュールが見つからなかった。そんなわけで、ちょっと自分で書いてみた。

セッション管理の最低限の仕組みがあればよいので、次のような要件のものを作ろうと思う。

• セッションの新規作成、更新、削除ができる（当然）
• セッション情報はテキストファイルで管理する
• セッションのタイムアウトが設定できる
• 簡単なパラメータを保持ができる
• セッションIDはアクセスごとに変更する

この要件から、次のように動作するものとする。

インスタンス生成時

• セッションIDない状態でインスタンス生成 --> 新規セッションの作成（セッションIDの発行+セッションID名のテキストファイルをセッションファイルとしてセッション保存用フォルダに作成)
• 有効期限内のセッションIDを引数にインスタンス生成 --> 既存のセッションを更新(新規セッションIDを発行+既存セッションファイルを新規のセッションIDに変更)
• 有効期限切れのセッションIDを引数にインスタンス生成 --> 既存セッションを削除し、新規セッションの作成

セッションパラメータの読み書き

セッションパラメータは「key=value」の形式でテキスト内に保存し、読み書きのためのメンバー関数を用意する。

• 読み込みはkeyを指定せず、辞書型で返す
• 書き込みはkeyと valueを指定する
• 削除はkeyを指定する

セッション削除

ログオフ時を想定し、任意のタイミングでセッションを削除できるようにする。セッション削除後はセッションパラメータへのアクセスを行おうとすると、「False」を受け取る。

セッションIDの生成

セッションIDとして、アルファベット(大文字、小文字)+数字からランダムな文字列を生成する。

コード

#/usr/local/bin/python
import os
import time
import random
import string
from stat import ST_ATIME, ST_MTIME

DEFAULT_EXPIRE_SECONDS = 36000 # Default Session Expire Seconds : 10hr
DEFAULT_SESSION_PATH = './' # Default Session Storage Path : Current Dir
NULL_SESSION_ID = '' # NULL Session ID
aZ09 = string.digits + string.letters # Charactor set for Session ID

class SimpleSession:
    def __init__(self, sessionPath = DEFAULT_SESSION_PATH, 
                 sessionID = NULL_SESSION_ID,
                 expire = DEFAULT_EXPIRE_SECONDS):
        self.sessPath = sessionPath
        self.sessID = sessionID
        self.expire = expire
        # If the session id is null, a new session id will be published.
        # If the session id exists, the session id will be renamed.
        if not self.sessID:
            self.sessID = self.publishSessID()
        else:
            self.sessID = self.checkSession()


    # If the session is expired, the session is closed and a new session
    # is created. If the session is valid, the session id is renamed.
    def checkSession( self ):
        if os.path.exists(self.sessPath + self.sessID):
            sessFile = self.sessPath + self.sessID
            sessAtime = os.stat(sessFile)[ST_ATIME]
            now = time.time()
            if now - sessAtime > self.expire:
                self.close()
            
            return self.publishSessID()
        else:
            return self.publishSessID()


    # A new session ID is created. If the session exists, it copy to a new
    # one. 
    def publishSessID( self ):
        newSessID = SimpleSession.mkSessID()
        while os.path.exists(self.sessPath + newSessID):
            newSessID = self.mkSessID()
        if self.sessID:
            if os.path.exists(self.sessPath + self.sessID):
                try:
                    os.rename(self.sessPath + self.sessID,
                              self.sessPath + newSessID)
                except OSError:
                    self.close()
                        
        try:
            fd = open(self.sessPath + newSessID, 'w')
        except:
            return NULL_SESSION_ID
        fd.write('')
        fd.close()
        return newSessID


    # The session ID is closed.
    def close( self ):
        if self.sessID:
            if os.path.exists(self.sessPath + self.sessID):
                try:
                    os.remove(self.sessPath + self.sessID)
                except OSError:
                    self.sessID = NULL_SESSION_ID
                    return False
        self.sessID = NULL_SESSION_ID
        return True


    # A Key-Value Parameter is set in the session data.
    def setParam( self, key, value = ''):
        if self.sessID:
            if os.path.exists(self.sessPath + self.sessID):
                try:
                    fd = open(self.sessPath + self.sessID, 'r')
                except:
                    return False
                sessParam = {}
                for line in fd:
                    lstStr = []
                    line = line.rstrip('\n')
                    lstStr = line.split('=')
                    sessParam[lstStr[0]] = lstStr[1]
                try:
                    fd = open(self.sessPath + self.sessID, 'w')
                except:
                    return False
                sessParam[key] = value
                for k, v in sessParam.items():
                    fd.write(k + '=' + v + '\n')
                fd.close()
                return True
        
        return False


    # The session parameters is got as a dictionary value.
    def getParam( self ):
        sessParam = {}
        if self.sessID:
            if os.path.exists( self.sessPath + self.sessID):
                try:
                    fd = open(self.sessPath + self.sessID, 'r')
                except:
                    pass
                for line in fd:
                    lstStr = []
                    line = line.rstrip('\n')
                    lstStr = line.split('=')
                    if not lstStr[1]:
                        lstStr[1] = ''
                    sessParam[lstStr[0]] = lstStr[1]
        return sessParam


    # The session parameter of key will be deleted.
    def deleteParam( self, key ):
        if self.sessID:
            if os.path.exists(self.sessPath + self.sessID):
                try:
                    fd = open(self.sessPath + self.sessID, 'r')
                except:
                    return False
                sessParam = {}
                for line in fd:
                    lstStr = []
                    line = line.rstrip('\n')
                    lstStr = line.split('=')
                    sessParam[lstStr[0]] = lstStr[1]
                fd.close()
                try:
                    fd = open(self.sessPath + self.sessID, 'w')
                except:
                    return False
                for k, v in sessParam.items():
                    if k != key:
                        fd.write(k + '=' + v + '\n')
                fd.close()
                return True
        
        return False


    # A Random string for a session id is returned.
    # This is a static function.
    @staticmethod
    def mkSessID( sets = aZ09, length = 36):
        random.seed()
        newSessID =  ''.join(random.choice(sets) for i in xrange(length))
        newSessID = 'Sess_' + newSessID
        return newSessID


if __name__ == '__main__':
    print "mkSessID : " + SimpleSession.mkSessID()
    sess1 = SimpleSession("../sess/", "")
    print "SessionPath = " + sess1.sessPath
    print "SessionID = " + sess1.sessID
    print "Expire = " + str(sess1.expire)
    print "Session Param : "
    print sess1.getParam()
    print "SetParam('KEY', 'VALUE')"
    sess1.setParam('KEY', 'VALUE')
    print "Session Param : "
    print sess1.getParam()
    print "Session Dir : "
    print os.listdir(sess1.sessPath)
    print "Refresh Session Using " + sess1.sessID
    sess1 = SimpleSession("../sess/", sess1.sessID, 1000)
    print "SessionPath = " + sess1.sessPath
    print "SessionID = " + sess1.sessID
    print "Expire = " + str(sess1.expire)
    print "Session Param : "
    print sess1.getParam()
    print "SetParam('KEY', 'VALUE2')"
    sess1.setParam('KEY', 'VALUE2')
    print "SetParam('KEY2', 'VALUE2')"
    sess1.setParam('KEY2', 'VALUE2')
    print "SetParam('KEY3')"
    sess1.setParam('KEY3')
    print "Session Param : "
    print sess1.getParam()
    print "DeleteParam('KEY')"
    sess1.deleteParam('KEY')
    print "Session Param : "
    print sess1.getParam()
    print "Session Dir : "
    print os.listdir(sess1.sessPath)
    print "Session Close"
    sess1.close()
    print "Session Dir : "
    print os.listdir(sess1.sessPath)

Kindleストアで本を出版してみる

アマゾンのKindleが日本でも販売され始めたので、Kindle向けの電子書籍を出版してみることにした。

実際にKindleストアに並ぶまでの流れは後に記載するとして、出版までの道のりは実に平坦でした。「Kindle自費出版ガイド 米アマゾンの先例から学ぶ電子書籍の作り方」、「キンドルで本を売る。―キンドル・ダイレクト・パブリッシングを通じた個人出版のアドバイス」みたいな電子書籍が、それこそKindleストアにあったりするわけですが、特にこだわりなくストアに本を並べるだけであれば、何となくでできてしまうレベルです。

まず、必要なものですが、ただストアに出すのを目的とするとたいしたものは必要ありません。

• Amazonアカウント (Amazon.co.jpのものでOK)
• サポートしているファイル形式を扱えるソフトウェア

Amazonアカウントについては買い物に利用しているものがあれば、そのまま使えます。ただし、既存のアカウントに支払先情報 (住所や振込先の口座情報) を追加する必要があります。アカウントを持っていないのであれば、作ってください。出版するための情報が不足している場合は、出版の管理ページでその旨の注意が表示されます。

次にサポートしているファイル形式を扱えるソフトウェアですが、現在以下のファイルタイプがサポートされています。

• Word
• HTML
• ePub
• XMDF

となっています。ただし日本語を含むWordファイルは試験運用となっています。詳細は「Kindle Direct Publishingヘルプ」に記載されています。

今回はSigilというオープンソースのePub編集ソフトを使い原稿を作成しました。Sigilの使い方は「Sigilの使い方」なんかをちらっと見ましたが、HTMLのWYSIWYGエディタの感覚で使用できました。

ePub形式で原稿さえできてしまえば、後はたいしたことありません。淡々とアップロードに必要な項目を埋めていくだけです。

では、Amazonに原稿をアップロードするまでの手順を紹介します。まずは、Amazon.co.jpの最下部から画像の部分をクリックします。

続いて、ページ中程から「Kindleダイレクト・パブリッシング」を選択します。

続いて表示されるページでサインインすると出版物の管理ページへとつながります。

ここで、アカウントの情報が不足している場合は、追加するようにページ上方にアカウント情報更新へのリンクが表示されます。面倒な方向けにリンクを張っておきます。「Kindleダイレクト・パブリッシング」

次に「新しいタイトルを追加」を押と、出版物情報を入力するページが表示されます。

以下の項目を入力することになります。「オプション」とあるものについては、特に入力しなくても何とかなります。

• タイトル
• フリガナ
• ローマ字

このあたりは基本情報です。海外のAmazonでも販売されるため、ローマ字情報も必要となります。

• この本はシリーズ中の一冊です
• シリーズのタイトル
• 巻
• フリガナ
• ローマ字
• 版 （オプション）

シリーズものの場合はチェックをつけて、追加情報を入力します。

• 版（半角数字） （オプション）
• 出版者 （オプション）
• 出版者（ローマ字） （オプション）

• 内容紹介（4000字）

内容紹介には何かは書かないとだめみたいです。

• 本の著者等 （詳細）
• 名前
• フリガナ
• ローマ字
• 著者等 「著者、編集、はしがき、イラスト、序論、写真、序文、翻訳」から選択

人数分の著者情報を入力します。

• 言語
• 発売日 （オプション）
• ISBN （オプション）
• 参照番号 （オプション）
• この本には18歳未満の方に不適切な表現内容（文章、シーン、画像など）が含まれていますか？

• 出版する権利のステータス
• これはパブリックドメインの作品です。
• これはパブリックドメインの作品ではなく、私は必要な出版する権利を保有しています。

自分で書いたものが前提なので、下のパブリックドメインではないという方を選択します。

• カテゴリー
• 検索キーワード（最大7つまで、オプション）
• 画像のアップロード （オプション）（表紙画像）

表示画像はなくても問題ないみたいですが、かっこわるくなります。

• デジタル著作権管理(DRM)のオプションを選択してください
• デジタル著作権管理を適用する
• デジタル著作権管理を適用しない

DRMは初回のみ設定可能で、出版後の変更はできないようなので、気をつけてください。DRMを適用すると購入者は無許可で再配布ができなくなります。

• ページめくりの方向を選択してください。
• 左から右（横書き）
• 右から左（縦書き）
• 本のコンテンツファイル

ここまで入力したら、作成したePubファイルをアップロードして完了です。

アップロードするとAmazonでのレビュー状態となり、1～2日程度で次のようにストアに追加されます。

まあ、ストアに並べるのが目的で作ったものなので、内容は推して知るべしと言うことで。

<

PlatformとDeivce選択をOpenCLラップクラスに追加

「複数のOpenCLデバイスを使ってみる」で試した複数のデバイスを使うために、「外部ファイル読み取り関数をOpenCLラップクラスに追加」のOpenCLラップクラスにPlatform選択とDevice選択機能を追加した。

引数にPlatformを指定するインデックスとDeviceを指定するインデックスをとるコンストラクタを追加した。

        OpenCL::OpenCL(int platformIndex, int deviceIndex)
        {
            this->setPlatform();
            this->setContext(CL_DEVICE_TYPE_ALL, platformIndex);
            this->setDevice();
            this->queue_ = cl::CommandQueue(this->context_, this->devices_.at(deviceIndex), 0, &this->error_);
        }

この変更に伴って、setContext関数の引数を変更している。setContext関数は次の2種類が存在する。

• cl_int setContext()
• cl_int setContext(cl_device_type deviceType)

両方の引数にplatformIndexを追加し、次のように変更した。

• cl_int setContext(int platformIndex)
• cl_int setContext(cl_device_type deviceType, int platformIndex)

それぞれ、内部でcl_context_properties生成時のplataform指定を"0"固定から、platformIndexを使用するように変更している。それに伴い既存のコンストラクタ内でも明示的にplatformIndexを指定している。

変更後のソースとドキュメントはこちら。

• DECS_OpenCL_20120415をダウンロード
• Doclpp_20120415をダウンロード

このブログの開発用PCはこちら

 

外部ファイル読み取り関数をOpenCLラップクラスに追加

「OpenCLのソースを外部ファイルから読み込む」で試した外部のOpenCLソースを読み込む機能を「OpenCLをクラスでラップしてみた」、「OpenCLのラッパークラスのドキュメントを書いてみる」で書いたクラスに追加した。

• OpenCL.hpp

#ifndef __OPENCL_HPP__
#define __OPDNCL_HPP__
 
#if defined(__APPLE__) || defined(__MACOSX)
    #include <OpenCL/cl.hpp> // For Mac OSX
#else
    #include <CL/cl.hpp> // For Windows/Linux
#endif
#include <iostream>

#if defined(__ENABLE_BOOST__)
    #include <boost/filesystem/fstream.hpp>
#else
    #include <fstream>
#endif

namespace DECS{
    namespace OpenCL{

        #if defined(__ENABLE_BOOST__)
            typedef boost::filesystem::path PATH;
            typedef boost::filesystem::ifstream IFSTREAM;
        #else
            typedef std::string PATH;
            typedef std::ifstream IFSTREAM;
        #endif

        const cl_device_type DEFAULT_DEVICE_TYPE = CL_DEVICE_TYPE_GPU;

        class OpenCL{
        private:
            cl_int error_;
            std::vector platforms_;
            std::vector devices_;
            cl::Context context_;
            cl::CommandQueue queue_;
            cl::Program::Sources source_;
            cl::Program program_;
            cl::Kernel kernel_;
            std::vector argument_;
        protected:
            cl_int setPlatform();
            cl_int setContext();
            cl_int setContext(cl_device_type deviceType);
            cl_int setDevice();
        public:
            OpenCL();
            OpenCL(cl_device_type deviceType);
            ~OpenCL();
            void setSource(std::string& source);
            int setSourceFromFile(PATH srcFilePath);
            cl_int buildProgram();
            cl_int setKernel(std::string functionName);
            cl_int setKernelInputArgument(cl_int index, float* argument, const int n);
            cl_int setKernelOutputArgument(cl_int index, const int n);
            cl_int enqueueNDRange(const int n);
            cl_int readBuffer(cl_int index, float* output, const int n);
        };
    }
}

#endif // __OPENCL_HPP__

太字部分が今回追加した部分となる。プリプロセッサマクロでBoost C++ Libraryのfilesystemを使用するかどうかを選択できるようにした。"__ENABLE_BOOST__"を定義すると"boost::filesystem"を使用する。定義しない場合は標準ライブラリのstd::stringとstd::ifstreamを使用してソースファイルを読み込む。setSourceFronFileの中身は次のようになっている。

        int OpenCL::setSourceFromFile(PATH srcFilePath)
        {
            IFSTREAM srcFile(srcFilePath);
            if(!srcFile){
                return 1;
            }
            std::string src;
            std::string buf;
            while(std::getline(srcFile, buf)){
                src.append(buf);
            }
            this->setSource(src);
            return 0;
        }

前回「OpenCLのソースを外部ファイルから読み込む」で試したものとほとんど変更はない。関数内部でsetSourceを呼び出し、プライベート変数にソースを保存している。

この変更を含むすべてのソースはこちら→ DECS_OpenCL_20120414をダウンロード

前回の「OpenCLのラッパークラスのドキュメントを書いてみる」で書いたドキュメントには既に含まれている。

このブログの開発用PCはこちら

 

OpenCLのラッパークラスのドキュメントを書いてみる

「OpenCLをクラスでラップしてみた」で作成したOpenCLのラッパークラスのドキュメントをDoxygenを使って作成した。Doxygenについては「doxygenを使ってドキュメントを作成してみる」や「doxygenで複数言語のドキュメントを生成してみる」あたりでも話題にしている。

今回は複数言語の切り替えは特に行っておらず、基本的に英語で記載している。

Doxygenコメント入りのヘッダファイルが次のようになる。

• OpenCL.hpp

/******************************************************************************
* Copyright (c) 2012 INCHOKI Studio
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions are met:
*
*  1. Redistributions of source code must retain the above copyright notice,
*     this list of conditions and the following disclaimer.
*  2. Redistributions in binary form must reproduce the above copyright notice,
*     this list of conditions and the following disclaimer in the documentation
*     and/or other materials provided with the distribution.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
* AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
* ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
* LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
* CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
* SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
* INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
* CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
* ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
* POSSIBILITY OF SUCH DAMAGE
*
******************************************************************************/

/*! \file
 *  \brief DECS CL for OpenCL C++ Bindings 1.1
 *  \author INCHOKI Studio
 *
 *  \version 0.1
 *  \date April 2012
 */

/*! \mainpage
 *  \section intro Introduction
 *  OpenCL enable to give effectivness to various kinds of software. It is good
 *  acpect of OpenCL. However, It has complex prosedure to use it for hobby
 *  programings and to go through its advantages. Many lines of code must be to
 *  written before the caluculations which you want are begun. The C++ classes
 *  including in this file help to use OpenCL easily.
 *
 *  The interface for wrapping OpenCL is contained with a single header file
 *  DECS/OpenCL/OpenCL.hpp and all definitions are within the namespace 
 *  DECS::CL. However, some definitions included DECS/OpenCL/OpenCL.hpp are
 *  implemented in a binary file, for example decsopencl.lib on Windows. The
 *  binary library should be lined when the software is built.
 *
 *  For detail or other information see:
 *
 *      WASABI Tokyo (Author's Blog in English)
 *      http://www.wasabi-tokyo.net/
 *
 *      INCHOKI Journal (Author's Blog in Japanese)
 *      http://blog.inchoki.com/
 *
 *  \section example Example
 *  The following example shows a simple use case for the DECS CL
 *  for OpenCL C++ Bindings.
 *
 *  \code
 * 
 *  #include "OpenCL.hpp"
 *  #include <iostream>
 *  #include <cstdlib>
 *
 *  const int nElements = 9000000;
 *  float input1[nElements];
 *  float input2[nElements];
 *  float output[nElements];
 *
 *  int main(int argc, char* argv[])
 *  {
 *      for(int i = 0; i < nElements; i++){
 *          input1[i] = (float)i * 10.0f;
 *          input2[i] = (float)i / 20.0f;
 *          output[i] = 0.0f;
 *      }
 *
 *      try{
 *          DECS::OpenCL::OpenCL ocl;
 *          ocl.setSourceFromFile("addVector.cl");
 *          ocl.buildProgram();
 *          ocl.setKernel(std::string("addVector"));
 *          ocl.setKernelInputArgument(0, input1, nElements);
 *          ocl.setKernelInputArgument(1, input2, nElements);
 *          ocl.setKernelOutputArgument(2, nElements);
 *          ocl.enqueueNDRange(nElements);
 *          ocl.readBuffer(2, output, nElements);
 *
 *          for(int i = 0; i < 20; i++){
 *              std::cout << "input1[" << i << "], input2[" << i << "], output[" << i << "] : ";
 *              std::cout << input1[i] << ", " << input2[i] << ", " << output[i] << std::endl;
 *          }
 *
 *      }catch(cl::Error err){
 *           std::cerr << "ERROR: " << err.what() << "(" << err.err() << ")" << std::endl;
 *  
 *      }
 *
 *      return EXIT_SUCCESS;
 *  }
 *
 *  \endcode
 */
#ifndef __OPENCL_HPP__
#define __OPDNCL_HPP__
 
#if defined(__APPLE__) || defined(__MACOSX)
    #include <OpenCL/cl.hpp> // For Mac OSX
#else
    #include <CL/cl.hpp> // For Windows/Linux
#endif
#include <iostream>

#if defined(__ENABLE_BOOST__)
    #include <boost/filesystem/fstream.hpp>
#else
    #include <fstream>
#endif

/*! \namespace DECS
 *  \brief The top level namespace for Discrete Element Calculation Suite(DECS)
 *
 *  DECS(\em D escrete \em E lement \em C alculation \em S uite) is the set of
 *  utilities and libraries for Granular and Powder pysical simulation. Their
 *  definitions are contained within this namespace or sub-namespace under this
 *  namespace.
 */
namespace DECS{

/*! \namespace OpenCL
 *  \brief The DECS CL for OpenCL C++ Bindings is defined within this namespace
 *
 *  All Classes, structures, constants and macros are defined within this
 *  namespace. This namespace is one of sub-namespace of DECS.
 */
    namespace OpenCL{

        #if defined(__ENABLE_BOOST__)
            typedef boost::filesystem::path PATH;
            typedef boost::filesystem::ifstream IFSTREAM;
        #else
            typedef std::string PATH;
            typedef std::ifstream IFSTREAM;
        #endif

/*! \var cl_device_type DEFAULT_DEVICE_TYPE
 *  \brief This variable is the constant default device type.
 *
 *  DEFAULT_DEVICE_TYPE is defined for the default constructor of OpenCL Class.
 *  CL_DEVICE_TYPE_GPU is set in this variable. When the default constructor of
 *  OpenCL Class is used, only GPU devices are selected.
 */
        const cl_device_type DEFAULT_DEVICE_TYPE = CL_DEVICE_TYPE_GPU;

/*! \class OpenCL
 *  \brief This class is the main class of DECS CL for OpenCL C++ Bindings.
 *
 *  The OpenCL class supply to the series of procedures to use OpenCL in the
 *  softwares from initialization to finalization. The initialization process,
 *  for example getting platforms, contexts, devices and creating command
 *  queues is run in constructor.
 *
 *  \todo Preparing the function to load binary program (Priority: High)
 *  \todo Preparing the function to make kernel with binary program
 *  (Priority: High)
 */
        class OpenCL{
        private:
            cl_int error_;
            std::vector platforms_;
            std::vector devices_;
            cl::Context context_;
            cl::CommandQueue queue_;
            cl::Program::Sources source_;
            cl::Program program_;
            cl::Kernel kernel_;
            std::vector argument_;
        protected:
            cl_int setPlatform();
            cl_int setContext();
            cl_int setContext(cl_device_type deviceType);
            cl_int setDevice();
        public:

/*! \fn OpenCL()
 *  \brief The default constructor
 *
 *  This is the default constructor of DECS::OpenCL::OpenCL class. This
 *  constructor is initialize the class as following steps.
 *   -# All platforms set into the private variabe 
 *      std::vector platforms_ using the setPlatform
 *      protected function.
 *   -# The Context with DEFAULT_DEVICE_TYPE set into the private
 *      variable cl::Context context_ using the setContext protected
 *      function.
 *   -# All devices related with context_ set into the private varible 
 *      std::vector device_ using the setDevice protected
 *      function.
 *   -# The Command Queue related the first device in device_ set into
 *      the private varible cl::CommandQueue queue_.
 *
 *  \todo Change due to select the device creating the command queue (Priority: High)
 *  \todo Change due to select the platform creating the context (Priority: Low)
 */

            OpenCL();

/*! \fn OpenCL(cl_device_type deviceType)
 *  \brief The constructor with the device type
 *
 *  This is the constructor of DECS::OpenCL::OpenCL class with the device type.
 *  This constructor is initialize the class as following steps.
 *   -# All platforms set into the private variabe 
 *      std::vector platforms_ using the setPlatform
 *      protected function.
 *   -# The Context with the argument deviceType set into the private
 *      variable cl::Context context_ using the setContext protected
 *      function.
 *   -# All devices related with context_ set into the private varible 
 *      std::vector device_ using the setDevice protected
 *      function.
 *   -# The Command Queue related the first device in device_ set into
 *      the private varible cl::CommandQueue queue_.
 *
 *  @param[in] deviceType The device type of context.
 *  
 *  All parameters are defined in OpenCL.
 *
 *  For detail see: The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  \todo Change due to select the device creating the command queue (Priority: High)
 *  \todo Change due to select the platform creating the context (Priority: Low)
 */
            OpenCL(cl_device_type deviceType);

/*! \fn ~OpenCL()
 *  \brief The default destructor
 *
 *  The default destructor does not have the specific process at this time. 
 */
            ~OpenCL();

/*! \fn setSource(std::string& source)
 *  \brief The function to set source strings for the program
 *
 *  This function set the source strings into the private variable
 *  cl::Program::Sources source_. cl::Program::Sources is
 *  std::vector for source string and its size. When this function
 *  call second time and more, the source strings push back the std::vector.
 *  The size of std::vector is managed in this function, so the user don't need
 *  to conserned with the size.
 *
 *  @param[in] source The source string refference.
 */

            void setSource(std::string& source);

/*! \fn setSourceFromFile(std::string& srcFile)
 *  \brief The function to load and set a source file
 *
 *  This function load a souce from source file and set its contents into the
 *  private variable source_. This function call the function setSource
 *  in its internal process. If you define the macro "__ENABLE_BOOST__", the
 *  type "PATH" is "boost::filesystem::path". And "std::string" is used in the
 *  default settings.
 *
 *  @param[in] srcFilePath The path to source file.
 *  \return 0: Success, 1: "srcFilePath cannot be opened
 */
            int setSourceFromFile(PATH srcFilePath);

/*! \fn buildProgram()
 *  \brief The function to build the program
 *
 *  This fucntion initialize the private variable program_ and build the
 *  programs from the sources stored in the private variable source_. If an
 *  error occur in initializing program, this function return the error code of
 *  cl::Program and throw the exception. If an error occur in building program,
 *  this function return the error code of cl::Program::Build and throw the
 *  exception. These exception is defined in OpenCL C++ Bindings.
 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  \return The error code of cl::Program or cl::Program::Build.
 */

            cl_int buildProgram();

/*! \fn setKernel(std::string functionName)
 *  \brief The function to set kernel
 *
 *  This function set the kernel with the OpenCL function included in the
 *  private variable program_instance into the private variable cl::Kernel
 *  kernel_. If an error occur in making the kernel instance, this function
 *  return the error code of cl::Kernel and throw the OpenCL C++ Bindings
 *  exception. 
 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  @param[in] functionName The function name to set into kernel.
 *  \return The error code of cl::Kernel.
 *
 *  \todo Change due to use binary program (Priority: High)
 */

            cl_int setKernel(std::string functionName);

/*! \fn setKernelInputArgument(cl_int index, float* argument, const int n)
 *  \brief The function to set the input argument into the private variable kernel_
 *
 *  This Function set the input argument for kernel into the private variable kernel_.
 *  The memory object is create based on the pointer of argument and the number
 *  of element of argument. In actual, this memory object is set into the kernel.
 *  If a error occur in making the memory object, this function return the error
 *  code of cl::Buffer and throw the exception. And if a error occur in setting
 *  the argument into the kernel, this function return the error code of
 *  cl::kernel::setArg and throw the exception.
 *
 *  \warning At this time this function can deal with only float type. And this 
 *  function can only deal with the memory object with only CL_MEM_READ_ONLY.

 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  @param[in] index The index of argument order of the kernel.
 *  @param[in] argument The pointer of a argument for the kernel.
 *  @param[in] n The number of element of the argument.
 *  \return The error code of cl::Buffer or cl::Kernel::setArg.
 *
 *  \todo Change due to deal with other type argument (Priority: High)
 *  \todo Change due to switch multi kernels (Priority: Mid)
 */

            cl_int setKernelInputArgument(cl_int index, float* argument, const int n);

/*! \fn setKernelOutputArgument(cl_int index, const int n)
 *  \brief The function to set the output argument into the private variable kernel_
 *
 *  This Function set the output argument for kernel into the private variable kernel_.
 *  The memory object is create based on the pointer of argument and the number
 *  of element of argument. In actual, this memory object is set into the kernel.
 *  If a error occur in making the memory object, this function return the error
 *  code of cl::Buffer and throw the exception. And if a error occur in setting
 *  the argument into the kernel, this function return the error code of
 *  cl::kernel::setArg and throw the exception.
 *
 *  \warning At this time this function can deal with only float type. And this 
 *  function can only deal with the memory object with only CL_MEM_WRITE_ONLY.
 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  @param[in] index The index of argument order of the kernel.
 *  @param[in] n The number of element of the argument.
 *  \return The error code of cl::Buffer or cl::Kernel::setArg.
 *
 *  \todo Change due to deal with other type argument (Priority: High)
 *  \todo Change due to switch multi kernels (Priority: Mid)
 */
            cl_int setKernelOutputArgument(cl_int index, const int n);

/*! \fn enqueueNDRange(const int n)
 *
 *  \brief The function to enqueue the kernel
 *
 *  This function enqueue the private variable kernel_ into the command queue
 *  which created as the private variable queue_. If a error occur, this
 *  function return the error code of cl::CommandQueue::enqueueNDRangeKernel,
 *  and throw the exception of its.
 *
 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  @param[in] n The number of index of the arguments.
 *
 *  \return The error code of cl::CommandQueue::enqueueNDRangeKernel
 */
            cl_int enqueueNDRange(const int n);

/*! \fn readBuffer(cl_int index, float* output, const int n)
 *  \brief The fucntion to read the buffer
 *
 *  This function read the output from the buffer into the argument
 *  "output". If a error occur duaring the reading buffer, this function
 *  return the error code of cl::CommandQueue::enqueueReadBuffer and throw
 *  the exception of cl::CommandQueue::enqueueReadBuffer and throw.

 *  For detail of the error code see:
 *
 *  The OpenCL Specification
 *  http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf
 *
 *  For Detail of the exception see:
 *
 *  The OpenCL C++ Wrapper API
 *  http://www.khronos.org/registry/cl/specs/opencl-cplusplus-1.1.pdf
 *
 *  \warning At this time this function can deal with only float type. 
 *
 *  @param[in] index The index of argument order of the kernel.
 *  @param[out] output The variable to store the output.
 *  @param[in] The number of element of the argument.
 *
 *  \return The error code of cl::CommandQueue::enqueueReadBuffer
 */
            cl_int readBuffer(cl_int index, float* output, const int n);
        };
    }
}

#endif // __OPENCL_HPP__

これからDoxygenを使って生成したドキュメントはこちら(doclpp_201204014.pdf)。

このブログの開発用PCはこちら

複数のOpenCLデバイスを使ってみる

IntelのOpenCLドライバもあるらしいので、インストールした流れで、AMD RadeonとCore i7両方を使って計算をしてみることにした。

そのまえに、ちゃんとIntelドライバを認識しているか、AMD Accelerated Parallel Processing (APP) SDKに付属している"clinfo.exe"で確認した。

C:\Program Files (x86)\AMD APP\bin\x86_64>clinfo.exe
Number of platforms:                             2
  Platform Profile:                              FULL_PROFILE
  Platform Version:                              OpenCL 1.1 AMD-APP (831.4)
  Platform Name:                                 AMD Accelerated Parallel Proces
sing
  Platform Vendor:                               Advanced Micro Devices, Inc.
  Platform Extensions:                           cl_khr_icd cl_amd_event_callbac
k cl_amd_offline_devices cl_khr_d3d10_sharing
  Platform Profile:                              FULL_PROFILE
  Platform Version:                              OpenCL 1.1
  Platform Name:                                 Intel(R) OpenCL
  Platform Vendor:                               Intel(R) Corporation
  Platform Extensions:                           cl_khr_fp64 cl_khr_global_int32
_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomi
cs cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_intel_pr
intf cl_ext_device_fission cl_intel_immediate_execution cl_khr_gl_sharing cl_khr
_icd

こんな感じの出力が出て、確かにIntel提供のプラットフォームが増えている。

「OpenCLで実際に計算してみる」のコードを単純にAMD用とIntel用に二重化して動かしてみた。

#include <cstdio>
#include <cmath>
#include <boost/chrono.hpp>
#include <boost/chrono/duration.hpp>

using namespace boost;

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

const int nElements = 9000000;
const int maxDevices  = 10;

float input1[nElements];
float input2[nElements];
float output[nElements];

void addVector(float* input1, float* input2, float* output){
	for(int i = 0; i < nElements; i++){
		output[i] = sin(input1[i]) * sin(input2[i]);
		output[i] = cos(output[i]);
		output[i] = pow(output[i], output[i]);
	}
}

int main(int argc, char* argv[])
{
    cl_int status;

    cl_platform_id platforms[10];
    cl_uint num_platforms;
    status = clGetPlatformIDs(sizeof(platforms) / sizeof(platforms[0]),
		platforms,
		&num_platforms);

    std::cout << "Number Of Platforms : " << num_platforms << std::endl;

    cl_context_properties amdProp[]
        = {CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[0], 0};

    cl_context_properties intelProp[]
        = {CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[1], 0};

    cl_context amdContext = clCreateContextFromType(amdProp,
        CL_DEVICE_TYPE_GPU,
        NULL,
        NULL,
        &status);

    cl_context intelContext = clCreateContextFromType(intelProp,
        CL_DEVICE_TYPE_CPU,
        NULL,
        NULL,
        &status);

    cl_device_id amdDev[maxDevices];
    size_t size_return;
    status = clGetContextInfo(amdContext,
		CL_CONTEXT_DEVICES,
		sizeof(amdDev),
		amdDev,
		&size_return);

    cl_device_id intelDev[maxDevices];
    status = clGetContextInfo(intelContext,
		CL_CONTEXT_DEVICES,
		sizeof(intelDev),
		intelDev,
		&size_return);

    cl_command_queue amdQueue = clCreateCommandQueue(amdContext,
        amdDev[0],
        0,
        &status);

    cl_command_queue intelQueue = clCreateCommandQueue(intelContext,
        intelDev[0],
        0,
        &status);

    const char *sources[] = {
        "__kernel void\n\
        addVector(__global const float *input1,\n\
        __global const float *input2,\n\
        __global float *output)\n\
        {\n\
        int index = get_global_id(0);\n\
		output[index] = sin(input1[index]) * sin(input2[index]);\n\
		output[index] = cos(output[index]);\n\
		output[index] = pow(output[index], output[index]);\n\
    }\n"};

    
    cl_program amdProgram = clCreateProgramWithSource(amdContext,
        1,
        (const char**)&sources,
        NULL,
        &status);

    cl_program intelProgram = clCreateProgramWithSource(intelContext,
        1,
        (const char**)&sources,
        NULL,
        &status);

    status = clBuildProgram(amdProgram, 1, amdDev, NULL, NULL, NULL);
    status = clBuildProgram(intelProgram, 1, intelDev, NULL, NULL, NULL);
    clUnloadCompiler();

    cl_kernel amdKernel = clCreateKernel(amdProgram, "addVector", &status);
    cl_kernel intelKernel = clCreateKernel(intelProgram, "addVector", &status);

    for (int i = 0; i < nElements; i++) {
        input1[i] = (float)i * 10.0f;
        input2[i] = (float)i / 20.0f;
        output[i] = 0.0f;
    }

    cl_mem amdMemInput1 = clCreateBuffer(amdContext,
		CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		sizeof(cl_float) * nElements,
		input1,
		&status);

    cl_mem amdMemInput2 = clCreateBuffer(amdContext,
        CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        sizeof(cl_float) * nElements,
        input2,
        &status);

    cl_mem amdMemOutput = clCreateBuffer(amdContext,
        CL_MEM_WRITE_ONLY,
        sizeof(cl_float) * nElements,
        NULL,
        &status);

    cl_mem intelMemInput1 = clCreateBuffer(intelContext,
		CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		sizeof(cl_float) * nElements,
		input1,
		&status);

    cl_mem intelMemInput2 = clCreateBuffer(intelContext,
        CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        sizeof(cl_float) * nElements,
        input2,
        &status);

    cl_mem intelMemOutput = clCreateBuffer(intelContext,
        CL_MEM_WRITE_ONLY,
        sizeof(cl_float) * nElements,
        NULL,
        &status);
    status = clSetKernelArg(amdKernel, 0, sizeof(cl_mem), (void *)&amdMemInput1);
    status = clSetKernelArg(amdKernel, 1, sizeof(cl_mem), (void *)&amdMemInput2);
    status = clSetKernelArg(amdKernel, 2, sizeof(cl_mem), (void *)&amdMemOutput);

    status = clSetKernelArg(intelKernel, 0, sizeof(cl_mem), (void *)&intelMemInput1);
    status = clSetKernelArg(intelKernel, 1, sizeof(cl_mem), (void *)&intelMemInput2);
    status = clSetKernelArg(intelKernel, 2, sizeof(cl_mem), (void *)&intelMemOutput);

	 boost::chrono::steady_clock::time_point start =
		 boost::chrono::steady_clock::now();

    size_t globalSize[] = {nElements};
    status = clEnqueueNDRangeKernel(amdQueue,
        amdKernel,
        1,
        NULL,
        globalSize,
        0, 0, NULL, NULL);

    status = clEnqueueReadBuffer(amdQueue,
        amdMemOutput,
        CL_TRUE,
        0,
        sizeof(cl_float) * nElements,
        output, 0, NULL, NULL);

	boost::chrono::duration sec =
		boost::chrono::steady_clock::now() - start;

    std::cout << "input1, input2, output" << std::endl;
    for (int i = 0; i < 100; i++) {
        std::cout << input1[i] << ", " << input2[i] << ", " << output[i] << std::endl;
    }
	
    std::cout << "OpenCL addVector(AMD) " << sec.count() << " seconds\n";

	 boost::chrono::steady_clock::time_point start_i =
		 boost::chrono::steady_clock::now();

    status = clEnqueueNDRangeKernel(intelQueue,
        intelKernel,
        1,
        NULL,
        globalSize,
        0, 0, NULL, NULL);

    status = clEnqueueReadBuffer(intelQueue,
        intelMemOutput,
        CL_TRUE,
        0,
        sizeof(cl_float) * nElements,
        output, 0, NULL, NULL);

	boost::chrono::duration sec_i =
		boost::chrono::steady_clock::now() - start_i;

    std::cout << "input1, input2, output" << std::endl;
    for (int i = 0; i < 100; i++) {
        std::cout << input1[i] << ", " << input2[i] << ", " << output[i] << std::endl;
    }
	
    std::cout << "OpenCL addVector(intel) " << sec_i.count() << " seconds\n";

    clReleaseMemObject(amdMemOutput);
    clReleaseMemObject(amdMemInput2);
    clReleaseMemObject(amdMemInput1);
    clReleaseKernel(amdKernel);
    clReleaseProgram(amdProgram);
    clReleaseCommandQueue(amdQueue);
    clReleaseContext(amdContext);
    clReleaseMemObject(intelMemOutput);
    clReleaseMemObject(intelMemInput2);
    clReleaseMemObject(intelMemInput1);
    clReleaseKernel(intelKernel);
    clReleaseProgram(intelProgram);
    clReleaseCommandQueue(intelQueue);
    clReleaseContext(intelContext);

	boost::chrono::steady_clock::time_point start2 =
		boost::chrono::steady_clock::now();

	addVector(input1, input2, output);

	boost::chrono::duration sec2 =
		boost::chrono::steady_clock::now() - start2;

	std::cout << "Normal addVector " << sec2.count() << " seconds\n";
    return 0;
}

amd*はRadeon用、intel*はCore i7用として、それぞれ実行時間を計測した。

実行時間部分の結果は以下のようになる。

OpenCL addVector(AMD) 0.129915 seconds
OpenCL addVector(intel) 0.114104 seconds
Normal addVector 1.12876 seconds

何回か実行したが、若干Intelプラットフォームの方が実行時間は短かった。

マルチスレッド化して、スレッドごとにデバイスを割り当てるなんてこともできそうな気がする。このあたりの結果を「OpenCLをクラスでラップしてみた」のクラスにでも反映していこうかなと思う。

このブログの開発用PCはこちら

微小時間刻み運動の要素パラメータ用データストア

前回の「微小時間刻み運動の要素をクラス化してみる」で微小時間刻み運動の要素をクラス化してみたが、どうにも使い勝手が悪そうなので、何かできないかと色々試した見ることにした。

ひとまず、データの柔軟性と将来的にOpenCLで並列処理をする際のデータの受け渡しなどを考えて、要素クラスからデータの実体を分離してみることにした。メインロジックとしては要素クラスを通してデータを操作するが、並列処理をする場合は全要素についてフラットなデータ空間の方が扱いやすいのではないかという予想してみた。なんか前にチラッと見たデザインパターンのフライウェイトとかにも近いような気がする。当然データを分離しているので、タイムステップの更新の際に実データの移動は発生せず、フライウェイトパターンと同じような利点もあると予想している。

実際のコードは次のようになっている。

• constant.hpp

#ifndef __CONSTANT_HPP__
#define __CONSTANT_HPP__

namespace DECS{
    // Constants for General
    typedef const unsigned int id_t;
    typedef const unsigned int dim_t;
    id_t DEFAULT_DIMENSION = 3;

    enum AXIS{
        X,
        Y,
        Z
    };

    enum RAXIS{
        RX,
        RY,
        RZ
    };

    // Constants for DataStore
    typedef unsigned int n_parameter_t;
    const n_parameter_t DEFAULT_PARAMETER_NUMBERS = 7;
    typedef unsigned int data_size_t;
    const data_size_t DEFAULT_DATASTORE_SIZE = 10;
    typedef unsigned int index_t;
}
#endif // __CONSTANT_HPP__

• DataStore.hpp

#ifndef __DATASTORE_HPP__
#define __DATASTORE_HPP__

#include <decs/constant.hpp>

namespace DECS{

    class DataStore{
    private:
        data_size_t     size_;
        data_size_t     wholeSize_;
        dim_t           dim_;
        n_parameter_t   nParam_;
        double*         data_;
    protected:
        void    fillZero();
    public:
        DataStore();
        DataStore(data_size_t size, dim_t dim, n_parameter_t nParam);
        DataStore(const DataStore& obj);
        ~DataStore();

        double* getStorePtr();
        double* getStorePtr(id_t id, index_t index);

        DataStore operator=(const DataStore& obj) const;
    };
}

#endif // __DATASTORE_HPP__

• DataStore.cpp

#include <decs/DataStore.hpp>

namespace DECS{
    DataStore::DataStore() : size_(DEFAULT_DATASTORE_SIZE),
                             dim_(DEFAULT_DIMENSION),
                             nParam_(DEFAULT_PARAMETER_NUMBERS)
    {
        this->wholeSize_ = this->size_ * this->dim_ * this->nParam_;
        this->data_ = new double[this->wholeSize_];
        this->fillZero();
    }

    DataStore::DataStore(data_size_t size, dim_t dim, n_parameter_t nParam) :
        size_(size), dim_(dim), nParam_(nParam)
    {
        this->wholeSize_ = this->size_ * this->dim_ * this->nParam_;
        this->data_ = new double[this->wholeSize_];
        this->fillZero();
    }

    DataStore::DataStore(const DataStore& obj) :
        size_(obj.size_), dim_(obj.dim_), nParam_(obj.nParam_), wholeSize_(obj.wholeSize_)
    {
        this->wholeSize_ = this->size_ * this->dim_ * this->nParam_;
        this->data_ = new double[this->wholeSize_];
        for(data_size_t i = 0; i < this->wholeSize_; i++){
            this->data_[i] = obj.data_[i];
        }
    }

    DataStore::~DataStore()
    {
        delete [] this->data_;
    }

    void    DataStore::fillZero()
    {
        for(data_size_t i = 0; i < this->wholeSize_; i++){
            this->data_[i] = 0.0;
        }
    }

    double* DataStore::getStorePtr()
    {
        return this->data_;
    }

    double* DataStore::getStorePtr(id_t id, index_t index)
    {
        if(id <= this->size_ && index <= this->nParam_){
            return this->data_ + id * this->nParam_ * this->dim_
                + index * this->dim_;
        }
        return 0;
    }

    DataStore DataStore::operator=(const DataStore& obj) const
    {
        DataStore rtn(obj);
        return rtn;
    }
}

constant.hppは単純に定数や列挙体をまとめて定義しているだけのヘッダーファイルとなっている。データ自体はdata_として1次元の配列として保持していて、getStorePtr(id_t　id, index_t index)でid番目の要素のindex番目の3次元ベクトル(デフォルトでは)パラメータのポインターを渡すようにした。

また、OpenCLにフラットなデータを渡せるようにgetStorePtr()でdata_の先頭のポインターを渡すようにしている。ただ、ポインターを渡してしまうならクラスじゃなくて構造体でいいのではないかとう気もする。うーん。

コピーコンストラクタとオペレータ"="は結果出力時にデータを対比することを想定している。

なんかおとなしくstd::vectorを使えばいいじゃないかという気もするが、std::vectorってそのままじゃOpenCLカーネルには渡せない気がするし、いちいちデータを作っていたらそれはそれでオーバーヘッドになるなと思うしで、色々迷うところだ。

このブログの開発用PCはこちら

微小時間刻み運動の要素をクラス化してみる

かなり前の記事「放物運動を微少時間刻みで考えてみる」で作成したコードの要素をクラス化してみようと思う。要素数が増えたときの対応や、違う現象を表現しようとするときにクラス化しているのと、していないのとでは、効率がかなり違う気がする。

クラスに含めるパラメータは

• 現在の位置
• 1ステップ前の位置
• 現在の速度
• 1ステップ前の速度

ひとまず、これらをひとまとめとして扱えるようにする。

• Element.hpp

#ifndef __ELEMENT_HPP__
#define __ELEMENT_HPP__

namespace DECS{

    enum RELATIVE_STEP{
        CURRENT,
        PREVIOUS
    };

    enum AXIS{
        X,
        Y
    };

    class Element{
    private:
        double curPos_[2];
        double prePos_[2];
        double curVel_[2];
        double preVel_[2];
    protected:
    public:
        Element();
        Element(double x, double y, double vx, double vy);
        ~Element();

        double getPos(RELATIVE_STEP rs, AXIS axis) const;
        double getVel(RELATIVE_STEP rs, AXIS axis) const;
        void setPos(RELATIVE_STEP rs, AXIS axis, double value);
        void setVel(RELATIVE_STEP rs, AXIS axis, double value);
    };
}

#endif // __ELEMENT_HPP__

• Element.cpp

#include "Element.hpp"

namespace DECS{
    Element::Element()
    {
        this->curPos_[0] = 0.0;
        this->curPos_[1] = 0.0;
        this->prePos_[0] = 0.0;
        this->prePos_[1] = 0.0;
        this->curVel_[0] = 0.0;
        this->curVel_[1] = 0.0;
        this->preVel_[0] = 0.0;
        this->preVel_[1] = 0.0;
    }

    Element::Element(double x, double y, double vx, double vy)
    {
        this->curPos_[0] = x;
        this->curPos_[1] = y;
        this->prePos_[0] = x;
        this->prePos_[1] = y;
        this->curVel_[0] = 0.0;
        this->curVel_[1] = 0.0;
        this->preVel_[0] = vx;
        this->preVel_[1] = vy;
    }

    Element::~Element()
    {
    }

    double Element::getPos(RELATIVE_STEP rs, AXIS axis) const
    {
        switch(rs){
        case CURRENT:
            return this->curPos_[axis];
        case PREVIOUS:
            return this->prePos_[axis];
        default:
            return 0.0;
        }
    }

    double Element::getVel(RELATIVE_STEP rs, AXIS axis) const
    {
        switch(rs){
        case CURRENT:
            return this->curVel_[axis];
        case PREVIOUS:
            return this->preVel_[axis];
        default:
            return 0.0;
        }
    }

    void Element::setPos(RELATIVE_STEP rs, AXIS axis, double value)
    {
        switch(rs){
        case CURRENT:
            this->curPos_[axis] = value;
            break;
        case PREVIOUS:
            this->prePos_[axis] = value;
            break;
        }
    }

    void Element::setVel(RELATIVE_STEP rs, AXIS axis, double value)
    {
        switch(rs){
        case CURRENT:
            this->curVel_[axis] = value;
            break;
        case PREVIOUS:
            this->preVel_[axis] = value;
            break;
        }
    }
}

アクセッサでは列挙RELATIVE_STEPを使用して、相対的なステップを指定できるようにしたが、1行の記述が長くなってしまうので、素直に相対ステップごとにアクセッサを用意した方がいいような気もする。

このクラスを使ってメイン関数を書き直すと次のようになる。

#include <fstream>
#include <cstdlib>

#include "Element.hpp"

using namespace std;

int main(int argc, char* argv[])
{
    DECS::Element element(0.0, 400.0, 30.0, 0.0);
    double dt = 0.0001; // 1ステップの微少時間
    double a[2] = {0.0, -9.81}; // 加速度
    int totalStep = 100000; // 総ステップ数
    double p[2][200]; // 保存用座標
    ofstream outf("./output.txt"); //結果出力ファイル
    int count = 0;
    p[0][0] = element.getPos(DECS::CURRENT, DECS::X);
    p[1][0] = element.getPos(DECS::CURRENT, DECS::Y);
    for(int i = 1; i < totalStep; i++){
        element.setVel(DECS::CURRENT, DECS::X, element.getVel(DECS::PREVIOUS, DECS::X) + a[0] * dt);
        element.setVel(DECS::CURRENT, DECS::Y, element.getVel(DECS::PREVIOUS, DECS::Y) + a[1] * dt);
        element.setPos(DECS::CURRENT, DECS::X, element.getPos(DECS::PREVIOUS, DECS::X) + 
            (element.getVel(DECS::CURRENT, DECS::X) + element.getVel(DECS::PREVIOUS, DECS::X)) * dt / 2);
        element.setPos(DECS::CURRENT, DECS::Y, element.getPos(DECS::PREVIOUS, DECS::Y) + 
            (element.getVel(DECS::CURRENT, DECS::Y) + element.getVel(DECS::PREVIOUS, DECS::Y)) * dt / 2);
        count++;
        if(count == 500){
            count = 0;
            p[0][i/500] = element.getPos(DECS::CURRENT, DECS::X);
            p[1][i/500] = element.getPos(DECS::CURRENT, DECS::Y);
        }
        element.setPos(DECS::PREVIOUS, DECS::X, element.getPos(DECS::CURRENT, DECS::X));
        element.setPos(DECS::PREVIOUS, DECS::Y, element.getPos(DECS::CURRENT, DECS::Y));
        element.setVel(DECS::PREVIOUS, DECS::X, element.getVel(DECS::CURRENT, DECS::X));
        element.setVel(DECS::PREVIOUS, DECS::Y, element.getVel(DECS::CURRENT, DECS::Y));
    }
    for(int i = 0; i < 200; i++){
        outf << p[0][i] << endl;
        outf << p[1][i] << endl;
    }
    return EXIT_SUCCESS;
} 

やはり、パラメータを使って値を更新していく記述がゴチャゴチャしてしまうので、なにかシンプルになるような仕様にした方がいいように思えてきた。また、現時点の状態では表現できないことが多すぎて、全然使い物にならない。とりあえずやってみ たレベルなので、ちょっと手を入れていこうと思う。

結果は「放物運動を微少時間刻みで考えてみる」と同じようになる。

 

このブログの開発用PCはこちら