J

2004 | 08 | 09 | 10 | 11 | 12 |
2005 | 01 | 02 | 03 | 04 | 05 | 06 | 07 | 08 | 09 | 10 | 11 | 12 |
2006 | 01 | 02 | 03 | 04 | 05 | 06 | 07 | 08 | 09 | 10 | 11 | 12 |
2007 | 01 | 02 | 03 | 04 | 05 | 06 | 12 |
2008 | 01 | 02 | 04 | 10 | 11 | 12 |
2009 | 01 | 02 | 03 | 04 | 05 | 06 | 07 | 08 | 09 | 10 | 11 | 12 |
2010 | 01 | 02 | 03 | 04 | 05 | 06 | 07 | 08 | 09 | 10 | 11 | 12 |
2011 | 01 | 02 | 03 | 04 | 05 | 06 | 07 | 08 | 09 | 10 | 11 | 12 |
2012 | 01 | 03 | 04 | 05 | 06 | 07 | 08 | 12 |
2013 | 01 | 02 | 03 | 05 | 06 | 07 | 08 | 09 | 10 | 11 |
2014 | 01 | 03 | 04 | 05 | 06 | 07 | 08 | 09 | 10 | 11 | 12 |
2015 | 02 | 03 | 04 | 05 | 06 | 07 | 10 | 11 | 12 |
2016 | 01 | 02 | 04 | 05 | 06 | 07 | 09 | 11 | 12 |
2017 | 01 | 02 | 03 | 05 | 06 | 07 | 08 | 09 |

ホーム

日記内の"morihyphen.hp.infoseek.co.jp"へのリンクは切れてます。必要な場合はお手数ですが int.main.jp へ書き換えをお願いします。

TODO: ファイル名確認を忘れないこと > 自分

twitter

 | 

2014-01-29

22:13

どうやったら組み込み業界の人々がweb業界の人々みたいに、業務内容をブログに書けるか、みないなことを考えていたのだが、web業界の人々はオープンソース(て何?)なフレームワークを使ってるし、トップレベルエンジニアは自社開発してるとこにいるから、まあ、大丈夫なんだという結論になってしまった。


業務内容に関することををNDAでガチガチに縛るのってどうなのかと思うね。

現代において、同業界内のライバルは真のライバルではない。本当のライバルは、別のもっと面白そうな業界である。

トラックバック - http://d.hatena.ne.jp/w_o/20140129

2014-01-28

プログラマが集中してる時に話しかけてあげよう 21:44

プログラマが集中してる時は邪魔すんな、みたいなのって、当然のように語られてるけど、あれ、正しくないこともよくあると思う。


まあ、真面目な話をすると、プログラマの障害の半分くらいはコミュニケーションロスから出てくるのだから、作業効率半分くらいになったとしても、話しかけんなオーラをアピールするのは良くない、とかある。


…が、まあそれはいいとして、寂しいから話しかけて欲しい、という人も、それなりにいると思う。

2ヶ月くらい前になんかの機会に8人ぐらいに聞いたところ、「集中してる時は話しかけんな」派は半分くらいだった。

多分、割合としては、

  • 話しかけんな派 : 50%
  • どっちでもいい派 : 25%
  • 寂しいから話しかけて欲しい派 : 25%

くらいなんではないですかね。(超適当)


僕とかは、人に仕事頼まれた直後とかは、承認欲求パワーによりそれなりにやる気あるけど、そこから2時間とか半日とか過ぎると、なんかどうでもよくなってきて、「自分が生きる理由とは…」とか考えだすから、あんま良くないんだよね。

集中してる時に作業中断されると嫌という気持ちもわかるけど、業務時間の95割くらいは集中してないから、まあ実際悪影響は無いと言える。

いやでもさすがに二時間に一回進捗確認されたら嫌かな…(めんどくさい系人間)


同じような人それなりにいると思う。が、そういう人がちゃんとアッピールしてないから、「プログラマ=集中してる時に話しかけたらいけない」とかみたいな常識ができてしまってよくないので、ちゃんと寂しいから構ってほしいとか主張すべきである。寂しくて死んでしまう人が出てからでは遅い。

トラックバック - http://d.hatena.ne.jp/w_o/20140128

2014-01-27

cuda-memcheck 21:39

cuda-memcheck とかいう CUDA で valgrind みたいなことができるツールがある。


↓こんなードを書いたとする。

#include <stdio.h>

__global__ void f(int *ptr)
{
    int tid = threadIdx.x;
    ptr[tid] = tid;
}

int
main()
{
    int *dmem;
    int hmem[32];

    memset(hmem, 0, sizeof(hmem));
    cudaMalloc(&dmem, 32);
    cudaMemcpy(dmem, hmem, 32, cudaMemcpyHostToDevice);
    f<<<1,32>>>(dmem);
    cudaMemcpy(hmem, dmem, 32, cudaMemcpyDeviceToHost);

    printf("%d\n", hmem[4]);
}

いちおう動いているように見えるが、間違っている。(どこが間違ってるかすぐわかる人偉い)


ビルドするときに、

$ nvcc -G -g hoge.cu

とかする(必須ではない)


んで、

$ cuda-memcheck ./a.out

とかやると

========= Invalid __global__ write of size 4
=========     at 0x000000d0 in /home/w0/test/cuda/memcheck.cu:6:f(int*)
=========     by thread (31,0,0) in block (0,0,0)
=========     Address 0x70014007c is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/lib64/libcuda.so (cuLaunchKernel + 0x331) [0xcd5d1]
=========     Host Frame:./a.out [0x1b698]
=========     Host Frame:./a.out [0x3ab63]
=========     Host Frame:./a.out [0x2937]
=========     Host Frame:./a.out [0x280d]
=========     Host Frame:./a.out [0x2827]
=========     Host Frame:./a.out [0x2750]
=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x21b75]
=========     Host Frame:./a.out [0x24e9]

メモリアクセスエラーしてるのがわかる。



あと cuda-gdb から set cuda memcheck しても同じようなことをデバッガからできる、とか書いてある。

http://docs.nvidia.com/cuda/cuda-memcheck/index.html

インストゥルメンテーション 21:50

全プロセッサメーカーはインストゥルメンテーションツールを用意すべき。

世界のIntelが公開してる pin をご覧ください http://software.intel.com/en-us/articles/pin-a-dynamic-binary-instrumentation-toolIntelがつくったvalgrindみたいなツール、Intel Inspectorはこれの上に実装されている。


CUDAのインストゥルメンテーションライブラリは公開されてないが、プロファイラや上のmemcheckの挙動を見るに、SASSレベルでのインストゥルメンテーションが存在すると思われる。


何故これが重要なのかというと、これがあると、機械語しか存在しないプログラムがどういう挙動をしてるか、というのが調べられるんだよね。

キャッシュサイズとか分岐予測変えたときのシミュレーションとかもできる。


つまり、こういうのを用意してるプロセッサメーカーは、ちゃんと実際のプログラム動かしてデータ集めてプロセッサ開発してるということである。

用意してないプロセッサメーカーは、脳内シミュレーションだけでアーキテクチャを決めてる可能性がある。

トラックバック - http://d.hatena.ne.jp/w_o/20140127

2014-01-26

自分に対してダメとか言ったらいけない 08:23

まあ最近のインターネットを活用した鬱研究によると、自分のこと悪く言うのが一番良くないと言われつつあった。


しかし、この点エンジにゃーは若干不利だよな。エンジにゃーは、それなりに、ものごとを悪く考えるのが仕事として有利な場面があって、日頃から物事を悪いほうに考える訓練をしてると言える。

もちろん、仕事で悪く考えるのと、日常のものごとを悪く考えるのは一緒ではないが、人間の精神はそんなにうまく切り離せるものではないので、こうしてエンジにゃーな人々は鬱っていくのだった。


まあ仕事より健全な精神のほうが大事なので、仕事でも悪く考えないようにするのが大事だと思う。

ファイル開けなかったらどうしよう
大丈夫。開けるよ。
ユーザーが操作まちがえたらどうしよう
まちがえるやつが悪いよ。僕は悪くないよ
SQL文字列が送られてきたらどうしよう
その人、SQL送ってデータベース見たかったんだよ。お客の要求に応えたんだよ。
スタックオーバーフローしたらどうしよう
任意のx86マシン語が実行できるwebサービスだよ。
こんなことやっててクビになったらどうしよう
エンジにゃーとか鬱になるからやめたほうがいいよ。クビになってよかったよ。
こんな文章書いてたらダメだよ
ダメだよ
トラックバック - http://d.hatena.ne.jp/w_o/20140126

2014-01-25

Kaveri 18:36

はかっていただきました。ありがとうございます。

https://twitter.com/neutral_tao/status/426943905012006912

https://twitter.com/neutral_tao/status/426945868516687873

enqueue kernel latency 80usec…

手元のマシンだとNVIDIAのは40usec、Intel HD Graphicsが60usecぐらい。

あとこれ見て思ったが、pref vector width char:4 ってそういう命令あったっけ…まあそのうち

トラックバック - http://d.hatena.ne.jp/w_o/20140125

2014-01-24

Silvermont 09:34

f:id:w_o:20140124092651p:image

madはなんか動かすと止まるので別にはからないといけなくて別にはかったら mad1 throughput 77GFlops だった

http://sdrv.ms/M0O9Q2

https://github.com/tanakamura/clminibench

前と比べて

  • madとfma間違ってたの修正
  • radeon用に調整
  • たくさんenqueueした場合の時間はかるように(kernel latency x16)

だれかKaveri持ってたらenqueue kernel latencyだけでもいいので教えてほしい。


CPUの計測はなんか思いついたらやる。(多分やらない)

追記 08:53

SilvermontはCPUコアだけをさすから、こういうGPUを含む文脈ではBay Trailと書くべきでは?だからダメなんだよお前は。

トラックバック - http://d.hatena.ne.jp/w_o/20140124

2014-01-22

clSetEventCallback()を正しく使いたい(書きかけ) 00:17

http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetEventCallback.html

これのコールバック(pfn_event_notify)正しく実装するの不可能では?という気がした



経緯 :

https://twitter.com/tanakmura/status/399251693092605952

https://twitter.com/tanakmura/status/399252471903567872

https://twitter.com/tanakmura/status/399252550756487168

https://twitter.com/tanakmura/status/399253075396788224

khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clSetEventCallback.html … これ正しく使うの不可能な気がする。

どうやって非同期で呼ばれるかが実装依存になってる。考えられるのは、

(1) Windows User Mode APCみたいなの : ブロックするシステムコールで呼ばれる

(2) POSIX SIGNALみたいなの : いつでも非同期で飛んでくる

(3) 別スレッドから呼ばれる

だが、

(1) の場合ビジーループでデバイスの完了を待つみたいな処理を入れるとデッドロックする

(2) の場合Mutex入れるとデッドロックする

(3) Mutex入れないとレースコンディション発生する

で、 (2) と (3) を両方起こさないようになんか処理するの無理だと思う。


複数platform の OpenCL を動かしたいとする(NVIDIA GPUIntel CPUみたいな)。

この時、複数デバイスを同期するためのスケジューラみたいなのが必要になる。それほど複雑ではないので、キューが繋がってて、それでやりとりする実装にしよう。


(あとで図を入れる)


僕は、ちょっとイベント処理するためだけに専用スレッド立ち上げるとか好きではないので、これをclSetEventCallback()を使ってなんとかしたいと考えた。

clSetEventCallbackは、eventとそのevent発生時に呼ばれる関数を関連付けるAPI

上の図でいう、notify complete の処理は、これで実装できるわけだ。


ところが、clSetEventCallbackの注意として、コールバックは非同期に呼ばれるかもしれないから、スレッドセーフにしろ、と書かれてあるのだけど、

This callback function may be called asynchronously by the OpenCL implementation. It is the application's responsibility to ensure that the callback function is thread-safe.

これ、非同期の呼びかたの実装方法が決まってないから、正しくスレッドセーフに実装するの無理だと思った、というのが言いたいことだった。


ぱっと思い付く非同期実装としては、シグナルとスレッドがあるが(上のtweetでは、Windows User Mode APCとか書いてるが、忘れてもいいので以後特に言及しない)、シグナルから呼ばれても大丈夫 && スレッドセーフにするの不可能だから、


(…と、ここまで書いてそもそもマルチスレッド関係無く Async-signal-safe 実装するの相当難易度高いのだから、clSetEventCallbackは別スレッドから呼ばれることを明記すべきという問題な気がしてきたな。もうちょっと考える)

いや 00:25

thread-safe にしろと明記してあるが、 Async-signal-safe にしろと明記されてないから、全然問題無いということなのかもしれない。

まあ、明記されてないから大丈夫というのも気持ち悪いが、Async-signal-safeを実現するのは、相当難易度高いので、明記されてないAPIは、unsafeだというのが世間(?)の常識(?)だから、まあ大丈夫という気がしてきたな…

(なので続き書かないかもしれない)

関東GPGPU勉強会#3 01:17

そう言えば書いてなかった。

http://int.main.jp/txt/ai-isa-jit.html

コード : https://github.com/tanakamura/ai-isa-jit/tree/gpgpu3 (多分僕以外が使えるものではない)

GCNだと、理論値出すには最低CUあたり512 item立ち上げないといけない理由はあとで読める文章にしたいな…


reveal.js 使ったのは、これの直近で↓見た影響

http://ezoeryou.github.io/kabukiza-tech2-slide/index.html

トラックバック - http://d.hatena.ne.jp/w_o/20140122
 |