Pyston試してみた

Pystonに関しては以下のニュースを参照。
http://sourceforge.jp/magazine/14/04/08/153000

ビルド方法がしっかりとドキュメントに記載されているので、その通りにやれば問題なかった。ちなみにLinux Mint15 x64で動かしている。

とりあえず"Hello World"やな、ってことで下記を実行する。

print "Hello World"

出力結果が下記です。

$ ./pyston ~/main.py

  715us to load stdlib
 1.8ms for initCodegen
1.8ms for jit startup
 25us parsing
Parsed code; ast:
print "Hello World"
==============
JIT'ing module with signature () -> void at effort level 0
CFG:
1 blocks
Block 0 'entry'; Predecessors: Successors:
    print "Hello World"
    return 
processing types for block 0
processing opt block 0
generated IR:

define void @module_e0_0() {
opt_block0:
  %0 = call i8* (...)* @_IO_printf(i8* inttoptr (i64 28776752 to i8*), i8* inttoptr (i64 28776592 to i8*)), !dbg !8
  %1 = call i8* (...)* @_IO_printf(i8* inttoptr (i64 28776784 to i8*)), !dbg !8
  ret void, !dbg !8
}

   282us in compileFunction
Compiling...
Compiled function to (nil)
  301us for _doCompile()
 320us for compileModule()
compiled module.main to machine code; running:
 5us to interpret
Hello World
finished running
377us to run
In teardownRuntime
23us joinRuntime
Stats:
OSR exits: 0
interpreted_runs: 1
num_compiles: 1
num_compiles_0_interpreted: 1
num_hidden_classes: 153
reopts: 0
us_compiling: 301
us_compiling_0_interpreted: 301
us_compiling_irgen: 282
us_interpreting: 6
us_parsing: 25
158us finishing up

おー、ちゃんとIRはいて実行してくれてる！

まだ開発段階でサポートしてない構文もあるみたいだが、ループくらいできるだろうと思って、一億回のループ回してみたら処理が返ってこなくなった。。。
100回だと返ってきたので、その結果も貼り付けておく。

a = 0
for i in range(100):
    a+=1

  458us to load stdlib
 1.3ms for initCodegen
1.3ms for jit startup
 5.0ms parsing
Parsed code; ast:
a = 0
<for loop>

==============
JIT'ing module with signature () -> void at effort level 0
CFG:
9 blocks
Block 0 'entry'; Predecessors: Successors: 1
    a = 0
    !0x2ec69a0 = range(100)
    #iter_0x2ec6870 = !0x2ec69a0:__iter__()
    goto 1
Block 1; Predecessors: 0 Successors: 2 3
    if #iter_0x2ec6870:__hasnext__() goto 2 else goto 3
Block 2; Predecessors: 1 Successors: 4
    goto 4
Block 3; Predecessors: 1 Successors: 7
    goto 7
Block 4; Predecessors: 2 5 Successors: 5 6
    i = #iter_0x2ec6870:next()
    a+=1
    if #iter_0x2ec6870:__hasnext__() goto 5 else goto 6
Block 5; Predecessors: 4 Successors: 4
    goto 4
Block 6; Predecessors: 4 Successors: 7
    goto 7
Block 7; Predecessors: 3 6 Successors: 8
    goto 8
Block 8; Predecessors: 7 Successors:
    return 
processing types for block 0
processing types for block 1
processing types for block 2
processing types for block 3
processing types for block 4
processing types for block 7
processing types for block 5
processing types for block 6
processing types for block 8
processing types for block 4
processing types for block 7
processing types for block 8
processing opt block 0
processing opt block 1
processing opt block 2
processing opt block 3
processing opt block 4
processing opt block 7
processing opt block 5
processing opt block 6
processing opt block 8
generated IR:

define void @module_e0_0() {
opt_block0:
  %0 = call %"class.pyston::Box"* @boxInt(i64 0), !dbg !8
  call void @setattr(%"class.pyston::Box"* inttoptr (i64 79188515504 to %"class.pyston::Box"*), i8* inttoptr (i64 49066800 to i8*), %"class.pyston::Box"* %0), !dbg !8
  %1 = call %"class.pyston::Box"* @getGlobal(%"struct.pyston::BoxedModule"* inttoptr (i64 79188515504 to %"struct.pyston::BoxedModule"*), %"class.std::basic_string"* inttoptr (i64 49048152 to %"class.std::basic_string"*), i1 true), !dbg !9
  %2 = call %"class.pyston::Box"* @boxInt(i64 100), !dbg !9
  %3 = call %"class.pyston::Box"* @runtimeCall(%"class.pyston::Box"* %1, i64 1, %"class.pyston::Box"* %2), !dbg !9
  %4 = call %"class.pyston::Box"* @callattr(%"class.pyston::Box"* %3, %"class.std::basic_string"* inttoptr (i64 49048384 to %"class.std::basic_string"*), i1 true, i64 0), !dbg !10
  br label %opt_block1, !dbg !10

opt_block1:                                       ; preds = %opt_block0
  %5 = call %"class.pyston::Box"* @callattr(%"class.pyston::Box"* %4, %"class.std::basic_string"* inttoptr (i64 49048896 to %"class.std::basic_string"*), i1 true, i64 0), !dbg !10
  %6 = call i1 @nonzero(%"class.pyston::Box"* %5), !dbg !10
  br i1 %6, label %opt_block2, label %opt_block3, !dbg !10

opt_block2:                                       ; preds = %opt_block1
  br label %opt_block4

opt_block3:                                       ; preds = %opt_block1
  br label %opt_block7

opt_block4:                                       ; preds = %opt_block5, %opt_block2
  %"#iter_0x2ec6870" = phi %"class.pyston::Box"* [ %4, %opt_block2 ], [ %"#iter_0x2ec6870", %opt_block5 ]
  %7 = call %"class.pyston::Box"* @callattr(%"class.pyston::Box"* %"#iter_0x2ec6870", %"class.std::basic_string"* inttoptr (i64 49049088 to %"class.std::basic_string"*), i1 true, i64 0), !dbg !10
  call void @setattr(%"class.pyston::Box"* inttoptr (i64 79188515504 to %"class.pyston::Box"*), i8* inttoptr (i64 49088640 to i8*), %"class.pyston::Box"* %7), !dbg !10
  %8 = call %"class.pyston::Box"* @getGlobal(%"struct.pyston::BoxedModule"* inttoptr (i64 79188515504 to %"struct.pyston::BoxedModule"*), %"class.std::basic_string"* inttoptr (i64 49047848 to %"class.std::basic_string"*), i1 true), !dbg !11
  %9 = call %"class.pyston::Box"* @boxInt(i64 1), !dbg !11
  %10 = call %"class.pyston::Box"* @augassign(%"class.pyston::Box"* %8, %"class.pyston::Box"* %9, i32 69), !dbg !11
  call void @setattr(%"class.pyston::Box"* inttoptr (i64 79188515504 to %"class.pyston::Box"*), i8* inttoptr (i64 49066800 to i8*), %"class.pyston::Box"* %10), !dbg !11
  %11 = call %"class.pyston::Box"* @callattr(%"class.pyston::Box"* %"#iter_0x2ec6870", %"class.std::basic_string"* inttoptr (i64 49048896 to %"class.std::basic_string"*), i1 true, i64 0), !dbg !10
  %12 = call i1 @nonzero(%"class.pyston::Box"* %11), !dbg !10
  br i1 %12, label %opt_block5, label %opt_block6, !dbg !10

opt_block5:                                       ; preds = %opt_block4
  %13 = load i64* @edgecount
  %14 = add i64 %13, 1
  store i64 %14, i64* @edgecount
  %15 = icmp sgt i64 %14, 100
  br i1 %15, label %onramp, label %opt_block4, !prof !12

opt_block6:                                       ; preds = %opt_block4
  br label %opt_block7

opt_block7:                                       ; preds = %opt_block6, %opt_block3
  br label %opt_block8

opt_block8:                                       ; preds = %opt_block7
  ret void

onramp:                                           ; preds = %opt_block5
  %16 = call i8* @"pyston::compilePartialFunc(pyston::OSRExit*)"(i8* inttoptr (i64 49094656 to i8*))
  %17 = bitcast i8* %16 to void (%"class.pyston::Box"*)*
  call void %17(%"class.pyston::Box"* %"#iter_0x2ec6870")
  ret void
}

   1.5ms in compileFunction
Compiling...
Compiled function to (nil)
  1.5ms for _doCompile()
 1.6ms for compileModule()
compiled module.main to machine code; running:
 5us to interpret
finished running
6.9ms to run
In teardownRuntime
52us joinRuntime
Stats:
OSR exits: 0
getglobal_builtins: 1
interpreted_runs: 1
nopatch_getglobal: 101
num_compiles: 1
num_compiles_0_interpreted: 1
num_hidden_classes: 155
reopts: 0
rewriter_nopatch: 706
slowpath_binop: 100
slowpath_callattr: 202
slowpath_getglobal: 101
slowpath_nonzero: 0
slowpath_resolveclfunc: 303
slowpath_runtimecall: 1
slowpath_setattr: 201
us_compiling: 1536
us_compiling_0_interpreted: 1536
us_compiling_irgen: 1520
us_interpreting: 127
us_parsing: 5020
154us finishing up

実行結果は6.9msec。

まだ「試しに動く」という段階っぽいので、今後の発展が楽しみ！

GO言語のビルド時に"imported and not used"を出力させない方法

GO言語をビルドするときに、使っていないパッケージや変数があると「使ってませんよ！」と怒られる。まぁその怒られる内容に行数書いてあるからコメントアウトして対応すればいいのだけど、たくさんあったりするとめんどくさい。そこでコメントアウトする以外の解決策を紹介する。

package main

import (
	"fmt"
	"os"
)

func main() {
	s := "Happy Programming"
	println("Hello World")
}

fmtとosは使っていないので、下記の内容で怒られる。

$ go build main.go
# command-line-arguments
.\main.go:4: imported and not used: "fmt"
.\main.go:5: imported and not used: "os"

そこでこんな記述をソースにする。

package main

import (
	"fmt"
	"os"
)

func main() {
	s := "Happy Programming"
	println("Hello World")

	// ビルドエラー回避
	var _ = fmt.Print
	var _ = os.Args
}

これでパッケージに関しては"imported and not used"で怒られなくなる。しかし、こんどは変数sを使っていないためまた怒られる。

$ go build main.go
# command-line-arguments
.\main.go:9: s declared and not used

変数の場合も同様の対応をとれば怒られなくなる。

package main

import (
	"fmt"
	"os"
)

func main() {
	s := "Happy Programming"
	println("Hello World")

	// ビルドエラー回避
	var _ = fmt.Print
	var _ = os.Args
	var _ = s
}

【2014/07/04追記】
最近知ったのだが、下記の方法でも怒られなくなる。

package main

import (
	_ "fmt"
	_ "os"
)

func main() {
	println("Hello World")

	// fmtもosも使ってないが怒られない！
}

毎週少しずつ

下は以前ぼくが関西GPGPU勉強会で発表した資料。

http://www.slideshare.net/ksakiyama134/javagpgpu-aparapi

毎週Slideshareからメールが届き、今週はあなたのスライドがx回読まれましたよ的な案内がくる。で、毎週+40くらいこのスライドのTotalViewが上がっている。。。GPGPUerだった頃が懐かしい。

最近はGPGPUとか並列プログラミングやってへんなー。仕事も一段落ついたし、久しぶりにOpenCLでコード書いたりしてみようか。もうOpenCL2.0とか動いたりするのか？

ATI Catalystを強制的にアンインストールする方法(Linux Mint)

私はよくLinux Mintを使っているが、アップデートをかけるとGUIが上手く起動しなかったりすることがよくある。そんなときは強制的にCatalystをアンインストールしてOSを再起動、そしてドライバを再インストールしている。Linuxよくわかってないからこんな復旧作業しか思い浮かばないけどメモっておく。
※自己責任でお願いします

【私の環境】
Linux Mint 15 Cinnamon x64
Radeon HD7970

1)
CTRL + ALT + F1
とキーボード入力しCUI画面へ、そしてログインする。

2)

cd /usr/share/ati/
sudo ./amd-uninstall.sh --force

「--force」オプションを指定して強制アンインストールをする。オプションを指定しなくてもアンインストールできる場合がある。しかし私の環境ではほとんどできた試しがない。

3)

sudo reboot

として再起動する。

4)
GUIのログイン画面は立ち上がるので、ここでまた
CTRL + ALT + F1
でCUI画面へ。

5)

sudo ./amd-driver-installer-catalyst-13-4-x86.x86_64.run

で再インストールをする。

※私はよくこの復旧作業を行うので、ドライバーのインストーラを常に保存してある。すでに削除している方は、wgetなどでダウンロードしてください。

6)

sudo reboot

これで終わり。

BOLTの使ってみた感想（2013年8月)

いつの間にかBOLT 1.0 GA Releaseとかなってた。

ちょっと見た。

ほぼ何にもかわってなかった。

マジでBOLT盛り上がってない感じ、、、
やはりCUDAには勝てないのか。

とりあえずドキュメントちゃんとしてほしい。
あとせめてLinuxで動くようにしてくれ。

Haswell買ったからclinfo貼り付けとく

とりあえずCPUをi7-4770K(Haswell)にした。
以下、clinfoを実行した結果。

  Platform Name:                                 Intel(R) OpenCL
Number of devices:                               1
  Device Type:                                   CL_DEVICE_TYPE_CPU
  Device ID:                                     32902
  Max compute units:                             8
  Max work items dimensions:                     3
    Max work items[0]:                           1024
    Max work items[1]:                           1024
    Max work items[2]:                           1024
  Max work group size:                           1024
  Preferred vector width char:                   1
  Preferred vector width short:                  1
  Preferred vector width int:                    1
  Preferred vector width long:                   1
  Preferred vector width float:                  1
  Preferred vector width double:                 1
  Native vector width char:                      32
  Native vector width short:                     16
  Native vector width int:                       8
  Native vector width long:                      8
  Native vector width float:                     4
  Native vector width double:                    4
  Max clock frequency:                           3500Mhz
  Address bits:                                  64
  Max memory allocation:                         4278918144
  Image support:                                 Yes
  Max number of images read arguments:           480
  Max number of images write arguments:          480
  Max image 2D width:                            16384
  Max image 2D height:                           16384
  Max image 3D width:                            2048
  Max image 3D height:                           2048
  Max image 3D depth:                            2048
  Max samplers within kernel:                    480
  Max size of kernel argument:                   3840
  Alignment (bits) of base address:              1024
  Minimum alignment (bytes) for any datatype:    128
  Single precision floating point capability
    Denorms:                                     Yes
    Quiet NaNs:                                  Yes
    Round to nearest even:                       Yes
    Round to zero:                               No
    Round to +ve and infinity:                   No
    IEEE754-2008 fused multiply-add:             No
  Cache type:                                    Read/Write
  Cache line size:                               64
  Cache size:                                    262144
  Global memory size:                            17115672576
  Constant buffer size:                          131072
  Max number of constant args:                   480
  Local memory type:                             Global
  Local memory size:                             32768
  Kernel Preferred work group size multiple:     128
  Error correction support:                      0
  Unified memory for Host and Device:            1
  Profiling timer resolution:                    292
  Device endianess:                              Little
  Available:                                     Yes
  Compiler available:                            Yes
  Execution capabilities:
    Execute OpenCL kernels:                      Yes
    Execute native function:                     Yes
  Queue properties:
    Out-of-Order:                                Yes
    Profiling :                                  Yes
  Platform ID:                                   00000000005526F0
  Name:                                          Intel(R) Core(TM) i7-4770K CPU @ 3.50GHz
  Vendor:                                        Intel(R) Corporation
  Device OpenCL C version:                       OpenCL C 1.2
  Driver version:                                1.2
  Profile:                                       FULL_PROFILE
  Version:                                       OpenCL 1.2 (Build 63463)
  Extensions:                                    cl_khr_fp64 cl_khr_icd cl_khr_global_int32_base_ato
mics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extende
d_atomics cl_khr_byte_addressable_store cl_intel_printf cl_ext_device_fission cl_intel_exec_by_local
_thread cl_khr_gl_sharing cl_intel_dx9_media_sharing cl_khr_dx9_media_sharing cl_khr_d3d11_sharing

Bolt Preview Version

BoltはAMDが開発したC++のテンプレートライブラリです。STLに似たインターフェースで並列アルゴリズムを提供しています。現在はPreview版で、AMDのサイトからダウンロードできます。
OpenCLベースで開発されていますが、他ベンダーのデバイスと互換性があるわけではなく、AMD APP SDKで動作するデバイスのみが対象となります。

例えばsortをする場合、STLやThrustとほぼ同じように使用できます。

std::vector<int> a(100000);

/* STL */
std::sort(a.begin(), a.end());

/* Thrust */
thrust::sort(a.begin(), a.end());

/* Bolt */
bolt::cl::sort(a.begin(), a.end());

transform()などにオリジナルの関数オブジェクトを渡す場合、Boltが提供しているマクロを使用して、OpenCL Cのコードを生成します。例えばSAXPYをbolt::cl::transform()を使用して書く場合

#include <bolt/cl/transform.h>

BOLT_FUNCTOR(SaxpyFunctor,
struct SaxpyFunctor
{
  float _a;
  SaxpyFunctor(float a) : _a(a) {}
  float operator()(const float &xx, const float &yy)
  {
    return _a * xx + yy;
  }
};
);

int main()
{
  unsigned int count = 8192;
  std::vector<float> x(count);
  std::vector<float> y(count);
  std::vector<float> z(count);

  SaxpyFunctor s(2.f);

  bolt::cl::transform(x.begin(), x.end(), y.begin(), z.begin(), s);
}

このようにSAXPYをSTLちっくに書けます。計算量が少ないんでGPUでやるメリットはほとんどありませんが...

また、BoltはAMDの拡張機能であるC++ Static Kernel Languageを利用しているため、Templateも使えます。Templateを使ってcount_if()に渡すオブジェクトを定義する場合は

#include <bolt/cl/bolt.h>
#include <bolt/cl/count.h>

std::string stdStrCountIfRange = 
BOLT_CODE_STRING(
template <typename T>
struct CountIfRange
{
  CountIfRange(const T low, const T high) : _low(low), _high(high) {}
  bool operator()(const T &x)
  {
    return (x >= _low) && (x <= _high);
  }
private:
  T _low;
  T _high;
};
);

BOLT_CREATE_TYPENAME(CountIfRange<int>);

int main()
{
  int a[12] = {10, 5, 4, 3, 1, 7, 11, -2, 9, 3, 2, 10};
  size_t countIf = bolt::cl::count_if(a, a + 12, CountIfRange<int>(5, 10), stdStrCountIfRange);
}

BOLT_CODE_STRINGの中に定義し、それをstd::stringで受け取っています。
BOLT_CREATE_TYPENAMEで型を明示的に宣言します。（C++ Static Kernel Languageもこれがなければけっこう使えると思うの）

OpenCLをベースとしているので、デバイスを選ぶことも可能です。（おそらくデフォルトではAMDのGPUになっているはず。）
Boltにはbolt::cl::controlというものが定義されており、これにコマンドキューを渡すことで、並列アルゴリズムをそのコマンドキューのデバイスで実行することができる。

#include <bolt/cl/bolt.h>
#include <bolt/cl/reduce.h>

int main()
{
  std::vector<cl::Platform> platforms;
  cl::Platform::get(&platforms);

  {
    std::string platformName;
    platforms[1].getInfo(CL_PLATFORM_NAME, &platformName);
    std::cout << platformName << std::endl;
  }

  std::vector<cl::Device> devices;
  platforms[1].getDevices(CL_DEVICE_TYPE_GPU, &devices);

  {
    std::string deviceName;
    devices[0].getInfo(CL_DEVICE_NAME, &deviceName);
    std::cout << deviceName << std::endl;
  }

  cl::Context context(devices);
  cl::CommandQueue queue(context, devices[0]);

  bolt::cl::control ctrl;
  ctrl.commandQueue(queue);

  std::vector<int> a(1024, 1);

  int sum = bolt::cl::reduce(ctrl, a.begin(), a.end(), 0);

  std::cout << "Result: " << sum << std::endl;
}

第一引数にctrlを渡すことでデバイスを指定します。ただし、intelのCPUのキューを作って渡した場合は上手くいかなかった。

Radeon GPGPUerにとってはかなり便利なので、正式版リリースが待ち遠しいです。
さてさて、ThrustかBolt、どちらが先に互換性をクリアするか...両方ともしないかもしれない！？