加治/日誌/2014-01-16
をテンプレートにして作成
[
トップ
] [
新規
|
一覧
|
単語検索
|
最終更新
|
ヘルプ
|
ログイン
]
開始行:
[[加治/日誌]]
もうどの日にやってんだか境目がない
*けんきゅ [#m8d3ec90]
**FTV環境構築 [#iabc3ba8]
***実験用マシンで [#o5ca500c]
-ていうかFTVのプログラムカメラ3つ用だった。むしろ4つだとエラー出た
--ハブも変換アダプタも買った意味なかったんじゃ・・・
-ドライバの入れ直しをしてもBSoD
--なお、キャッシュを消してからでもだめだった。
--ドライババージョン1.8
-個別に刺しても認識しない時がある
-カメラのビューワだけでもBSoD
***ノートで [#ec012270]
-ビルドは通る
-ビューワは見れる、けど一つ一つじゃないと見れない(どっちのマシンでも)
-BSoDにはならないが、FTVの画面が立ち上がってこない
-ドライバ(SDK)のバージョンを確認したら2.1 Release 5だった
--先生にDLしてもらったものの中にこのバージョンはない
--矢口先生にSDK貰いに行くのと、状況を見せてアドバイスを貰うつもり
**2DCDP [#zbb37195]
***デバッグ [#sb8cc1e3]
-バックトレース on GPGPU実装前のaccmCalc()スレッド数をinp.height, inp.widthにしたらスポッティングポイントがおかしくなった
--Nsightで見てみても特にエラーが出ているわけではない
---スレッド数の限界を超えるときはエラーが出てた
--バックトレースGPGPU実装後はスレッド数を変えたからといって結果が正しくなる、おかしくなるという風には切り替わらない
-initBufferでCUDA streamのDestroyのし忘れを解消
-タイマーの問題
--後述する連続データ入力対応により、プログラムの実行中に何度も累積計算を行えるようになり、累積計算にかかる時間を正しく測定出来るかと思ったが、何度測っても同じ
-結果画像がおかしい問題
--バックトレースGPGPU実装前のプログラムを一度実行してから実装後のプログラムを実装すると結果が正しくでる(今現在)
-考えられる原因
--スレッド数の問題
---元の実装のままでもおかしかったから多分違うと思う
--変数初期化せずに値を足しちゃってる
---インクリメントや初期化が必要なデータをひと通り見たが特に問題はなさそう
--変数名の食い違い(rw=ref.width+1なのにref.widthと勘違いしていたり)
---これも確認したが大丈夫そう
--デバイスメモリのセグフォ
---これもひと通り確認したが、確保された領域と使用している領域はあっていると思うけど、強いていうならこれがちょっと不安なので再確認
-これからのデバッグ
--ビルド中に出るWARNINGを頼りに怪しいところを見つける
***高速化、最適化 [#s8774d89]
-ホスト側のaccmbufはもはや必要ないので消去。
--ref.height * ref.width * inp.height * inp.width * sizeof(float)分ホストメモリ空いた
-pickbufはcallocで、確保と同時に0初期化されていたため、確保と初期化を分離するためにmallocで確保して初期化はmemsetで0で埋めることにした
-mallocはcudaMallocHostに変える予定
-Nsightでのプロファイリングをし、初期化処理のcu_rankbufInit()、cu_rlabufInit()、pickupInit()、resbufInit()をCUDA Streamを使った並列化
-さらにその中でもcu_rankbufInit()とcu_rlabufInit()の初期化部分が大きかったため、1スレッドで実行していたものを要素数分並列化させて大きく高速化できた
--両関数で初期化される領域(要素数)はinp.height+3 * inp.width+3 * ref.height
--これをブロック、スレッドに分けたいが1ブロックあたりのスレッド数は1024
--dim3 block((inp.width+3)/32, inp.height+3) /32, ref.height+1) thread(32,32)で3次元ブロックを1024スレッドでアクセスしようとしたが、gridDim.zは最高が64ブロックのためにこの構成はできない
--半端な数になるが、スレッドも3次元構成にしてgridDim.zのブロック数を減らすことにした
---この場合スレッド最大数1024の3乗根の最大整数10で、thread(10,10,10)とした。(16,16,4)とかでもいいかも
--関数内でのインデックスアクセスは3次元スレッドを1次元にしなければならないが頭が働かず式が立てられない。調べた所こんな素晴らしいサイトが!
---http://www.martinpeniak.com/index.php?option=com_content&view=article&catid=17:updates&id=288:cuda-thread-indexing-explained
---1D grid of 1D blocks
__device__ int getGlobalIdx_1D_1D()
{
return blockIdx.x *blockDim.x + threadIdx.x;
}
---1D grid of 2D blocks
__device__ int getGlobalIdx_1D_2D()
{
return blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
}
---1D grid of 3D blocks
__device__ int getGlobalIdx_1D_3D()
{
return blockIdx.x * blockDim.x * blockDim.y * blockDim.z
+ threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;
}
---2D grid of 1D blocks
__device__ int getGlobalIdx_2D_1D()
{
int blockId = blockIdx.y * gridDim.x + blockIdx.x;
int threadId = blockId * blockDim.x + threadIdx.x;
return threadId;
}
---2D grid of 2D blocks
__device__ int getGlobalIdx_2D_2D()
{
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;
return threadId;
}
---2D grid of 3D blocks
__device__ int getGlobalIdx_2D_3D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x)
+ threadIdx.x;
return threadId;
}
---3D grid of 1D blocks
__device__ int getGlobalIdx_3D_1D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * blockDim.x + threadIdx.x;
return threadId;
}
---3D grid of 2D blocks
__device__ int getGlobalIdx_3D_2D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y)
+ (threadIdx.y * blockDim.x)
+ threadIdx.x;
return threadId;
}
---3D grid of 3D blocks
__device__ int getGlobalIdx_3D_3D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x)
+ threadIdx.x;
return threadId;
}
---今回の自分の場合は3D grid of 3D blocksなので一番下の式を使ってインデックスを求めた。
--cu_rankbufInit()は、パディング要素への代入も行わなければならないが、こちらはほとんど負荷がかからなそうなのでインデックス0のスレッドに処理させるようにした
---バックトレースGPGPU実装前
https://drive.google.com/file/d/0B7qfGAF9dTCXTXZlNjFEREVHMnc/edit?usp=sharing
---バックトレースGPGPU実装後+Stream並列化
https://drive.google.com/file/d/0B7qfGAF9dTCXT3ZzMDdWZXhPelk/edit?usp=sharing
---バックトレースGPGPU実装後+Stream並列化+初期化処理並列化
https://drive.google.com/file/d/0B7qfGAF9dTCXSS1LR3dsZmMwYWc/edit?usp=sharing
-その他「要素数がスレッド数より多い場合どうするか」など、CUDAプログラムガイド
--http://accc.riken.jp/secure/4467/cuda-programming_main.pdf
***FTV化 [#nee164ec]
-連続データ入力に対応した
--initBuffer()内のメモリ確保処理を独立化しmallocBuffer()として定義
--連続データは次のcdpInit()が呼ばれる前にinitBuffer()を呼べば初期化される
**論文 [#n54ad72f]
-矢口先生に添削して頂いたところを修正、提出
終了行:
[[加治/日誌]]
もうどの日にやってんだか境目がない
*けんきゅ [#m8d3ec90]
**FTV環境構築 [#iabc3ba8]
***実験用マシンで [#o5ca500c]
-ていうかFTVのプログラムカメラ3つ用だった。むしろ4つだとエラー出た
--ハブも変換アダプタも買った意味なかったんじゃ・・・
-ドライバの入れ直しをしてもBSoD
--なお、キャッシュを消してからでもだめだった。
--ドライババージョン1.8
-個別に刺しても認識しない時がある
-カメラのビューワだけでもBSoD
***ノートで [#ec012270]
-ビルドは通る
-ビューワは見れる、けど一つ一つじゃないと見れない(どっちのマシンでも)
-BSoDにはならないが、FTVの画面が立ち上がってこない
-ドライバ(SDK)のバージョンを確認したら2.1 Release 5だった
--先生にDLしてもらったものの中にこのバージョンはない
--矢口先生にSDK貰いに行くのと、状況を見せてアドバイスを貰うつもり
**2DCDP [#zbb37195]
***デバッグ [#sb8cc1e3]
-バックトレース on GPGPU実装前のaccmCalc()スレッド数をinp.height, inp.widthにしたらスポッティングポイントがおかしくなった
--Nsightで見てみても特にエラーが出ているわけではない
---スレッド数の限界を超えるときはエラーが出てた
--バックトレースGPGPU実装後はスレッド数を変えたからといって結果が正しくなる、おかしくなるという風には切り替わらない
-initBufferでCUDA streamのDestroyのし忘れを解消
-タイマーの問題
--後述する連続データ入力対応により、プログラムの実行中に何度も累積計算を行えるようになり、累積計算にかかる時間を正しく測定出来るかと思ったが、何度測っても同じ
-結果画像がおかしい問題
--バックトレースGPGPU実装前のプログラムを一度実行してから実装後のプログラムを実装すると結果が正しくでる(今現在)
-考えられる原因
--スレッド数の問題
---元の実装のままでもおかしかったから多分違うと思う
--変数初期化せずに値を足しちゃってる
---インクリメントや初期化が必要なデータをひと通り見たが特に問題はなさそう
--変数名の食い違い(rw=ref.width+1なのにref.widthと勘違いしていたり)
---これも確認したが大丈夫そう
--デバイスメモリのセグフォ
---これもひと通り確認したが、確保された領域と使用している領域はあっていると思うけど、強いていうならこれがちょっと不安なので再確認
-これからのデバッグ
--ビルド中に出るWARNINGを頼りに怪しいところを見つける
***高速化、最適化 [#s8774d89]
-ホスト側のaccmbufはもはや必要ないので消去。
--ref.height * ref.width * inp.height * inp.width * sizeof(float)分ホストメモリ空いた
-pickbufはcallocで、確保と同時に0初期化されていたため、確保と初期化を分離するためにmallocで確保して初期化はmemsetで0で埋めることにした
-mallocはcudaMallocHostに変える予定
-Nsightでのプロファイリングをし、初期化処理のcu_rankbufInit()、cu_rlabufInit()、pickupInit()、resbufInit()をCUDA Streamを使った並列化
-さらにその中でもcu_rankbufInit()とcu_rlabufInit()の初期化部分が大きかったため、1スレッドで実行していたものを要素数分並列化させて大きく高速化できた
--両関数で初期化される領域(要素数)はinp.height+3 * inp.width+3 * ref.height
--これをブロック、スレッドに分けたいが1ブロックあたりのスレッド数は1024
--dim3 block((inp.width+3)/32, inp.height+3) /32, ref.height+1) thread(32,32)で3次元ブロックを1024スレッドでアクセスしようとしたが、gridDim.zは最高が64ブロックのためにこの構成はできない
--半端な数になるが、スレッドも3次元構成にしてgridDim.zのブロック数を減らすことにした
---この場合スレッド最大数1024の3乗根の最大整数10で、thread(10,10,10)とした。(16,16,4)とかでもいいかも
--関数内でのインデックスアクセスは3次元スレッドを1次元にしなければならないが頭が働かず式が立てられない。調べた所こんな素晴らしいサイトが!
---http://www.martinpeniak.com/index.php?option=com_content&view=article&catid=17:updates&id=288:cuda-thread-indexing-explained
---1D grid of 1D blocks
__device__ int getGlobalIdx_1D_1D()
{
return blockIdx.x *blockDim.x + threadIdx.x;
}
---1D grid of 2D blocks
__device__ int getGlobalIdx_1D_2D()
{
return blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
}
---1D grid of 3D blocks
__device__ int getGlobalIdx_1D_3D()
{
return blockIdx.x * blockDim.x * blockDim.y * blockDim.z
+ threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;
}
---2D grid of 1D blocks
__device__ int getGlobalIdx_2D_1D()
{
int blockId = blockIdx.y * gridDim.x + blockIdx.x;
int threadId = blockId * blockDim.x + threadIdx.x;
return threadId;
}
---2D grid of 2D blocks
__device__ int getGlobalIdx_2D_2D()
{
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;
return threadId;
}
---2D grid of 3D blocks
__device__ int getGlobalIdx_2D_3D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x)
+ threadIdx.x;
return threadId;
}
---3D grid of 1D blocks
__device__ int getGlobalIdx_3D_1D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * blockDim.x + threadIdx.x;
return threadId;
}
---3D grid of 2D blocks
__device__ int getGlobalIdx_3D_2D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y)
+ (threadIdx.y * blockDim.x)
+ threadIdx.x;
return threadId;
}
---3D grid of 3D blocks
__device__ int getGlobalIdx_3D_3D()
{
int blockId = blockIdx.x
+ blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x)
+ threadIdx.x;
return threadId;
}
---今回の自分の場合は3D grid of 3D blocksなので一番下の式を使ってインデックスを求めた。
--cu_rankbufInit()は、パディング要素への代入も行わなければならないが、こちらはほとんど負荷がかからなそうなのでインデックス0のスレッドに処理させるようにした
---バックトレースGPGPU実装前
https://drive.google.com/file/d/0B7qfGAF9dTCXTXZlNjFEREVHMnc/edit?usp=sharing
---バックトレースGPGPU実装後+Stream並列化
https://drive.google.com/file/d/0B7qfGAF9dTCXT3ZzMDdWZXhPelk/edit?usp=sharing
---バックトレースGPGPU実装後+Stream並列化+初期化処理並列化
https://drive.google.com/file/d/0B7qfGAF9dTCXSS1LR3dsZmMwYWc/edit?usp=sharing
-その他「要素数がスレッド数より多い場合どうするか」など、CUDAプログラムガイド
--http://accc.riken.jp/secure/4467/cuda-programming_main.pdf
***FTV化 [#nee164ec]
-連続データ入力に対応した
--initBuffer()内のメモリ確保処理を独立化しmallocBuffer()として定義
--連続データは次のcdpInit()が呼ばれる前にinitBuffer()を呼べば初期化される
**論文 [#n54ad72f]
-矢口先生に添削して頂いたところを修正、提出
ページ名: