2011-02-28

Boost.勉強会 #4で発表しました

Boost

資料は Boost.勉強会#4 Boost.Proto にあります。他の方の発表資料については、id:thincaさんがいい感じにまとめてくれています。感謝。

さて、僕はBoost.Protoネタで発表をしてきたわけですが、前半後半と分けたために、ちょっと微妙な感じになってしまったかもしれません。あと、ぺろぺろマスターはどうみても田中さんです。また機会があったら発表したいですねー。
あと、昼御飯のつけ麺屋に15人超で入れたのにはびっくりです。

2010-12-28

gccのインラインアセンブラでレジスタに別名をつける

x86

register __m128* tmp asm ("xmm0");

asm volatile (
    "pxor %[tmp], %[tmp]"
    : [tmp] "+x" (tmp)
);

なんだかまわりくどいですね…。

このコードは、

pxor %xmm0, %xmm0

に展開されます。

2010-12-27

MOVDQUの特性について

x86

IntelのOptimization Manualの2.2.5.1 Efficient Handling of Alignment Hazardsによると、Nehalem系CPUからMOVDQUの性能が向上したとあるので、実験してみました。

4パターンのmovを実行して速度差を見ます。alignedはアクセス先のアドレスが16バイトアライメントが取られているの意味、splitはアクセス先の領域(先頭+16バイト)がcache lineを跨ないの意味です。

aligned MOVDQA
aligned MOVDQU
not-aligned split MOVDQU
not-aligned not-split MOVDQU

実行環境はCore i5 2.66GHzと、Core2 Duo 1.6GHzです。続きにあるアセンブラ片を300000000回実行しています。

	Core i5	Core2 Duo
aligned MOVDQA	0.653 sec	1.13 sec
aligned MOVDQU	0.633 sec	2.26 sec
not-aligned split MOVDQU	2.802 sec	14.947 sec
not-aligned not-split MOVDQU	0.636 sec	2.26 sec

確かにNehalem系CPUは、以前のCPUと比べてMOVDQUの性能が改善されています。cache lineを跨がなければMOVDQAと同じ速度が出ています。

以下ソースコード(ただしアセンブラ部のみ)

2010-12-22

Boost.Protoぺろぺろ

Boost Boost.Proto

最初に

このエントリは C++ Advent Calendar jp 2010 : ATND 16日目の記事です．

Boost.Protoぺろぺろ

Boost.ProtoとはBoostに含まれているExpression Template(ET)のためのライブラリです．Xpressiveや，Spirit，Phoenix v3のベースとなっています．このエントリではETについては詳しく触れませんが，式の形を型として保持する手法です．

Boost.Protoについては公式のUsers' Guide - 1.45.0がよくまとまっています．

Boost.Protoは非常に拡張性の高いライブラリです．設計を見習いたいです．ぺろぺろ．

MiniLambda

今回のネタはMiniLambdaです．Boost.Lambdaの簡易版を作ります．まぁ，User's Guideにもそういう話が載っていて，それをベースにしています．

ソースコードはhttps://gist.github.com/751465にあります．コンパイルと動作確認はg++ 4.5.1 + Boost 1.45で行いました．

今回作るminilambdaの機能(制限)は，

型はintのみ
operator +のみを提供
プレースホルダーをサポートする

としています．

includeとか

// minilambda.hpp:4
#include <boost/proto/proto.hpp>
#include <boost/preprocessor.hpp>
#include <boost/fusion/include/vector.hpp>
#include <boost/fusion/include/make_vector.hpp>
#include <boost/fusion/include/at.hpp>
#include <boost/mpl/int.hpp>
#include <boost/utility/enable_if.hpp>

namespace minilambda {

namespace proto = boost::proto;
namespace fusion = boost::fusion;
namespace mpl = boost::mpl;

using proto::lit;

今回作るminilambdaの名前空間はminilambdaとします．proto::litは，ライブラリを使う側のコードで使う可能性があるので，minilambda名前空間にusingしておく方が使いやすいと思います．boost::spirit::litもproto::litのusingです．

expression

// minilambda.hpp:23
template <class Expr>
struct expression;

struct grammar
    : proto::or_<
          proto::plus<grammar, grammar>,
          proto::terminal<proto::_>
      > {
};

struct domain
    : proto::domain<proto::generator<expression>, grammar> {
};

// minilambda.hpp:40
template <class Expr>
struct expression
    : proto::extends<Expr, expression<Expr>, domain> {
    typedef proto::extends<Expr, expression<Expr>, domain> base;

    expression(Expr const& expr)
    : base(expr) {
    }

    template <class... Args>
    int operator ()(Args&&... args) const {
        return proto::eval(*this, context<Args...>(std::forward<Args>(args)...));
    }
};

まず，domainを定義します．domainによって式と文法が関連付けられます．そしてminilambda用の式の型としてexpressionを用意します．Protoの式(この場合Expr)をラップする形になり，expressionに対して独自の動作を加えていきます．独自の動作として，operator ()によってevalを実行できるようにしています．(eval, contextについては次章) grammarはその名の通り文法を表します．grammarに定義されていない式はコンパイルエラーとなります．(_1 - _2とか)

context

// minilambda.hpp:55
template <class... Args>
struct context {
    explicit
    context(Args&&... args)
    : args_(fusion::make_vector(std::forward<Args>(args)...)) {
    }

    /* 中略 */

private:
    decltype(fusion::make_vector(std::declval<Args>()...)) args_;
};

contextは評価器の状態+評価器そのものを表します．minilambdaの場合，プレースホルダーを実装するために，ラムダ式が呼び出された時の引数を状態として持つ必要があります．せっかくのC++0xなのでVariadic Templatesで何個でも持てるようにしましょう(実際はfusion::vectorの要素数に上限がありますが)．fusion::vectorとは書けないのでfusion::make_vectorの返り値型を使って変数を定義します．

予め定義されたcontextとして，proto::default_contextやproto::callable_contextがあります．詳しくはUser's Guideを(略)．実は今回のminilambdaの用途ではcallable_contextを使った方が楽です．が，それだとUser's Guideと同じものになってしまうので…．

// minilambda.hpp:62
    template <
        class Expr,
        class Enable = void
    >
    struct eval;

    template <class Expr>
    struct eval<
        Expr,
        typename boost::enable_if<
            proto::matches<
                Expr,
                proto::terminal<proto::_>
            >
        >::type
    > {
        typedef int result_type;

        result_type
        operator ()(Expr& expr, context const& ctx) const {
            return as_value(proto::value(expr), ctx);
        }

    private:
        int as_value(int value, context const&) const {
            return value;
        }

        template <class I>
        auto as_value(placeholder<I> const&, context const& ctx) const -> decltype(fusion::at<I>(ctx.args_)) {
            return fusion::at<I>(ctx.args_);
        }
    };

    template <class Expr>
    struct eval<
        Expr,
        typename boost::enable_if<
            proto::matches<
                Expr,
                proto::plus<proto::_, proto::_>
            >
        >::type
    > {
        typedef int result_type;

        result_type
        operator ()(Expr& expr, context const& ctx) const {
            return proto::eval(proto::left(expr), ctx) +
                proto::eval(proto::right(expr), ctx);
        }
    };

一番ややこしい部分です．proto::evalを呼ぶと引数のcontextから，context::evalの型が評価関数として使われます．Exprの型によって処理を分岐していきたいのでSFINAE先生の出番となります．

proto::matchesは式に対するパターンマッチを書く事ができます．例えば，proto::_は任意の式にマッチし，proto::terminalは任意の終端記号にマッチします．このあたりのシステムは，元々はパターンマッチというより，式の変形(Transform)に使うもののようですが，僕はよく理解していません．

proto::valueはterminalの値を取る関数です．fusion::result_ofのように，proto::result_of::valueという返り値型が取れるメタ関数がありますが，C++0xではdecltypeが使えるので，あまり必要ではないかもしれません．proto::left, proto::rightは0番目の子，1番目の子ノードを取る関数です．こちらもresult_of::left, result_of::rightが定義されています．

placeholder

宣言は20行目にあります．placeholder >のように，MPL定数を抱えて使うようにします．

// minilambda.hpp:20
template <class I>
struct placeholder : I {};

そして定数は119行目から定義しています．

// minilambda.hpp:119
#define MINILAMBDA_DEFINE_PH(z, n, data)                                \
    auto const BOOST_PP_CAT(_, BOOST_PP_INC(n)) = proto::make_expr<proto::tag::terminal, domain>( \
        placeholder<mpl::int_<n> >()                                    \
    );
BOOST_PP_REPEAT(5, MINILAMBDA_DEFINE_PH, data)

_1, _2, _3, ... を終端記号(terminal)として定義します．protoの式を作るにはmake_expr関数を使います．terminal以外にもplusなど全ての式を作ることができます．make_exprにdomainを指定するとdomainが持つgeneratorの式でextendしてくれます．今回の場合，minilambda::expressionで囲われた型で式が作られます．

最後にmain

#include <iostream>
#include <boost/assert.hpp>
#include "minilambda.hpp"

int main() {
    namespace m = minilambda;

    BOOST_ASSERT(
        (m::_1 + m::_1)(100) == 200
    );

    BOOST_ASSERT(
        (m::_1 + 1)(100) == 101
    );

    BOOST_ASSERT(
        (m::_1 + m::_2)(100, 200) == 300
    );

    return 0;
}

テストコードを兼ねたmainです．

最後に

僕は説明とか，文書を書くとか，まとめを書くとかが下手なので，わかりにくい部分が多々あると思います．つっこみなどはTwitterかコメント欄でお願いします．

文章がボロボロすぎですね!

2010-11-27

CURANDライブラリを使う [Device API]

CUDA

CUDA Toolkit 3.2から，CURANDという疑似乱数ジェネレータライブラリが導入されました．GPUによる乱数の生成を行えるライブラリです．

CURANDは，2つのAPIカテゴリから構成されています．

Host API: ホストからの使用を目的としたAPI群．Device APIの上に構築されている．
Device API: デバイスからの使用を目的としたAPI群．要するに__device__修飾された関数群．

Host APIについてはCURANDライブラリを使う [Host API] - fjnlの生存記録のような何かを参照してください．

以下，サンプルコード．モンテカルロ法による円周率を求めるプログラムです．エラー処理の類は省略しています．

// -*- c++ -*-
#include <numeric>
#include <iostream>
#include <curand_kernel.h>

__global__
void monte_pi(unsigned long seed, int n_try, float* out) {
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    curandState s;

    curand_init(seed, id, 0, &s);

    int inside = 0;
    for (int i = 0; i < n_try; ++i) {
        float x = curand_uniform(&s);
        float y = curand_uniform(&s);
        if (x * x + y * y < 1) {
            inside++;
        }
    }

    out[id] = inside / static_cast<float>(n_try);
}

int main() {
    cudaSetDeviceFlags(cudaDeviceMapHost);

    int const n_thread = 256;
    int const n_block = 256;
    int const n = n_thread * n_block;
    int const n_try = 1000;

    float* out;
    float* d_out;

    cudaHostAlloc(&out, n * sizeof(float), cudaHostAllocMapped);
    cudaHostGetDevicePointer(&d_out, out, 0);

    monte_pi<<<n_block, n_thread>>>(0, n_try, d_out);

    cudaThreadSynchronize();

    std::cout
        << 4.0 * std::accumulate(out, out + n, 0.0) / n
        << std::endl;

    cudaFreeHost(out);

    return 0;
}

CURANDライブラリを使う [Host API]

CUDA

CUDA Toolkit 3.2から，CURANDという疑似乱数ジェネレータライブラリが導入されました．GPUによる乱数の生成を行えるライブラリです．

CURANDは，2つのAPIカテゴリから構成されています．

Host API: ホストからの使用を目的としたAPI群．Device APIの上に構築されている．
Device API: デバイスからの使用を目的としたAPI群．要するに__device__修飾された関数群．

今回のエントリではHost API+Pseudo Random Generatorの組み合わせのみを扱います．Quasirandom Generatorも作れるそうですが，乱数について詳しいわけではないので割愛します．

Host APIの流れは，

初期化(curandCreateGenerator)
Seed設定(curandSetPseudoRandomGeneratorSeed)
乱数列生成(curandGenerate*)
解放(curandDestroyGenerator)

となっています．CPU上でのライブラリでも似たような手順を踏む作りが多いと思います．

生成関数(curandGenerate*)は，いくつかの種類があるので，詳しくは公式ドキュメントを参照してください．今回の例で使用しているcurandGenerateUniformは0から1の間の乱数列を生成します．

なお，疑似乱数列は，生成を分割しても同じ乱数列が生成されます．1万個の乱数を得るのに1000個の生成を10回しても，10個の生成を1000回しても，結果として得られる乱数列は同じものとなります．ただ，手元の環境では生成関数の呼び出し関数を少なくした方が高速でした．

以下，サンプルコード．エラー処理の類は省略しています．

#include <iostream>
#include <iterator>
#include <curand.h>

int main() {
    curandGenerator_t g;

    cudaSetDeviceFlags(cudaDeviceMapHost);

    curandCreateGenerator(&g, CURAND_RNG_PSEUDO_DEFAULT);
    curandSetPseudoRandomGeneratorSeed(g, 0);

    int const n = 100;
    float* p;
    float* dp;
    cudaHostAlloc(&p, n * sizeof(float), cudaHostAllocMapped);
    cudaHostGetDevicePointer(&dp, p, 0);

    curandGenerateUniform(g, dp, n);
    cudaThreadSynchronize();

    std::copy(p, p + n, std::ostream_iterator<float>(std::cout, "\n"));

    curandDestroyGenerator(g);
    cudaFreeHost(p);

    return 0;
}

fstabをUUID指定にする方法

Gentoo

Blogは一ヶ月overぶりです。生きてます。

最近のUbuntuでは、/etc/fstabのパーティションの指定にUUIDを使うようになっています。そのため、従来の/dev/sda1といった指定でたびたび発生していた問題が解消されています。例えば、HDDが増えた時にsdaがsdbになったり、数年前にあったhd*がsd*へ大移動したりしてもパーティションを識別することができます。

ただ、人間はUUIDを見ただけでは、どのディスクのどのパーティションかを判断できないので、ちょっと非直感的と言えるかもしれません。

今回、あるPCのHDDを別のPCに繋いだ時に起動しなくて面倒になったので書き換えを決意しました。

UUIDの作り方

UUIDの作成にはuuidgenを使います。sys-apps/util-linuxに含まれているようです。実行するたびにUUIDを1つ生成するので、適当にシェルで、

$ for p in sda1 sda2 sda3 sda5; do echo "${p}: `uuidgen`"; done

として生成しました。

UUIDの設定方法(ext2/ext3/ext4)

tune2fsコマンドで後付けできます。

# tune2fs -U ${UUID} ${DEVICE}

UUIDの設定方法(swap)

swapはちょっと手間です。作成時しかUUIDを指定できなさそうなので、一回swapを止めて作り直しました。

# swapoff ${DEVICE}
# mkswap -U ${UUID} ${DEVICE}
# swapon ${DEVICE}

UUIDの確認方法

一番簡単なのは、/dev/disk/by-uuidを見ることだと思います。

$ cd/dev/disk/by-uuid && ls -l
lrwxrwxrwx 1 root root 2010-11-21 20:21 UUID1 -> /dev/sda1
lrwxrwxrwx 1 root root 2010-11-21 20:21 UUID2 -> /dev/sda2
...

でも…

手元の環境のパーティションはUUIDを付けてしまったので、もうわからないんですが、UUIDは生成して付与しなくてもデフォルトで付いてるんじゃないかという疑惑が…。

fjnlの生存記録のような何か

Boost.勉強会 #4で発表しました

gccのインラインアセンブラでレジスタに別名をつける

MOVDQUの特性について

Boost.Protoぺろぺろ

最初に

Boost.Protoぺろぺろ

MiniLambda

includeとか

expression

context

placeholder

最後にmain

最後に

CURANDライブラリを使う [Device API]

関連エントリ

CURANDライブラリを使う [Host API]

関連エントリ

fstabをUUID指定にする方法

UUIDの作り方

UUIDの設定方法(ext2/ext3/ext4)

UUIDの設定方法(swap)

UUIDの確認方法

でも…