2010-10-04

[CUDA] GPUデバイスからprintfする

CUDA Toolkit 3.1からデバイスからprintfを発行できるようになりました．ただし，Device Capability 2.0以降のデバイスでなければサポートされていないので，GTX480といったいわゆるFermiアーキテクチャのカードが必要です．

#include <stdio.h>

__global__
void test() {
    printf("Hello world from %d of %d\n", threadIdx.x, blockDim.x);
}

int main() {
    test<<<1, 32>>>();
    cudaThreadSynchronize();
    return 0;
}

printfが使えるようになったとはいえ，CPUのように使うわけにはいきません．制限事項については，CUDA Programming ManualのB.14のFormatted Outputを参照してください．

制限があるとはいえ，簡単にログが取れるようになったのは大きいと思います．

2010-10-03

Oven meets Boost.Spirit.Karma

Oven Boost

Boost.Spiritは元々が*1がパーサージェネレータであったこともあり、Qiが一番有名であると思いますが、Boost.Spirit一家にはKarmaという次男坊(?)がいます。

Spirit.KarmaはSpirit.Qiの逆で、データから文字列を生成するコンビネータ郡です。Performance of Numeric Generators - 1.44.0によると、iostreamやsprintfやBoost.Formatよりも高速であるとされています。また、sprintfと違い型安全な設計となっています。難点は書式指定がやや面倒なことと、コンパイル時間がboostされることですが…。

Boost.Karmaは入力に対して柔軟な設計がされており、Ovenのレンジと組合せて使うことができます。

int main() {
    std::vector<int> const xs = { 1, 2, 3, 4, 5 };

    karma::generate(
        std::ostreambuf_iterator<char>(std::cout),
        spirit::karma::int_ % ' ',
        xs | oven::transformed(std::negate<int>()));

    return 0;
}

たとえば、上記のような記述をすれば、std::vectorの中身の符号を反転さしつつ、スペースで区切って出力をすることができます。

また、oven::fuzippedといったアダプタを使えばより複雑なケースに対応することができます。fuzippedはzippedのfusion版です。zippedは型がboost::tupleとなりますが、fuzippedはboost::fusion::vectorとなります。

int main() {
    std::vector<int> xs = { 1, 2, 3, 4, 5 };

    karma::generate(
        std::ostreambuf_iterator<char>(std::cout),
        (karma::lit("negate: ") << karma::int_ << " -> " << karma::int_) % karma::eol,
        egg::pack(
            xs,
            xs | oven::transformed(std::negate<int>())
        ) | oven::fuzipped
    );

    std::cout << std::endl;

    return 0;
}

このプログラムを実行すると、

negate: 1 -> -1
negate: 2 -> -2
negate: 3 -> -3
negate: 4 -> -4
negate: 5 -> -5

という結果が得られます。

Rangeを文字列に変換できるという特徴は、他のsprintfやiostreamによる実装にはない特徴です。

*1:現在のClassic

2010-09-28

Boost.Buildで複数のバージョンのgccを切り替えて使う方法

Boost

$ bjam toolset=gcc

のようにして実行をすると、システムデフォルトのgcc, g++が使われます。しかしながら、複数のgccのバージョンを切り替えてビルドやテストを行いたいケースもあります。autoconfによるビルドを採用しているソフトウエアならば、一般的に、

$ CXX=g++-4.5 ./configure

のようにすれば利用するコンパイラを指定することができますが、bjamではこのような指定ではコンパイラを切り替えることができません。

bjamでコンパイラの指定を行うには、site-config.jamを利用します。どこのディレクトリが参照されるかはドキュメントを参照してください (http://www.boost.org/doc/tools/build/doc/html/bbv2/overview/configuration.html)。

site-config.jamに以下のようにusingの記述を追加します。左から順にコンパイラ名、バージョン、コマンド名です。オプションなども指定できるようです。

using gcc : 4.5.1 : g++-4.5.1 ;
using gcc : 4.4.4 : g++-4.4.4 ;

そして、bjam実行時に使いたいコンパイラをtoolset引数で指定します。複数のtoolsetを同時に指定することもできます。

bjam toolset=gcc-4.5.1 toolset=gcc-4.4.4

複数のコンパイラを指定した場合は順々にビルドが進みます。コンパイラAでビルドした後にコンパイラBでビルドというようになります。また、オブジェクトファイルなどのコンパイラの出力ファイルはバージョン毎に分けて出力され、まざることはありません。

2010-09-10

tupleのサイズ

ちょっと気になる所があったので，

std::tuple (C++0x)
boost::tuple
boost::fusion::vector

について大きさを調べてみました．fusion::vectorはタプルというよりも，タプル+αという感じですが…．

コンパイラによっても差が出そうだなーと思いつつも，VCがすぐに動く環境がないのでg++ 4.5.1のみです．

#include <iostream>
#include <tuple>
#include <boost/tuple/tuple.hpp>
#include <boost/fusion/include/vector.hpp>

template <class T>
void size() {
    std::cout << sizeof(T) << std::endl;
}

struct T {};
struct U {};

int main() {
    size< std::tuple<T, U> >();
    size< boost::tuple<T, U> >();
    size< boost::fusion::vector<T, U> >();

    std::cout << std::endl;

    size< std::tuple<int, T, U> >();
    size< boost::tuple<int, T, U> >();
    size< boost::fusion::vector<int, T, U> >();

    return 0;
}

このコードをコンパイルして実行すると，

という出力が得られました．

std::tupleはEBOがかかっている感じです．boost::tupleはstructに要素を並べたサイズに，fusion::vectorはそれに+1要素という印象を受けます．

2010-09-07

AES-NIについて(頓挫)

最近のIntel CPUにはAES-NIと呼ばれるAES暗号化を加速させるための専用命令群が実装されています。対応CPUは、32nmプロセスで製造されたi7, i5となっています。いろいろややこしいので、詳しくはIntelのサイトかwikipediaなどを参照してください。

最終的には動く物を作ろうと思っていたのですが、アクセス可能でAES-NI対応の環境がなかったので放置しています。環境を手に入れたら続きを書こうとは思いますが…。

AES-NIでは、以下の6命令が追加されています。x86系の命令にしては、ニーモニックが読みやすいです。なお、Intelのマニュアルでは用語の定義をfips-197から参照しています。

AESENC
AESENCLAST
AESDEC
AESDECLAST
AESIMC
AESKEYGENASSIST

コンパイラの対応状況についてはintrinsicのみ調べました。インラインアセンブラの対応状況は見ていません。

AES-NI命令のintrinsicが使えるコンパイラですが、clangが対応しています。ライセンスがないので試せませんが、Intelマニュアルには記述があるのでIntel C/C++ Compilerでも使えるようです。gcc4.5.1には後述するwmmintrin.hは含まれていないので、まだ対応していないと思われます。

clangでは、wmmintrin.hというヘッダに_mm_aesenc_si128の様に命令が宣言されています。その後コンパイルする時に-march=corei7や-maesというようにAES命令を使うように指定する必要があります。

clang -o aes aes.c -O2 -W -Wall -maes

2010-08-11

BrainfuckをJITしてみた

C++

BFだとJITしやすいので作ってみた。簡単な最適化をしています。

add/sub [rdi], immの直後にcmp [rdi], 0が来るのは無駄でしかないので、そこの最適化ぐらいはしてみたいですね。

http://gist.github.com/518146

ソースコードは400行ほどあるのでgist上に置いています。Core i5@2.66GHzの環境でhttp://esoteric.sange.fi/brainfuck/bf-source/にあるmandelbrot.bを1.4秒程度で実行する能力があります。

x86_64限定です。Mac OS XとLinux上で動く事を確認しました。DEPがONな環境だと、どうなるんでしょうね。

2010-07-29

Parsecで簡単な電卓を作ってみた

Haskell Parsec

LanguageDef と TokenParser と ExpressionParser を使ってみたかったので。。。

空白の扱いについて、ごちゃごちゃ書く必要がなくて楽でいいですね。

module Main (main) where

import Text.ParserCombinators.Parsec as Parsec
import qualified Text.ParserCombinators.Parsec.Token as Token
import qualified Text.ParserCombinators.Parsec.Language as Lang
import qualified Text.ParserCombinators.Parsec.Expr as Expr

languageDef :: Lang.LanguageDef ()
languageDef = Lang.emptyDef
              {
                Lang.opLetter = oneOf "+-*/"
              }

lexer :: Token.TokenParser ()
lexer = Token.makeTokenParser languageDef

term :: Parser Integer
term = parens expr <|> natural
    where parens = Token.parens lexer
          natural = Token.natural lexer

expr :: Parser Integer
expr = Expr.buildExpressionParser table term
    where table = [
                    [op "*" (*), op "/" div],
                    [op "+" (+), op "-" (-)]
                  ]
          op sym fn = Expr.Infix (Token.reservedOp lexer sym >> return fn) Expr.AssocLeft

program :: Parser Integer
program = Token.whiteSpace lexer >> expr >>= (\x -> eof >> return x)

main :: IO ()
main = rep >> main
    where rep = getLine >>= print . eval
          eval = either (error . show) id . parse program "input"