ComputeShader.Dispatchの真の実行タイミングはいつ？【Unity】

# はじめに
ComputeShader.Dispatch命令はGPUで計算を行なうためのものです。
[ComputeShader.Dispatch - Unityスクリプトリファレンス](https://docs.unity3d.com/ja/2020.3/ScriptReference/ComputeShader.Dispatch.html)

この命令はGPU計算の完了を保証するものではありません。つまり「GPUの計算がおわる前にCPUは次の命令にうつる」となります。
ここまではまぁ知られていると思います。

普通に考えればGPUが計算している最中にCPUも計算できてラッキーとなるのですが、実は<B>GPUで計算がはじまることも保証されてません。</B>

結論からかくとGPUで計算ははじまらず
`
GLFlush();
`
を行なうことで即時計算がはじまります。

# CPUもGPUも無駄な時間をかかえている！

1フレーム内の動きを考えます。
あるタイミングでComputeShader.Dispatchを実行します。
その後GPUで計算が走っていることを期待しCPUの重めの処理を実行し、GetDataでGPUから計算結果をとってきます。

この間、GPUもCPUも動いているので効率がいいです(と私は思っていました)。
![f3186b25.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/06c18f6d-5564-5f5a-c649-b9fe8b2d9cc9.png)
可能な限りフレーム内の最初のほうでComputeShader.Dispatchを行ない、フレーム内の最後のほうでGetDataを行なうことでCPUとGPUをフルに活用できるのではないか、と思っていました。

しかし実際は、こんなことになっていました。↓
![gre252.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/b976e44d-6623-fdab-b4a7-e8cbb4dc3af9.png)
なんとCPUがGetDataの行に到達してやっとGPUが動き出し、GPUの計算がおわってデータがCPUに転送されてはじめてCPUは次の命令にうつれます。
CPUもGPUも無駄な時間をかかえてしまっているのです！

## 検証

上記結論にいたるのに以下のコードで検証しました。

C#

```DispatchTest.cs
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

```HeavyComputeShader.c
#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が伸びます。

```DispatchTest.cs
        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()なし
![amdnasi.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/5a26ab80-9942-31e2-cb87-2bd522d4b761.png)

##### GL.Flush()あり
![amdari.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/aa1c8b71-2120-1f2f-3dcd-5b31aa3a5d13.png)

注目すべきは
「time34」のところです。
GL.Flush()なし:61ms
GL.Flush()あり:13ms

time3～time4の間の時間を測っていますが、そこにあるのは
`
computeBuffer.GetData(host, 0, 0, 4);
`
だけです。
ここでGPUの結果まちをしているのですが、DispatchのGPU負荷が約60msであり、<INS>GL.Flush()ありのほうではSleep(50)の間にもGPUの計算が進んでいたことがわかります。</INS>

念のためPIXでDispatchの負荷を測定してみました。
##### GL.Flush()なし
![glflushなし.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/e97cc137-d575-5e97-2112-4110d321b21e.png)
##### GL.Flush()あり
![GLFlushあり.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/0085f094-cc75-2e6a-a724-c2e958a08d13.png)
当然どちらも約60msです。(測定の影響かやや66msと負荷が増え気味？)

#### NVIDIA RTX 2080Ti
##### GL.Flush()なし
![FlushなしRTX2080Ti.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/4ece7d9d-9a3d-0a59-e48c-6d8b84cd7a2d.png)
##### GL.Flush()あり
![FlushありRTX2080Ti.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/ec09f6eb-0f15-6e8e-3d8a-e1c8db41dbd4.png)
RTX2080Tiでも60msくらいの負荷になるよう調整し
GL.Flush()なし:59ms
GL.Flush()あり:9ms

NVIDIAでも同じ傾向がみられます。


#### Intel(R) Iris(R) Xe Graphics
##### GL.Flush()なし
![inteflushなしい.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/4d95b1e6-6716-9692-52df-6794d5bbca7f.png)
##### GL.Flush()あり
![intelflush.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/6375f470-886d-3f1b-45f4-4159e4baf935.png)

GL.Flush()なし:93ms
GL.Flush()あり:48ms

やはりIntelでも同じ傾向がみられます。


### じゃあ実行タイミングはいつ？

3主要ベンダーのGPUでGLFlush()をしないとDispatchの計算がはじまらないことがわかりました。
しかしまだタイトルの回収ができていません。
今度は1フレームだけUpdate()内でDispatchを実行するようにして、それ以外のフレームでは何も行ないません。その条件下で各フレームのタイミングを測ってみます。
ここにはのせてませんが、Start()内で一度DispatchとGetDataを行って初回カーネル起動のオーバーヘッドを解消しています。また起動直後のフレームは不安定なので100フレーム目から計測です。

```DispatchTest1.cs
    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++;
    }
```

#### 結果
![flames.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/a5cf4940-6458-5715-417f-72842da1f452.png)

()内の数字は前フレームとの差のmsです。
101フレーム目のUpdate()内でDispatchを実行、次の102フレーム目のUpdate()で時間を測るとすでに60ms経過しています。
101フレーム目と102フレーム目の間でたしかにDispatchが実行完了していることがわかります。

#### Dispatchの完了がどこか探す
さて、そもそも101フレーム目の中でDispatchが実行完了しているのでしょうか。
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/0ffc9e21-8bca-8b8c-8dbd-4401e5473933.png)

~~これをみるとOnGUIのタイミングが遅いのでそこで時間を取得してみることにします。~~
→OnGUIよりさらに遅いタイミングで実行するOnEndOfFrame()を見つけたのでこれで実行してみます。
[【Unity】LateUpdate()やOnGUI()よりもタイミングが遅い、OnEndOfFrame()を作ってみた](https://qiita.com/yuiyoichi/items/88a6df9c247e607c4311) 
  


```DispatchTest1.cs

    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の最初にかきました。
結果は
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/9faa8d04-426a-f9b8-a3e9-66c92bef74fe.png)

やはり101フレーム目の最後の時点では完了してないようです。

つまりDispatchはそれを実行したフレームのOnGUI(～OnEndOfFrame())の後まで持ちこされ、GPUで実行完了してから次のフレームがはじまる、という認識で良さそうです。

最後に淡い期待をもって、DispatchがOnEndOfFrame()の手前で「実行開始されている」可能性について検証してみましょう。

#### Dispatchの開始がどこか探す

101フレーム目内のどこかでDispatchの開始がされているならば、OnEndOfFrame()でSleepとGetDataを駆使して計測できるはず！

```DispatchTest1.cs

    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で確認してみます。
![image.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/227361/61bf915e-342d-4610-3702-0f275110af60.png)

残念・・・。Dispatch自体が約60msの負荷なのでSleep(50)の間にGPU実行されてなかったようです。
そしてUpdate()内のGL.Flush();のコメントアウトをはずしたら見事に12msと表示されました。

### 結論:真の実行タイミングはフレームの一番最後
ComputeShader.DispatchによりGPUで計算が開始されるタイミングはGLFlush()をしなければそのフレームの一番最後、GLFlush()をすれば即時開始。
ComputeShader.DispatchによりGPUで計算が終了するタイミングは、GLFlush()をしなければそのフレームの一番最後、GLFlush()をすればCPU負荷に関係なくGPUが計算を終えた瞬間。

#### 検証で使ったソースコードはこちら

https://github.com/toropippi/DispatchGLFlush

## 参考文献

Unity Forumで2018年に議論されたようです
[【コンピューティング シェーダー非同期の問題 [解決済み]】](https://forum.unity.com/threads/compute-shader-async-problem-resolved.513599/)

公式DocにはDX11ならID3D11DeviceContext::Flushが呼ばれ、DX12ならpending command lists が実行され、OpenGLならgl.Flushが呼ばれると書いてあります。
[Unity - Scripting API:GL.Flush](https://docs.unity3d.com/ScriptReference/GL.Flush.html)


Microsoft公式
[ID3D12CommandQueue::ExecuteCommandLists method (d3d12.h)](https://learn.microsoft.com/en-us/windows/win32/api/d3d12/nf-d3d12-id3d12commandqueue-executecommandlists)

## 補足
私はOpenCLでGPGPUをやって遊んでいます。OpenCLにおいてもDispatchのような命令(clEnqueueNDRangeKernel)があり、これが即時実行でないのをおぼろげに知っていました。またclFlush()というGLFlush()と同じようなものが存在することも頭の片隅にあったのですが、Unity Compute Shaderでも同じことが起こっているとは思いませんでした・・。
DX11やOpenGLでも同様の現象はあるようで、GPU計算全般にいえることなのかもしれません。