はじめに
ComputeShader.Dispatch命令はGPUで計算を行なうためのものです。
ComputeShader.Dispatch - Unityスクリプトリファレンス
この命令はGPU計算の完了を保証するものではありません。つまり「GPUの計算がおわる前にCPUは次の命令にうつる」となります。
ここまではまぁ知られていると思います。
普通に考えればGPUが計算している最中にCPUも計算できてラッキーとなるのですが、実はGPUで計算がはじまることも保証されてません。
結論からかくとGPUで計算ははじまらず
GLFlush();
を行なうことで即時計算がはじまります。
CPUもGPUも無駄な時間をかかえている!
1フレーム内の動きを考えます。
あるタイミングでComputeShader.Dispatchを実行します。
その後GPUで計算が走っていることを期待しCPUの重めの処理を実行し、GetDataでGPUから計算結果をとってきます。
この間、GPUもCPUも動いているので効率がいいです(と私は思っていました)。
可能な限りフレーム内の最初のほうでComputeShader.Dispatchを行ない、フレーム内の最後のほうでGetDataを行なうことでCPUとGPUをフルに活用できるのではないか、と思っていました。
しかし実際は、こんなことになっていました。↓
なんとCPUがGetDataの行に到達してやっとGPUが動き出し、GPUの計算がおわってデータがCPUに転送されてはじめてCPUは次の命令にうつれます。
CPUもGPUも無駄な時間をかかえてしまっているのです!
検証
上記結論にいたるのに以下のコードで検証しました。
C#
using System;
using UnityEngine;
using System.Threading;
public class DispatchTest : MonoBehaviour
{
const int width = 1920;
const int height = 1080;
//compute shader周り
[SerializeField] ComputeShader cs;
private ComputeBuffer computeBuffer = null;
private int kernelHeavyFunc;
//その他
private int flamecnt;
private int n;
private int[] host;
private int lasttime;
private int deltatime;
private int time12, time23, time34;
private float avgdeltatime;
[SerializeField] int loopnum = 600;//強いGPUならもっと増やして
private void Awake()
{
Application.targetFrameRate = 60;
flamecnt = 0;
avgdeltatime = 0;
lasttime = Gettime();
n = width * height;
computeBuffer = new ComputeBuffer(n, 4);
host = new int[n];
kernelHeavyFunc = cs.FindKernel("HeavyFunc");
cs.SetInt("WX", width);
cs.SetInt("WY", height);
cs.SetBuffer(kernelHeavyFunc, "buffer", computeBuffer);
}
private void Update()
{
var time1 = Gettime();
cs.SetInt("loopnum", loopnum);
cs.Dispatch(kernelHeavyFunc, 1, height, 1);
GL.Flush();
var time2 = Gettime();
Thread.Sleep(50);//50ms待つ。CPUの重い処理を想定
var time3 = Gettime();
computeBuffer.GetData(host, 0, 0, 4);//GPUの計算が終わっていれば一瞬のはず
var time4 = Gettime();
time12 = time2 - time1;
time23 = time3 - time2;
time34 = time4 - time3;
deltatime = time3 - lasttime;
lasttime = time3;
flamecnt++;
}
private void OnDestroy()
{
computeBuffer.Release();
}
void OnGUI()
{
avgdeltatime = avgdeltatime * 0.99f + 0.01f * deltatime;
GUILayout.Label($"flamecnt:\t{flamecnt}");
GUILayout.Label($"1フレームのtime:\t{deltatime}ms");
GUILayout.Label($"平均FPS:\t{1000.0f / avgdeltatime}");
GUILayout.Label($"time12:\t{time12}ms");
GUILayout.Label($"time23:\t{time23}ms");
GUILayout.Label($"time34:\t{time34}ms");
// グラフィックデバイス名
GUILayout.Label($"グラフィックデバイス名:{UnityEngine.SystemInfo.graphicsDeviceName}");
// グラフィックスAPIタイプ(Direct3D11とか)
var s0 = UnityEngine.SystemInfo.graphicsDeviceVersion;
GUILayout.Label($"グラフィックスデバイスバージョン:{s0}");
// シェーダレベル
GUILayout.Label($"シェーダレベル:{UnityEngine.SystemInfo.graphicsShaderLevel}");
// コンピュートシェーダが使えるか
GUILayout.Label($"コンピュートシェーダが使えるか:{UnityEngine.SystemInfo.supportsComputeShaders}");
}
//現在の時刻をms単位で取得
int Gettime()
{
return DateTime.Now.Millisecond + DateTime.Now.Second * 1000
+ DateTime.Now.Minute * 60 * 1000 + DateTime.Now.Hour * 60 * 60 * 1000;
}
}
ComputeShader
#pragma kernel HeavyFunc
#define TPB 64
RWStructuredBuffer<uint> buffer;
int loopnum, WX, WY;
[numthreads(TPB, 1, 1)]
void HeavyFunc(uint2 id : SV_DispatchThreadID)
{
float tt = 0.0;
for (int j = 0; j < WX / TPB; j++)
{
for (int i = 0; i < loopnum; i++)
{
tt = cos(id.x + tt + 3.0 / (1.2 + sin(0.0011 / (id.x + 3.42134))));
tt = 3.4 / sqrt((id.y + 0.123) / (tt + 2.0));
tt = cos(tt * 1.2) * sin(0.3 + tt) + exp(2.3 * tt);
}
//uint outt = (uint)tt;
buffer[(j * TPB + id.x) + id.y * WX] = tt;
}
}
解説
Update()では各命令がどのくらいの時間を消費しているか確認しています。
Thread.Sleep(50)で50ms待っていますが、ここはゲームロジックなどCPU側の重い処理がくることを想定しています。
computeBuffer.GetDataでは4*4byteしか転送していないので命令オーバーヘッドの分しか時間を食いません。0-1ms程度の負荷と想定されます。GPUの計算が終わってなければ、それに応じてmsが伸びます。
var time1 = Gettime();
cs.Dispatch(kernelHeavyFunc, 1, height, 1);
GL.Flush();
var time2 = Gettime();
Thread.Sleep(50);//50ms待つ。CPUの重い処理を想定
var time3 = Gettime();
computeBuffer.GetData(host, 0, 0, 4);//GPUの計算が終わっていれば一瞬のはず
var time4 = Gettime();
このtime1~time4から命令の実行時間がわかります。
このコードではGL.Flush();がかかれていますが、ここをコメントアウトしたりして速度を測りました。
もしGPUで計算が走りながらThread.Sleep(50)が実行されていれば、GetDataの行はその分時間が短くなるはずです。
もしGPUで計算が走らずにThread.Sleep(50)が実行されていれば、GetDataの行はGPU時間のぶんだけ待たされるはずです。
結果(AMD,NVIDIA,Intelで測定)
AMD Radeon RX 6600
GL.Flush()なし
GL.Flush()あり
注目すべきは
「time34」のところです。
GL.Flush()なし:61ms
GL.Flush()あり:13ms
time3~time4の間の時間を測っていますが、そこにあるのは
computeBuffer.GetData(host, 0, 0, 4);
だけです。
ここでGPUの結果まちをしているのですが、DispatchのGPU負荷が約60msであり、GL.Flush()ありのほうではSleep(50)の間にもGPUの計算が進んでいたことがわかります。
念のためPIXでDispatchの負荷を測定してみました。
GL.Flush()なし
GL.Flush()あり
当然どちらも約60msです。(測定の影響かやや66msと負荷が増え気味?)
NVIDIA RTX 2080Ti
GL.Flush()なし
GL.Flush()あり
RTX2080Tiでも60msくらいの負荷になるよう調整し
GL.Flush()なし:59ms
GL.Flush()あり:9ms
NVIDIAでも同じ傾向がみられます。
Intel(R) Iris(R) Xe Graphics
GL.Flush()なし
GL.Flush()あり
GL.Flush()なし:93ms
GL.Flush()あり:48ms
やはりIntelでも同じ傾向がみられます。
じゃあ実行タイミングはいつ?
3主要ベンダーのGPUでGLFlush()をしないとDispatchの計算がはじまらないことがわかりました。
しかしまだタイトルの回収ができていません。
今度は1フレームだけUpdate()内でDispatchを実行するようにして、それ以外のフレームでは何も行ないません。その条件下で各フレームのタイミングを測ってみます。
ここにはのせてませんが、Start()内で一度DispatchとGetDataを行って初回カーネル起動のオーバーヘッドを解消しています。また起動直後のフレームは不安定なので100フレーム目から計測です。
private void Update()
{
var time3 = Gettime();
if ((flamecnt >= 100) &(flamecnt < 116))
{
times[flamecnt - 100] = Gettime();
if (flamecnt == 101)
cs.Dispatch(kernelHeavyFunc, 1, height, 1);
}
deltatime = time3 - lasttime;
lasttime = time3;
flamecnt++;
}
結果
()内の数字は前フレームとの差のmsです。
101フレーム目のUpdate()内でDispatchを実行、次の102フレーム目のUpdate()で時間を測るとすでに60ms経過しています。
101フレーム目と102フレーム目の間でたしかにDispatchが実行完了していることがわかります。
Dispatchの完了がどこか探す
さて、そもそも101フレーム目の中でDispatchが実行完了しているのでしょうか。
これをみるとOnGUIのタイミングが遅いのでそこで時間を取得してみることにします。
→OnGUIよりさらに遅いタイミングで実行するOnEndOfFrame()を見つけたのでこれで実行してみます。
【Unity】LateUpdate()やOnGUI()よりもタイミングが遅い、OnEndOfFrame()を作ってみた
private void Update()
{
flamecnt++;
var time3 = Gettime();
if (flamecnt == 101)
cs.Dispatch(kernelHeavyFunc, 1, height, 1);
deltatime = time3 - lasttime;
lasttime = time3;
}
//中略
private void OnEndOfFrame
{
if ((flamecnt >= 100) & (flamecnt < 116))
{
times[flamecnt - 100] = Gettime();
}
}
フレームカウント(flamecnt)はOnEndOfFrameの後にインクリメントされてほしいので今度はUpdateの最初にかきました。
結果は
やはり101フレーム目の最後の時点では完了してないようです。
つまりDispatchはそれを実行したフレームのOnGUI(~OnEndOfFrame())の後まで持ちこされ、GPUで実行完了してから次のフレームがはじまる、という認識で良さそうです。
最後に淡い期待をもって、DispatchがOnEndOfFrame()の手前で「実行開始されている」可能性について検証してみましょう。
Dispatchの開始がどこか探す
101フレーム目内のどこかでDispatchの開始がされているならば、OnEndOfFrame()でSleepとGetDataを駆使して計測できるはず!
private void Update()
{
flamecnt++;
var time3 = Gettime();
if (flamecnt == 101)
{
cs.Dispatch(kernelHeavyFunc, 1, height, 1);
//GL.Flush();
}
deltatime = time3 - lasttime;
lasttime = time3;
}
//中略
private void OnEndOfFrame()
{
if (flamecnt == 101)
{
Thread.Sleep(50);
var tm1 = Gettime();
computeBuffer.GetData(host, 0, 0, 4);
var tm2 = Gettime();
Debug.Log(tm2 - tm1);
}
}
また同じように
Thread.Sleep(50);
のあとにGetDataをかいてGettime()ではさみます。
tm2-tm1がGetData命令の実行時間をあらわします。Debug.Logで確認してみます。
残念・・・。Dispatch自体が約60msの負荷なのでSleep(50)の間にGPU実行されてなかったようです。
そしてUpdate()内のGL.Flush();のコメントアウトをはずしたら見事に12msと表示されました。
結論:真の実行タイミングはフレームの一番最後
ComputeShader.DispatchによりGPUで計算が開始されるタイミングはGLFlush()をしなければそのフレームの一番最後、GLFlush()をすれば即時開始。
ComputeShader.DispatchによりGPUで計算が終了するタイミングは、GLFlush()をしなければそのフレームの一番最後、GLFlush()をすればCPU負荷に関係なくGPUが計算を終えた瞬間。
検証で使ったソースコードはこちら
参考文献
Unity Forumで2018年に議論されたようです
【コンピューティング シェーダー非同期の問題 [解決済み]】
公式DocにはDX11ならID3D11DeviceContext::Flushが呼ばれ、DX12ならpending command lists が実行され、OpenGLならgl.Flushが呼ばれると書いてあります。
Unity - Scripting API:GL.Flush
Microsoft公式
ID3D12CommandQueue::ExecuteCommandLists method (d3d12.h)
補足
私はOpenCLでGPGPUをやって遊んでいます。OpenCLにおいてもDispatchのような命令(clEnqueueNDRangeKernel)があり、これが即時実行でないのをおぼろげに知っていました。またclFlush()というGLFlush()と同じようなものが存在することも頭の片隅にあったのですが、Unity Compute Shaderでも同じことが起こっているとは思いませんでした・・。
DX11やOpenGLでも同様の現象はあるようで、GPU計算全般にいえることなのかもしれません。