A10-7850KとかいうクソザコCPUを用意した。

これの hsa_queue_store_write_index_relaxed() 〜 hsa_signal_wait_acquire() の間の時間が10usecぐらい。

CPU間でも同期オブジェクト経由でスレッド間通信したら10usecぐらいかかるから、GPUがCPUと同じレイテンシで呼べるようになってると言って良いのでは。

hsa_memory_registerはパフォーマンスのために必要なだけ(とどっかに書いてあった気がする)で、消しても動く。

つまり、これまでのAPU()と違ってまともに使って便利なレベルになったと理解してよいと思う。

あとHSAでは、メモリ経由で通信しても読み書きできるのが保証されてるはずだからメモリ経由で通信したときのレイテンシを計測すべき。(そのためにはHSAIL書かないといけないのだが…)

■

あと https://bwidawsk.net/ これ誰か読んで内容教えて。

いや

10usec はおかしい気がする(4MB read write してるので理論値越えてる)。

あとで調べるからちょっと待って。

調べた。10usecでいい。まあプログラム間違ってたけど直しても10usecだった。