CUDA

http://www.nvidia.com/content/GTC-2010/pdfs/2238_GTC2010.pdf
Thread-level parallelism (TLP) にばっかり目を向けるんじゃなくてInstruction-level parallelism (ILP)もちゃんと見た方が良いというお話

PTXの確認方法

• • keep --source-in-ptx -G

を付ける。デバッグ情報を付加するとCのコードも確認出来るけど実行速度がデバッグと同じぐらいになるような…。。

cudaMallocPitch

アラインメントを合わせたらパフォーマンスが上がった。
しかしデバイス側でピッチのバイト数でオフセットする下記のようなコードを毎回書くのが大変だ。

p_hoge = (const uint16_t*)((const char*)p_hoge + hoge_pitch);

template <typename T>
__device__ __forceinline__
const T* offset_ptr(const T* ptr, int num_bytes)
{
	return (const T*)((const char*)(ptr) + num_bytes);
}

こんなテンプレート関数を作って使ってみたら実行時間が元のより掛かってしまうようになった。何故だ。。

#define OFFSET_PTR(ptr, num_bytes) ((decltype(ptr))((const char*)ptr + num_bytes))

こういうマクロなら処理速度が変わらなかった。