Google Homeでのやり取りの履歴を見る

Google Homeに話しかけた内容、また、Google Homeの返答の履歴を見る方法です。

  1. こちらのページを開く
  2. 検索の下の"日付とサービスでフィルタ"を選択
  3. アシスタントにチェック
  4. 検索の虫眼鏡ボタンを押すと、Google Homeとのやり取りの履歴が出てきます

スマートフォン等で利用したGoogleアシスタントの履歴も含まれます。

f:id:dfukunaga:20171029224454p:plain

f:id:dfukunaga:20171029224513p:plain

こんな感じで履歴が表示されます。自分の声も聞くことができます。

f:id:dfukunaga:20171029224445p:plain

【CUDA】カーネル内での動的メモリ確保

バイス上のメモリは、ホストからのcudaMalloc()で確保するのが通常だが、 カーネル内で動的にメモリ確保することもできる*1カーネル内での動的メモリ確保には、C/C++と同じようにmalloc()/free()new/deleteが使える。

__global__
void testKernel() {
  size_t size = 1024;

  // Dynamic allocation with malloc/free
  char *a = (char*)malloc(size);
  free(a);

  // Dynamic allocation with new/delete
  char *b = new char[size];
  delete[] b;
}

int main() {
  testKernel<<<1, 1>>>();
  cudaDeviceSynchronize();
  return 0;
}

ただし、ホストからのcudaMalloc()カーネル内でのmalloc()はメモリ確保方法が異なる。 ホストからのcudaMalloc()はデバイスメモリの未使用の領域にメモリを確保するのに対して、 カーネル内でのmalloc()はあらかじめ用意されたデバイス上のヒープにメモリを確保する。 そのため、ホストからのcudaMalloc()で確保したメモリはカーネル内では解放できないし、 カーネル内で確保したメモリはホストからは解放できない。 以下のコードはエラーになるパターンを表している。

__global__
void mallocInKernel(char **ptr, size_t size) {
  *ptr = (char*)malloc(size);
}

__global__
void freeInKernel(char *ptr) {
  free(ptr);
}

int main() {
  size_t size = 1024;

  // ホスト上でcudaMalloc()したメモリをカーネル内でfree()
  char *a;
  cudaMalloc(&a, size);
  freeInKernel<<<1, 1>>>(a);  // Error!!
  cudaDeviceSynchronize();

  // カーネル内でmalloc()したメモリをホストからcudaFree()
  char *hb, **db;
  cudaMalloc(&db, sizeof(char*));
  mallocInKernel<<<1, 1>>>(db, size);
  cudaMemcpy(&hb, db, sizeof(char*), cudaMemcpyDeviceToHost);
  cudaFree(hb);  // Error!!
  cudaDeviceSynchronize();

  return 0;
}

バイスのヒープはデフォルトで8MBしか確保されない。 つまり、カーネル内のmalloc()で8MB以上のメモリを確保するとエラーになる。 この設定を変えるにはcudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)を使えばよい。 例えば、以下のように書くと1GBまでカーネル内でmalloc()できる。

size_t size = 1024 * 1024 * 1024;
cudaDeviceSetLimit(cudaLimitMallocHeapSize, size);

この設定は全てのカーネル実行が行われる前に書かれている必要があり、途中で設定を変えることはできない。 また、現在のヒープサイズの設定がどうなっているか確認するためには cudaDeviceGetLimit(size_t* size, cudaLimitMallocHeapSize)を使えばよい。

size_t size;
cudaDeviceGetLimit(&size, cudaLimitMallocHeapSize);
printf("Heap Size=%ld\n", size);

Compute Capability 3.5以降ではカーネル内でcudaMalloc()/cudaFree()を使うことができるが、 これはカーネル内でのmalloc()/free()と同じ効果をもたらす

バイスのヒープサイズを無駄に拡げないためにも、できるだけホストからのcudaMalloc()で メモリを確保したいが、場合によってはカーネル内での動的メモリ確保が有効なケースもあるだろう。

*1:Compute Capability 2.x以降でのみサポート

NVIDIA Jetson TX2を買ったのでセットアップ

届きました!

f:id:dfukunaga:20170902213325p:plain

意外と大きいです。一辺20cmくらいでしょうか。 ディスプレイ(HDMI)、キーボード、マウス、電源を接続します。USBポートは一基しかないので、キーボードとマウスを両方接続する場合は USBハブを使うか、付属のMicroUSB-USB変換ケーブルを使います。

f:id:dfukunaga:20170902213757p:plain

電源ボタンを押すとUbuntu 16.04のコンソールが開きます。 まずは以下のコマンドでドライバをインストールします。 (sudo時のパスワードはnvidia

$ cd NVIDIA-INSTALLER
$ sudo ./installer.sh
$ sudo reboot

再起動するとGUIが立ち上がります。 Wi-Fiもしくは有線でネットワークの設定を行っておきます。 (後で使うので、割り振られたIPアドレスを確認しておきます。)

ここから、JetPackのインストールです。

JetPackとは、CUDAやcuDNNを始め、Jetsonを使うのに必要そうなソフトウェアをまとめたものっぽいです。 筆者はMacbook Proを使っているので、VirtualBox上のUbuntu 16.04をホストとしてJetPackをインストールすることにしました。

こちらからUbuntu 16.04のイメージをダウンロードして、 VirtualBoxVMを作ります。JetPackのインストールに最低10GBのディスクスペースが必要とのことなのですが、 ディスクサイズは余裕を持って30GBに設定しました。

VMを起動してこちらからJetPackをダウンロードします。 NVIDIA Developer Programへの登録が必要になります。ダウンロードしたら以下のコマンドで実行。

$ chmod +x JetPack-L4T-3.1-linux-x64.run
$ ./JetPack-L4T-3.1-linux-x64.run

途中の"Component Manager"画面でFlash OS Image to TargetのActionをno actionに変更します。 各種パッケージをインストールした後、 “Device Information - Jetson TX2"画面でJetsonのIPアドレスとユーザ名、パスワード(共にnvidia)を入力します。

新しいウィンドウが開き、Jetson側にいろいろインストールされていき、 Installation of target components finished, close this window to continue.と表示されたら終わりです。 /home/nvidia/NVIDIA_CUDA-8.0_SamplesにCUDAのサンプルが入っているので、 これがうまく動けば少なくともCUDAはちゃんと入っているはずです。

ちなみに、JetsonではNVIDIAGPUでおなじみのnvidia-smiコマンドは対応していないみたいです。 代わりにホームディレクトリにあるtegrastatsGPUの状態を確認します。

nvidia@tegra-ubuntu:~$ sudo ./tegrastats 
RAM 1103/7851MB (lfb 1501x4MB) cpu [0%@1728,off,off,0%@1728,0%@1727,0%@1728] EMC 0%@1600 APE 150 GR3D 0%@114
RAM 1103/7851MB (lfb 1501x4MB) cpu [0%@345,off,off,1%@345,0%@346,5%@345] EMC 8%@40 APE 150 GR3D 0%@114
RAM 1103/7851MB (lfb 1501x4MB) cpu [1%@345,off,off,1%@345,0%@345,5%@345] EMC 8%@40 APE 150 GR3D 0%@114
RAM 1103/7851MB (lfb 1501x4MB) cpu [1%@345,off,off,2%@345,0%@345,4%@339] EMC 8%@40 APE 150 GR3D 0%@114

GR3Dの右がGPU使用率だそうです。GPUメモリはシステムと共有らしいので一番左の列を見ればOKです。

ZeroMQ: 複数のEndpointに接続する

f:id:dfukunaga:20170419234028p:plain

ちゃんとドキュメントを読めという話かもしれないが、ZeroMQ のソケットが 複数のEndpointにbind/connectできることを最近知った。そこで、複数のEndpointに対してのbind/connectを 使ったメッセージングのパターンをいくつか考えてみた。

1. ローカル/リモート共用ソケット

f:id:dfukunaga:20170418233138p:plain

同じプロセス内でのソケット間通信をするなら、inprocを使ったほうが ローカルループバックに対するtcpを使うよりも圧倒的に速い。このパターンでは、 一つのソケットに同一プロセスからの接続用のアドレス(ローカル: inproc)と、 別マシンからの接続用のアドレス(リモート: tcp)の二つをbindし、接続するソケットの位置 によってconnectするアドレスを使い分ける。これにより、ローカルのソケット間通信はより高速に行われ、また、 受信する側では、ローカルから来たメッセージとリモートから来たメッセージを区別することなく処理を行うことができる。

2. ブロードキャストクラスタ

f:id:dfukunaga:20170418233145p:plain

各ノードでZMQ_PUBZMQ_SUBを一つずつ用意します。 ZMQ_PUBはbindしてZMQ_SUBは各ノードのZMQ_PUBにconnectすると(逆でもたぶんOK)、 ZMQ_PUBにsendすることで他の全ノードにメッセージが送信されるようになります。 複数のソケットを用意すれば同じことが実現できますが、一つのソケットで複数のEndpointにconnectした方が、 recv/pollするのが一つのソケットで済むので楽になるかと思います。 このパターンを実装するときはZMQ_SUBSCRIBEを設定するのを忘れずに。

3. P2Pクラスタ

f:id:dfukunaga:20170418233158p:plain

bindとconnectを併用することもできます。このパターンでは、各ノードにZMQ_ROUTERを一つ用意し、 一つのアドレスに対してbindするのと同時に、他のノードのZMQ_ROUTERにconnectします。 このように繋ぐことで、ZMQ_ROUTERのルーティング機能を使えばクラスタ内のどのノードへも メッセージを送信することができます。ZMQ_ROUTER同士を繋ぐときはidentityの設定が厄介ですが、 bindする前にZMQ_IDENTITYを設定し、connectする際にZMQ_CONNECT_RIDを設定しておけば問題ないでしょう。

以上、複数のEndpointへのbind/connectを用いたメッセージングパターンでした。

通信時のソースアドレスを取得する

f:id:dfukunaga:20170417005216p:plain

Linuxです。

なんて言えばよいのかわからなかったのですが、ip route get xxx.xxx.xxx.xxxと打ったときに出てくる ソース(送信元)アドレスをC/C++内で取得したかった。すぐに思いついた方法は、

  • popenを使ってip route getを実行し、標準出力を取得
  • 一時的にソケットを作って送信元IPアドレスを取得

です。以下のようなサンプルコードを書きました。

get route source address

どちらも同じ出力が得られますが、手元の環境で試したところ、 1000回実行するのにかかる時間がpopenを使った方法では約1秒、 ソケットを使った方法では約0.005秒だった。(コードの改善点はあるかもしれませんが。)

他にもっと良い方法があれば教えてください!

ソケットに結び付けられたアドレスを取得する

f:id:dfukunaga:20170412233423p:plain

Linuxです。

netstatしたときのLocal AddressとForeign AddressにあたるものをC/C++で取得したかった。 Local Addressはgetsockname()、Foreign Addressはgetpeername()を使えば取得できるらしい。 C++でサンプルを書くと以下のような感じ。

gist6c0628fe0ad9790fe7a1fbc5cba9b13c

vagrant@vagrant-ubuntu-trusty-64:~$ ./a.out 192.168.0.2 2551
Local Address:   192.168.0.1:46402
Foreign Address: 192.168.0.2:2551

ちゃんと取得できてますね!(↑のサンプルコードを使うときは対応するアドレスをlistenしてください)

ZeroMQのheartbeatを試してみた

ZeroMQのversion 4.2からheartbeat機能が追加されました。 元々コネクションの切断は検知してzmq_socket_monitor(3)で通知してくれたのですが、 TCPエラーを吐かずにコネクションが死んだり、Out of Memory等でプロセスが止まったりした場合は検知できないので heartbeat機能が追加されたようです。(参考:ZeroMQのRFC

heartbeatに関するソケットのオプションは以下の3つです。

  • ZMQ_HEARTBEAT_IVL
  • ZMQ_HEARTBEAT_TIMEOUT
  • ZMQ_HEARTBEAT_TTL

ZMQ_HEARTBEAT_IVLミリ秒の間隔でpingを送り、ZMQ_HEARTBEAT_TIMEOUTミリ秒経っても pongが返ってこなかったらコネクションが死んだとみなすという感じです。また、pongを返す側では ZMQ_HEARTBEAT_TTLミリ秒経っても次のpingが来なかったらコネクションが死んだとみなします。 ZMQ_HEARTBEAT_IVLを設定することで初めてheartbeatが動作します。

試しに以下のような設定をしてみます。基本的にどちらか片方(connect側もしくはbind側)のソケットにだけ設定しておけばOKだと思います。

int interval = 100, timeout = 1000, ttl = 1000;
zmq_setsockopt(sock, ZMQ_HEARTBEAT_IVL, &interval, sizeof(int));
zmq_setsockopt(sock, ZMQ_HEARTBEAT_TIMEOUT, &timeout, sizeof(int));
zmq_setsockopt(sock, ZMQ_HEARTBEAT_TTL, &ttl, sizeof(int));
zmq_setsockopt(sock, ZMQ_HANDSHAKE_IVL, &timeout, sizeof(int));

ZMQ_HEARTBEAT_TTLZMQ_HEARTBEAT_IVLよりも十分大きい値にしておく必要があります。 そうしないと接続が頻繁に切れてしまいます。ZMQ_HEARTBEAT_TIMEOUTも短すぎると 偽陽性が増えてしまうので注意です。

また、ZMQ_HANDSHAKE_IVLを設定しています。ZeroMQではソケットが接続した後に ソケットの設定情報を交換する(これをhandshakeと呼ぶ)そうなのですが、止まっているソケットに再接続した際に ZMQ_HEARTBEAT_TIMEOUTよりもZMQ_HANDSHAKE_IVLが優先されるため、両方共同じ値にしておきます。 ZMQ_HANDSHAKE_IVLのデフォルト値は30秒なので、これを設定しないと ZMQ_HEARTBEAT_TIMEOUTをどれだけ短くしても再接続時に30秒待ってしまいます。

実際にこんな感じのコードを書いて動かします。 bind/connectが完了した時点で、connect側のプロセスにSIGTSTPを送るとbind側のzmq_socket_monitor(3)から ZMQ_EVENT_DISCONNECTEDが通知されます。逆に、bind側のプロセスにSIGTSTPを送ると、connect側の zmq_socket_monitor(3)からZMQ_EVENT_DISCONNECTEDが通知された後に、再接続を試みます。 bind側のlistenが残っているからか、再接続は成功しますが、handshakeが失敗して再びZMQ_EVENT_DISCONNECTEDが 通知されるというような挙動になります。

また、heartbeatを有効にしても前回の調査同様、unbind/disconnect時には 相手側のソケットには何も通知されません。disconnectしても再接続しようとするし、unbind/disconnectとは何なのだろう…。 とりあえずソケットを切断するときは必ずcloseした方が良さそうです。

※ZeroMQ v4.2.1を使用