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?
- GoogleのOSS GitHub - sparsehash/sparsehash: C++ associative containers
- その中に含まれているC++のハッシュテーブル
- STLと同じインターフェイス
- メモリ使用量は多いが、高速(らしい)
使い方
- ヘッダオンリーなのでソースをダウンロードして#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()
はバケット数もリセットしてしまう- STLの
clear()
と同じ挙動をするのはclear_no_resize()
clear_no_resize()
するときはset_deleted_key(key)
が必要
- STLの
パフォーマンス
- 使ってみると
std::unordered_map
より確かに速い- 特にlookup、数倍速くなる
- ベンチマークは以下のページなどに載っている
【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; }
サンプルコード
2017/12/11追記
warpSize
はどうやらコンパイル時定数ではないみたいです。
そのため、サンプルコードのように書くとループ展開が行われない可能性があります。
パフォーマンスを気にするのであれば warpSize
→ 32
とした方が良さそうです。
基数を取得する
最下位ビット/バイトからk番目の基数を取得する関数です。
- radix1bit → 2進数表現で右からk番目の値を取り出す (0 or 1)
- radix4bit → 16進数表現で右からk番目の値を取り出す (0 to 15)
- radix8bit → 1バイト区切りで右からk番目の値を取り出す (0 to 255)
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は全体的にサンプルコードが不足しているような気がするので、 まず普通にハッシュテーブルを使うためのサンプルコードを載せておきます。
上記のサンプルではエラーチェックをしていません。 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);
このハッシュテーブルにはcudppMultivalueHashGetAllValues
と
cudppMultivalueHashGetValuesSize
という特別な関数が用意されています。
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実装なので ぜひ試してみてください。
【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以降でのみサポート