JetsonのGPU/ファン/温度設定

いつも忘れてしまうのでメモ。

  • 最大パフォーマンスにしたいとき
$ sudo ./jetson_clocks.sh
  • 今の設定を保存したいとき
$ sudo ./jetson_clocks.sh --store
  • 元の設定に戻したいとき
$ sudo ./jetson_clocks.sh --restore
  • 今の設定を見たいとき
$ sudo ./jetson_clocks.sh --show
  • ファン速度を変えたいとき
$ echo [0-255] | sudo tee /sys/kernel/debug/tegra_fan/target_pwm
  • 温度を確認したいとき
$ cat /sys/devices/virtual/thermal/thermal_zone*/temp

google::dense_hash_mapを使う

google::dense_hash_map?

使い方

  • ヘッダオンリーなのでソースをダウンロードして#includeするだけ
  • 使う前にset_empty_key(key)を実行しなければいけない ←これ重要
    • ここで設定したキーは、ハッシュマップに格納できなくなる
    • これがstd::unrodered_mapと比べたときの欠点
  • また、要素を削除する前にset_deleted_key(key)を実行する必要あり
    • 要素を削除しない(挿入するだけ)場合は不要
    • set_empty_key(key)とは別のキーを設定しなければならない
  • その他は基本的にstd::unordered_mapと同じ

サンプルコード

#include <iostream>
#include <string>

#include "google/dense_hash_map"

int main() {
  google::dense_hash_map<std::string, std::string> smap;
  smap.set_empty_key("Mori");
  smap.insert({"Kimura", "Takuya"});
  smap.insert({"Kusanagi", "Tsuyoshi"});
  smap.insert({"Katori", "Shingo"});
  std::cout << smap["Kusanagi"] << std::endl;
  return 0;
}

注意点

  • バケット数の調整はreserve()ではなくresize()
  • clear()バケット数もリセットしてしまう
    • STLclear()と同じ挙動をするのはclear_no_resize()
    • clear_no_resize()するときはset_deleted_key(key)が必要

パフォーマンス

【CUDA】Warp Sum & Warp Scan

Warp Sum

Warpの持つ値の合計を計算する。 例:[1, 2, 3, 4] -> [10, 10, 10, 10]

for (int i = 1; i < warpSize; i *= 2)
  value += __shfl_xor(value, i);

Warp Scan

Warpの持つ値の累積を計算する。 例:[1, 2, 3, 4] -> [1, 3, 6, 10]

for (int i = 1; i < warpSize; i *= 2) {
  int n = __shfl_up(value, i);
  if (laneId >= i) value += n;
}

サンプルコード

CUDA Warp's Sum and Scan

2017/12/11追記

warpSizeはどうやらコンパイル時定数ではないみたいです。 そのため、サンプルコードのように書くとループ展開が行われない可能性があります。 パフォーマンスを気にするのであれば warpSize32 とした方が良さそうです。

基数を取得する

最下位ビット/バイトからk番目の基数を取得する関数です。

  • radix1bit → 2進数表現で右からk番目の値を取り出す (0 or 1)
  • radix4bit → 16進数表現で右からk番目の値を取り出す (0 to 15)
  • radix8bit → 1バイト区切りで右からk番目の値を取り出す (0 to 255)

Select k-th (from LSB) radix of the number.

CUDAのハッシュテーブル CUDPP編

インストール

まずCUDPPのインストール手順です。事前にcmakeのインストールが必要です。

$ git clone -b 2.2 https://github.com/cudpp/cudpp.git
$ cd cudpp
$ git submodule init
$ git submodule update
$ mkdir build && cd build
$ cmake ..
$ make
$ sudo make install

自分の環境だとgit submodule updateに失敗してしまったので、 以下のコマンドを追加で実行しました。

$ cd ext/cub
$ git reset --hard HEAD

更に、ビルド時にcudpp_config.hが無いと言われてしまったので、以下のコマンドも。

$ sudo cp include/cudpp_config.h /usr/local/include

サンプルコード

CUDPPは全体的にサンプルコードが不足しているような気がするので、 まず普通にハッシュテーブルを使うためのサンプルコードを載せておきます。

Example of CUDPP Hash Table

上記のサンプルではエラーチェックをしていません。 CUDPPのエラーチェックをする場合は以下のようになります。

CUDPPResult result = cudppCreate(&cudpp);
if (result != CUDPP_SUCCESS) {
  fprintf(stderr, "Error in cudppCreate.\n");
  return 1;
}

コンパイルは以下のようなコマンドで行います。

$ nvcc cudpp_hash_sample.cc -lcudpp -lcudpp_hash

ハッシュテーブル

CUDPPには以下の3種類のハッシュテーブルがあります。

  • CUDPP_BASIC_HASH_TABLE … 通常のキー・値のハッシュテーブル
  • CUDPP_COMPACTING_HASH_TABLE … キーに対する一意なIDを生成するハッシュテーブル
  • CUDPP_MULTIVALUE_HASH_TABLE … 一つのキーに対して複数の値を格納できるハッシュテーブル

どのハッシュテーブルを使うかは、CUDPPHashTableConfigで設定します。

CUDPPHashTableConfig config;
config.type = CUDPP_BASIC_HASH_TABLE;

ハッシュテーブルへのデータの挿入および取得は ホストからのみ可能で、 複数のキー・値をまとめて挿入、まとめて取得します。 データの挿入、取得の際に与えるキー・値の配列は バイス上に確保 されている必要があります。

// d_keys → 挿入するキーの配列、d_vals → 挿入する値の配列、N → 挿入するキー・値の数
cudppHashInsert(hash_table_handle, d_keys, d_vals, N);

// d_input → 値を取得するキーの配列、d_output → 取得した値の書き込み先、N → 取得する値の数
cudppHashRetrieve(hash_table_handle, d_input, d_output, N);

キーおよび値には32ビットのデータしか使えない ので注意してください。ここからは各ハッシュテーブルの使い方を見ていきます。

CUDPP_BASIC_HASH_TABLE

一つのキーに対して一つの値を格納できるハッシュテーブルです。 キーは重複してはいけません。また、格納していないキーの値を取得しようとすると CUDPP_HASH_KEY_NOT_FOUNDが返ってきます。

// keys = {2, 4, 0, 3, 1}, vals = {0, 1, 2, 3, 4};
cudppHashInsert(hash_table_handle, keys, vals, 5);

// input = {0, 1, 2, 3, 4, 0, 2, 4} --> output = {2, 4, 0, 3, 1, 2, 0, 1}
cudppHashRetrieve(hash_table_handle, input, output, 8);

CUDPP_COMPACTING_HASH_TABLE

キーに対する一意なIDを生成するハッシュテーブルです。 例えば{16, 34, 81, 16, 52}という入力に対して、{16=>0, 34=>1, 81=>2, 52=>3}のような ハッシュテーブルを作ることができます。キーは重複していても問題ありません。 また、値の配列は不要なのでNULLを指定すればOKです。

// keys = {16, 34, 81, 16, 52}
cudppHashInsert(hash_table_handle, keys, NULL, 5);

// input = {16, 34, 81, 16, 52} --> output = {2, 3, 1, 2, 0}
cudppHashRetrieve(hash_table_handle, input, output, 5);

CUDPP_MULTIVALUE_HASH_TABLE

一つのキーに対して複数の値を格納できるハッシュテーブルです。 キーが重複していた場合に、同じキーに対して複数の値が設定されます。

// keys = {3, 1, 3, 3, 2, 1}, vals = {0, 1, 2, 3, 4, 5};
cudppHashInsert(hash_table_handle, keys, vals, 6);

このハッシュテーブルにはcudppMultivalueHashGetAllValuescudppMultivalueHashGetValuesSizeという特別な関数が用意されています。 cudppMultivalueHashGetAllValuesでは全ての値が格納された配列を返します。 値は以下の図のように、各キー毎にまとめて並べられています。

|     key1    | key2 |        key3        | ...
| val1 | val2 | val1 | val1 | val2 | val3 | ...

cudppMultivalueHashGetValuesSizeでは、その配列のサイズを取得できます。

// hash_table = {1 => {1, 5}, 2 => {4}, 3 => {0, 2, 3}}

unsigned int size;
cudppMultivalueHashGetValuesSize(hash_table_handle, &size);
// size = 6

unsigned int *d_all_vals;
cudppMultivalueHashGetAllValues(hash_table_handle, &d_all_vals);
// d_all_vals = {1, 5, 4, 0, 2, 3}

このハッシュテーブルでは他の二つと異なり、cudppHashRetrieveで64ビットのデータを返します。 前半32ビットには先ほど取得した配列上のインデックス、後半32ビットには値の数が表されています。

// hash_table = {1 => {1, 5}, 2 => {4}, 3 => {0, 2, 3}}

uint2 *d_output;
cudaMalloc((void**)&d_output, sizeof(uint2) * 3);
cudaMemset(d_output, 0, sizeof(uint2) * 3);

// input = {1, 2, 3} --> output = {{0, 2}, {2, 1}, {3, 3}}
cudppHashRetrieve(hash_table_handle, d_input, d_output, 3);

// 値を表示
uint2 h_output[3];
cudaMemcpy(h_output, d_output, sizeof(int) * 3, cudaMemcpyDeviceToHost);
for (int i = 0; i < 3; ++i) {
  for (int j = 0; j < output[i].y; ++j) {
    printf("key = %d, val = %d\n", input[i], all_vals[output[i].x + j]);
  }
}

まとめ

CUDPPの3種類のハッシュテーブルの使い方を見ていきました。 少し使い方にクセがあるかもしれませんが、貴重なハッシュテーブルのCUDA実装なので ぜひ試してみてください。

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以降でのみサポート