10
5

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?

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

Last updated at Posted at 2023-03-10

はじめに

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も動いているので効率がいいです(と私は思っていました)。
f3186b25.png
可能な限りフレーム内の最初のほうでComputeShader.Dispatchを行ない、フレーム内の最後のほうでGetDataを行なうことでCPUとGPUをフルに活用できるのではないか、と思っていました。

しかし実際は、こんなことになっていました。↓
gre252.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

GL.Flush()あり

amdari.png

注目すべきは
「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()なし

glflushなし.png

GL.Flush()あり

GLFlushあり.png
当然どちらも約60msです。(測定の影響かやや66msと負荷が増え気味?)

NVIDIA RTX 2080Ti

GL.Flush()なし

FlushなしRTX2080Ti.png

GL.Flush()あり

FlushありRTX2080Ti.png
RTX2080Tiでも60msくらいの負荷になるよう調整し
GL.Flush()なし:59ms
GL.Flush()あり:9ms

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

Intel(R) Iris(R) Xe Graphics

GL.Flush()なし

inteflushなしい.png

GL.Flush()あり

intelflush.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

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

Dispatchの完了がどこか探す

さて、そもそも101フレーム目の中でDispatchが実行完了しているのでしょうか。
image.png

これをみるとOnGUIのタイミングが遅いのでそこで時間を取得してみることにします。
→OnGUIよりさらに遅いタイミングで実行するOnEndOfFrame()を見つけたのでこれで実行してみます。
【Unity】LateUpdate()やOnGUI()よりもタイミングが遅い、OnEndOfFrame()を作ってみた

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

やはり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

残念・・・。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計算全般にいえることなのかもしれません。

10
5
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
10
5

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?