How to Avoid Branching on the GPU 如何在GPU避免分支

在GPU避免分支的方法

Authored by Brandon Fogerty(XR Graphics Engineer at Unity Technologies.) with Additional Organized by JP.Lee(李正彪）[email protected]

概要

這邊來稿文章中，希望了解如何編寫GPU着色器友好型代碼，以避免和分支相關的性能費用。

"分支"有着什麼意義？

使用明确的 if / then随時都可以産生分支。編譯器遇到條件會作出指令。GPU可去的地方有兩處。

是以需要決定要使用哪種代碼路徑。以下示例展示了在彙編GPU裡面運作的示例。

運算變量設定為幻數7，則添加該數字。否則的話要減掉。

為什麼要避免分支？

為了增加CPU使用率，大多數計算機都嘗試在管線上執行盡可能多的任務。彙編指令按順序執行，CPU會嘗試在盡可能多的CPU核心上執行盡可能多的指令。舉個簡單的例子。

想象一下，在Twilight Zone内，我們就是世界上最高效的程式員!

用兩行代碼寫了電子遊戲！

CPU會盡可能在每一個CPU核心上各執行一個，共計兩個光榮的指令。

這是計算機能夠更有效的運作代碼。

但是在分支時，計算機可能會花時間準備未運作的代碼結果。從結論上看處理器時間被浪費，影響了遊戲或應用程式的應答能力。

CPU 側分枝誤診通常會導緻周期遺漏超過40次。– Charles Sanglimsuwan (內建開發人員相關工程師)

運氣不錯，最新的CPU處理器速度驚人，實際分支預測出色，是以分支錯失幾乎不會有問題。

但是GPU仍然存在性能問題。

GPU嘗試并行解決大量計算，是以大部分GPU不支援分支預測

GPU為什麼會發生性能問題？

GPU為生成美麗的圖像，喜歡并行進行很多工作！

GPU為進行多種固有的結果計算，進行了精心設計，解決了通過單一濾鏡（例如：着色器程式）實施的多種輸入和相關的問題。

這就是渲染經常使用GPU(Graphics processor units)的原因。通過幾個Shader程式運作位置，法線，貼圖坐标等具有不同屬性的固有頂點，該Shader程式輸出在畫面上表現的大量固有像素顔色。

舉個例子，顯示器是1080 × 960這樣的一般HD像素的情況下，GPU在一組輸入時，将計算1080 × 960 = 1,036,800個的固有像素顔色值！

那是很多的計算結果！

想象一下在1080x960的高清遊戲體驗。遊戲嘗試每秒渲染60個固有的圖像或幀。是以GPU在1秒之内需要計算1080x960x60以上的固有像素顔色值！

結論上看一秒之内可以計算62,208,000個固有像素值！

撰寫本文的節點，NVDIA最現金的GPU是GeForce RTX 2080 Ti。它有着比英特爾最先進的CPU處理器i9-7980XE更多4,352個Cuda核心，18個核心。

非常大的差異吧？

GPU有大量的處理核心，但是與最新的CPU相比，其核心處理數量仍不足百萬。

至少目前還沒有！顯示卡在實體上太大了，可能會很昂貴。

是以，GPU将嘗試對要解決的問題類型進行特定假設。GPU充分利用SIMD。SIMD表示單個指令的多種資料。SIMD允許以并行方式為多個輸入運作計算。SIMD通常希望輸入和輸出位于相鄰的記憶體塊中。是以SIMD操作不必加載每個輸入并單獨儲存每個結果，而是加載輸入并将結果儲存為單個加載/儲存操作，進而減少昂貴的記憶體加載和儲存。SIMD的使用要求記憶體布局嚴格，并在應用程式設計中維持更好地記憶體緩存一緻性上有積極作用。SIMD在CPU和GPU均可使用.

在CPU使用SIMD的示例可參考此處。與CPU程式不同，GPU程式幾乎始終使用SIMD。為了充分利用SIMD，GPU通常每個核心都有很多ALU。ALU是算術邏輯裝置的縮寫。ALU執行數學指令。舉個例子，使用“add 4, 3”指令計算結果7。是以，單一核心上有8個ALU的話，就可以同僚并行運作8個計算。即，單一GPU核心有8個ALU，則單個GPU核心則可以同僚計算8個像素值！但是，着色器階段使用動态分支會發生什麼呢？核心是必須執行所有代碼路徑，并最終丢棄不滿足條件的代碼路徑。這意味着ALU會浪費時間執行任何未使用的操作。也就是說渲染圖像需要更長的時間。

使用示例：

在Unity裡，使用立體shader以保證雙眼能夠準确渲染。左眼以綠色表示渲染，右眼渲染的顯示為紅色。下面是使用立體shader渲染雙眼鎖渲染的單個球體。

針對此的純接近方式是基于unity_StereoEyeIndex變數的條件方法。

Shader "XR/StereoEyeIndexColor"
{
    Properties
    {
        _LeftEyeColor("Left Eye Color", COLOR) = (0,1,0,1)
        _RightEyeColor("Right Eye Color", COLOR) = (1,0,0,1)
    }

    SubShader
    {
        Tags { "RenderType" = "Opaque" }

        Pass
        {
            CGPROGRAM
            #pragma vertex vert
            #pragma fragment frag

            float4 _LeftEyeColor;
            float4 _RightEyeColor;

            #include "UnityCG.cginc"

            struct appdata
            {
                float4 vertex : POSITION;
                UNITY_VERTEX_INPUT_INSTANCE_ID
            };

            struct v2f
            {
                float4 vertex : SV_POSITION;
                UNITY_VERTEX_OUTPUT_STEREO
            };

            v2f vert (appdata v)
            {
                v2f o;

                UNITY_SETUP_INSTANCE_ID(v);
                UNITY_INITIALIZE_OUTPUT(v2f, o);
                UNITY_INITIALIZE_VERTEX_OUTPUT_STEREO(o);

                o.vertex = UnityObjectToClipPos(v.vertex);

                return o;
            }

            fixed4 frag (v2f i) : SV_Target
            {
                UNITY_SETUP_STEREO_EYE_INDEX_POST_VERTEX(i);
                if(unity_StereoEyeIndex == 0)
                {
                	return _LeftEyeColor;
                }

                return _RightEyeColor;
            }
            ENDCG
        }
    }
}
           

分段shader的集合輸出如下:

Shader hash 084f3be2-e9e6b34f-66f52d0a-1b95fab0

ps_4_0
      dcl_constantbuffer cb0[2], immediateIndexed
      dcl_input_ps_siv v1.x, rendertarget_array_index
      dcl_output o0.xyzw
   0: if_z v1.x
   1:   mov o0.xyzw, cb0[0].xyzw
   2:   ret
   3: endif
   4: mov o0.xyzw, cb0[1].xyzw
   5: ret
           

注意從0行開始的分支。分段shader stage上指定的核心将強制所有ALU執行兩個條件結果。

但是不滿足條件的結果将被丢棄。這種情況，可能是一部分ALU執行了被浪費的操作。我們剛剛讓幾種ALU很傷心，現在他們提出了他們為什麼需要存在。

不要讓可憐的小ALU做很忙的工作了。(✖╭╮✖)

你肯定不喜歡浪費時間吧？用ALU做一些遊泳的事情吧。

如果能夠去掉條件部，可以減少浪費，讓ALU更有效率，更加幸福！

以下示例示範了分段着色器的實作方式：

Shader "XR/StereoEyeIndexColor"
{
    Properties
    {
        _LeftEyeColor("Left Eye Color", COLOR) = (0,1,0,1)
        _RightEyeColor("Right Eye Color", COLOR) = (1,0,0,1)
    }

    SubShader
    {
        Tags { "RenderType" = "Opaque" }

        Pass
        {
            CGPROGRAM
            #pragma vertex vert
            #pragma fragment frag

            float4 _LeftEyeColor;
            float4 _RightEyeColor;

            #include "UnityCG.cginc"

            struct appdata
            {
                float4 vertex : POSITION;
                UNITY_VERTEX_INPUT_INSTANCE_ID
            };

            struct v2f
            {
                float4 vertex : SV_POSITION;
                UNITY_VERTEX_OUTPUT_STEREO
            };

            v2f vert (appdata v)
            {
                v2f o;

                UNITY_SETUP_INSTANCE_ID(v);
                UNITY_INITIALIZE_OUTPUT(v2f, o);
                UNITY_INITIALIZE_VERTEX_OUTPUT_STEREO(o);

                o.vertex = UnityObjectToClipPos(v.vertex);

                return o;
            }

            fixed4 frag (v2f i) : SV_Target
            {
                UNITY_SETUP_STEREO_EYE_INDEX_POST_VERTEX(i);
                return lerp(_LeftEyeColor, _RightEyeColor, unity_StereoEyeIndex);
            }
            ENDCG
        }
    }
}
           

沒有明确的條件，仔細看下HLSL元件輸出。

Shader hash 7f0a4d98-21be8f11-77007603-2899b3a0

ps_4_0
      dcl_constantbuffer cb0[2], immediateIndexed
      dcl_input_ps_siv v1.x, rendertarget_array_index
      dcl_output o0.xyzw
      dcl_temps 2
   0: utof r0.x, v1.x
   1: add r1.xyzw, -cb0[0].xyzw, cb0[1].xyzw
   2: mad o0.xyzw, r0.xxxx, r1.xyzw, cb0[0].xyzw
   3: ret
           

再一次确認也是！沒有分支！

我們現在使用的不是明确的條件部，而是Lerp。

lerp在兩個值之間執行線型插值。

Lerp本質上是：

float lerp( float a, float b, float t)
{
            return (1.0f-t)*a + t*b;
}
           

如元件輸出所示，lerp會轉換一些比乘法和分支更好地附加輸出。基本上會向lerp函數提供兩個任意值，傳遞第三個參數"t value"。t值一般是易于操作的空間數0.0和1.0之間的數字。

t值為0.0，結果為a。如果t值為1，則結果為b。

t值介于0.0和1.0之間的話，結果将在a和b之間線性顯示。

對于我們來說，我們知道unity_StereoEyeIndex是左眼或者右眼分别是0或者1。

是以，這可以很好的操作，并可以避免明确的分支。

避免分支的其他方法是什麼？

預設情況下，我們希望得到的是将條件轉換為數學方程式。為了幫助實作這個目标，我們繼續介紹HLSL硬體加速功能。

step(y,x)

x 參數比y參數大或者一樣的話，1；反之則是0。

lerp(x,y,t)

線性插值結果。

min(x,y)

x或者y參數中最小的值。

max(x,y)

x 或者y參數中最大的值。

Lerp嘗試在某些條件下傳回兩個不同的任意值時最友善。

這一步驟對于Boolean運算非常有用。如果這是真的，你可能還會想要傳回一些東西，例如假設你有以下功能：

// The enable parameter should be either 0 or 1
// If enable is 1, then the color blue is returned.  Otherwise the color black is returned.
// If you add the result of AddBlueTint to your final color and black is returned,
// your original color will remain unchanged.
float3 AddBlueTint( float enable )
{
            return step(0.5, enable) * float3(0,0,1);
}

float3 finalColor = someColor + AddBlueTint(1.0);
           

也可以修改立體聲shader并使用step代替lerp。

fixed4 frag (v2f i) : SV_Target
{
    UNITY_SETUP_STEREO_EYE_INDEX_POST_VERTEX(i);
    return step(unity_StereoEyeIndex, 0.5) *_LeftEyeColor + step(0.5, unity_StereoEyeIndex) * _RightEyeColor;
}
           

那樣的話，HLSL彙編如下所示：

Shader hash e56cb1bf-ce5441de-72c92b9d-336cc956

ps_4_0
      dcl_constantbuffer cb0[2], immediateIndexed
      dcl_input_ps_siv v1.x, rendertarget_array_index
      dcl_output o0.xyzw
      dcl_temps 2
   0: utof r0.x, v1.x
   1: ge r0.y, r0.x, l(0.500000)
   2: ge r0.x, l(0.500000), r0.x
   3: and r0.xy, r0.xyxx, l(1.000000, 1.000000, 0.000000, 0.000000)
   4: mul r1.xyzw, r0.yyyy, cb0[1].xyzw
   5: mad o0.xyzw, r0.xxxx, cb0[0].xyzw, r1.xyzw
   6: ret
           

再一次…

沒有分支！SIMD shader彙編指令"ge"不是分支。基本條件下，HLSL階段指令在shader彙編中進行轉換。ge是計算輸入值是否大于或等于其他輸入值的指令，并根據結果傳回值1或者0。使用多階段函數，可以根據輸入值是都在特定範圍内兒傳回結果。這一般稱為"脈沖功能"(“pulse function” : The pulse function may also be expressed as a limit of a rational function)

例如可以考慮以下條件：

if(value >= 0.5 && value <= 1.0)
{
	return float3(0,0,1);
}
           

使用2 階段功能構成的脈沖功能。

float pulse(float value, float minValue, float maxValue)
{
	return step(minValue,value) - step(maxValue, value);
}

float3 finalColor = float3(0,0,1) * pulse(0.4, 0.1, 0.5);
           

可以使用最小值/最大值去選擇兩個值中大或較小的值。這樣的内置功能都轉換為硬體加速功能或者與lerp相同的硬體加速功能。即，這些方法通過顯示卡的矽膠中烤制的邏輯回路進行轉換或被轉換。

删除條件的更多方法請關注Orange Duck的以下文章。GPU核心即工作方式的詳細内容請參考一下UC Davis的示範。

附錄. 1

GPU 運算方式

1. SIMD, MIMD, SIMT

(a) SIMD(Single Instruction, Multiple Data): 用一個指令加工多個資料的計算。

(b) MIMD(Multiple Instruction, Multiple Data): 幾個處理器核心運作各自獨立的程式（指令 列），并各自處理資料。

© SIMT(Single Instruction, Multiple Thread): 用一個指令(NVIDIA GPU時)并行運作32個線程

*P(Predicate): 有一個寄存器用于記住上述的比較指令條件代碼，使用各個指令的Fredicate比特和條件代碼的比對/不比對控制是否忽略在運算其中執行指令。

SIMD

• 因為涉及到4各元素的向量，是以用相同的指令将4個運算器結合在一起的方法很有效。
• 如果隻輸入一個元素，運算器4個裡面隻會運作1個，剩餘3個空閑，效率低。
• 另外，必須讓按照一般順序運作的程式執行不同的向量運算。
• 用CUDA或OpenCL編寫程式時，SIMD用矢量化或詳細的指令執行順序等是編譯器負責的，是以不需要程式員考慮。編譯器無法保證可以生成所有的SIMD代碼，可能會使性能降低。

• 指令供應是一個整體，各運算器都不需要指令供應，是以所需的半導體數量很少，對比MIMD，面積要更小。

  需要的硬體數量： MIMD > SIMT > SIMD

  運作自由度： MIMD > SIMT > SIMD

GPU有從運作SIMD轉變為運作SIMT的趨勢。

1. 一般情況下，CPU和GPU都有各自的記憶體。

• 通過DMA(Direct Memory Access)引擎，在兩個記憶體之間傳送資料。
• 這樣的情況下，每次資料傳送都需要在記憶體之間複制資料，但也會發生與‘深層複制’的情況一樣的問題。
• 是以制作共同記憶體的空間成為了GPU制造商的重要研發目标。

CPU => 需要大容量記憶體(CPU數十GB / GPU數GB)

GPU => 需要高帶寬記憶體以支援高計算性能

4. GPU的3D圖像處理需要大量計算，是以GPU需要搭載比通用CPU更多的浮點運算器

是以産生了将這樣的運算能力運用到科技計算的需求，GPU的适用範圍更廣了。

e.g.) 利用生物分子系統功能控制，建立個新的制約基礎時

• 生物體組織等按照分子原子的水準模型化
• 模拟與藥物的互動産生怎樣的結果
• 原子層面上，為計算移動的力量，因為原子數量大，需要大量的計算

諸如科技計算，圖像識别，語音識别等研究趨勢将資料分成更小的機關進行分析，精密度準确度也越來越高。

是以需要更多的計算，使用GPU的并行計算已經成為處理的重心。

附錄 2.

Warp Execution

Warp是執行GPU Instruction的最基本機關。Warp的所有Thread隻執行一個相同的Instruction。但是Warp的各Thread都可以讀取不同的資料值。是以可以認為Warp的每個Thread使用相同的Instruction不同的資料值來執行運算。這種運算方法被稱為SIMD (Single Instruction Multiple Data) 或者SIMT (Single Instruction Multiple Thread)。

SIMD和SIMT有一點差異。SIMD可以看做是執行Vector運算。要執行SIMD運算，為了執行Vector運算需要Instruction。而SIMT則是在編寫程式時，根據一個Thread來編寫。而且硬體是将幾個Thread組合在一起來執行運算。

從Volta Architecture開始，Warp的Thread可以執行互相獨立的運算。沒有詳細的閱讀Volta Architecture，是以操作方法不是很清楚下次打算整理一下Volta或Turing Architecture(出處1, 2)

圖1顯示了使用SIMT運算方法執行ADD運算。32個Thread有獨立的存儲空間(Register)。為了執行ADD運算，每個Thread從兩個Source Register加載Data值。執行ADD運算之後得到的結果存儲在其他32個Destination Register。結果上看一個Warp将通路64個Source Register，并将運算結果值存儲在32個Register裡面。Warp的所有Thread都為了執行同樣的運算，之家在一個Instruction。準确的說，在被排程時要執行相同的運算的Instruction執行了Broadcasting，同時所有的Thread執行被Broadcasting的一個Instruction。

通常一個SM (Streaming Multiprocessor)有48個64個的Warp。每個SM一般有數十數百個的運算器(ALU)，可同時進行2~4個Warp Instruction。每個SM的Warp總數，ALU個數等因GPU Architecture的不同而略有不同。GPU與CPU一樣使用Pipeline執行Instruction。通常GPU由5個Pipeline Stage組成，分别是Instruction Fetch, Decode, Issue, ALU&Memory, Writeback格式。(出處 3)

從Turing GPU Architecture開始 Integer, Floating運算器已經分離了。以前Architecture的情況下，Integer, Floating Point運算都在一個ALU裡面執行了。

GPGPU-Sim是對NVIDIA GPU進行模組化的電腦Architecture模拟器。用C, C++寫的。雖然模拟器陳舊，但是因為沒有比這個更好地選擇，是以很多研究室都是用相應模拟器編寫論文。

圖2展示了5個Pipeline Stage。以下是對各Stage的說明。

Fetch: Fetch階段使用PC (Program Counter)的值加載Instruction。最近的GPU，與CPU一樣使用64-Bit Instruction。Fetch Instruction的順序是Round-Robin方式。用Warp 0, 1, 2, ….的順序Fetch Instruction。

Decode: Decode階段是為執行在記憶體裡面加載的Instruction執行ALU/Memory運算而查找運算Type及需要的Register的階段。 Decode Instruction存儲在Instruction Buffer中。GPU由不同的邏輯計算，如Integer, Floating Point, Special Function, Memory。

Issue: 是用運算器發送Decode Instruction的階段。在此階段中，運算所需的資料值在Source Register讀取。在Source Register讀取的值發送到ALU，ALU将執行運算。一般情況下，Instruction Fetch/Decode的速度(?)因為比Instruction Issue的速度快，是以到了這個階段Decode的(可能的Issue)Warp個數多的機率會很高。是以Warp Scheduler選擇Issue的Warp。衆所周知Warp Scheduling技術有Round-Robin (RR)和Greedy-Then-Oldest (GTO)方式。SM平均有2~4個Warp Scheduler，是以每個周期會進行2~8個Warp Instruction Issue。Issue Instruction可在ALU/LSU執行運算。根據運算類型，Instruction會在數十個周期内需要數百周期的運算時間。

RR Warp Scheduling是從幾個Warp中用ALU組成的Warp的順序，依次選擇的方法。例如，假設SM中有24個Warp，并且所有Warp都已經準備好了運作Instruction。Warp 0, Warp 1, …, Warp 23在ALU裡面依次出現一個Instruction，是Issue的形态。

GTO Warp Scheduling方法已經執行了在Warp 0号可執行Instruction。如果沒有準備好的Warp 0的可執行Instruction(Stall發生的情況)，則選擇下一個準備好的Warp來進行Instruction Issue。下一個選擇的Warp選擇現在Stall Warp的下一個ID的 Warp。例如，Warp 3号沒有可執行的 Instruction，則4号Warp Instruction将成為Issue。

ALU&Memory: 使用從Source Register讀取的值執行運算。根據Instruction類型，使用Load/Store Unit (LSU), ALU, Special Function Unit (SFU)運算器。ALU執行Integer, Floating Point運算。SFU執行一些有些複雜的指令，如Sin, Cos等。LSU執行記憶體運算。…

Writeback: 是将在ALU, LSU, SFU運算的值存儲在Destination Register的階段。

GPU執行上述5個階段的Warp Instruction。一般Computer Architecture書裡面介紹的5階段Pipeline Stage和它幾乎相同。

出處

https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf

https://www.nvidia.com/content/dam/en-zz/Solutions/design-visualization/technologies/turing-architecture/NVIDIA-Turing-Architecture-Whitepaper.pdf

http://www.gpgpu-sim.org/

附錄3.

GPU Branch Divergence

GPU将Thread綁定為32或64個，執行一個相同的Instruction。Branch Divergence發生在屬于Warp的Thread必須執行不同的操作時。例如，屬于Warp的Thread偶數ID必須執行與IF語句相對應的代碼，而為奇數ID的Thread必須執行ELSE相對應的代碼，就會發生Branch Divergence。以下代碼是檢查Thread ID是否為奇數，并運作不同代碼的示例（未驗證）。

__global__ void mkTest(){
	int threadID = threadIdx.x
	if((threadIdx.x % 2) == 0){
		Even number Thread code...
	}
	else{
		odd number Thread code...
	}
}
           

圖1表示了發生Branch Divergence的情況下，Warp執行的順序。簡而言之，Warp依次執行IF和ELSE等代碼。執行與IF相應的代碼運算時，不需要IF運算的Thread運算除外。相反，執行ELSE相應的運算時，不需要ELSE運算的Thread将從運算中除外。不知道運算中去除的正确方法。也有不完全執行運算的方法，相反也有執行運算但是不在Register使用結果值的方法。有效的方法是不執行計算本身。

圖 1: Branch Divergence發生的情況下，執行Warp的順序

是以當發生Branch Divergence時，GPU将按序執行IF/ELSE相應的代碼。因為Warp的所有Thread不會同時執行運算，是以GPU的Utilization下降。是以在使用GPU程式設計時，最好減少Branch Divergence。有事是在将IF/ELSE除以Thread ID/Block ID/Grid ID時，最好添加一到兩個不同的運算來删除Branch Divergence。GPGPU早期(2010年左右)，進行了很多從HW中删除Branch Divergence的研究。例如，當發生Branch Divergence，将Warp的Thread變成其他的Warp Thread。在研究所學生時期讀論文的時候，我覺得這是個很酷的想法。但是，現在回想起來在性能上應該不會有很大的益處。最大的原因是Register值讀取的速度急速變慢，是以可能實際性能并不會像論文中說的那麼好。此外也有很多在編譯階段解決Branch Divergence問題的研究釋出。很多論文發表意味着Branch Divergence對GPU的性能有很大的影響。

根據GPGPU-Sim的說法，使用SIMT Stack硬體邏輯Handling Branch Divergence(出處 1)。每個Warp都有一個名為SIMT Stack的存儲空間。SIMT Stack由幾個Entry構成。根據論文内容，SIMT Stack由4個Entry構成(出處2)。每個Entry由PC (Program Counter) + 32Bit存儲空間構成。如果PC是64 Bit，每個Entry由92 Bit構成。PC後面的32 Bit用于存儲True/False值，即屬于Warp的Thread是執行還是不執行與PC相對應的Instruction。對于NVIDIA GPU，由于Warp由32個Thread構成，是以使用32 Bit檢查Thread的Instruction是否執行。但是如果Warp由64個Thread構成，則使用64 Bit檢查是否執行Thread Instruction。

想了一下如何解釋運作原理，似乎使用示例進行解釋是最有效的。圖2展示了Branch Divergence發生的情況下，使用SIMT Stack執行Instruction的方法。為友善起見，我們假設Warp是由8個Thread構成的。如前面示例代碼所說，假設偶數ID Thread執行IF語句，奇數ID Thread執行ELSE語句。

Step 1: 在IF/ELSE之前，SIMT Stack有一個Entry。該Entry存儲了PC值 + 11111111 值。PC值存儲了Warp必須執行的下一個Instruction位置值。随後存儲的 “11111111”意味着所有現存的Thread都必須執行與PC相對應的Instruction運算。

Step 2: 如果遇到IF/ELSE語句，總共會生成3個Entry。最上面存儲IF語句相應的IF Instruction PC值 + “10101010”值，表示需要進行偶數Thread運算（因為是從0開始，第一個Thread标記為1）。以下Entry存儲與ELSE語句Instruction相對應的PC值 + “01010101”值。最後存儲完成了與IF/ELSE相對應的Instruction運算的PC 值 (RPC) + “111111111”。IF/ELSE語句完成後，Warp的所有Thread為了再次執行相同的Instruction而合并，稱為 “Reconverge”。

Step 3: Stack這個名稱可以得知，使用POP運算在SIMT Stack存儲的Entry值。首先在SIMT Stack執行POP運算會得出與IF語句相對應的Entry。判斷出IF語句相對應的PC值和Warp的Instruction運算是否需要後執行運算。标記為“1”的Thread執行其Instruction運算。反之，标記為“0”的Thread不執行。

Step 4: 與Step 3一樣執行以下Entry運算。以下Entry是對應ELSE的Instruction的運算。

如果IF/ELSE的Instruction個數多于1，則更新PC值并将可執行的Thread資訊再次儲存至SIMT Stack。Stack遵循Last-In-First-Out (LIFO)順序，是以執行所有相應于IF運算的Instruction運算，并按序執行ELSE相應的Instruction運算。

Step 5: IF/ELSE相應的Instruction全部執行後，SIMT Stack将隻剩下具有Reconverge PC (RPC)值的Entry。該Entry被POP，Warp的所有Thread都執行相同的運算。

如上所述，Branch Divergence發生的話，使用SIMT Stack按序執行Instruction運算。發生Branch Divergence時，由于屬于Warp的Thread無法執行Instruction，導緻GPU的Utilization下降。此外，由于執行了IF/ELSE等相應的所有Instruction，GPU必須要執行的Instruction的個數增加了。總之，最重要的是要確定在程式設計過成功不發生Branch Divergence。

出處

http://www.gpgpu-sim.org/

Stack-less SIMT Reconvergence at Low Cost

About JP

連結: Website.

在這裡插入圖檔描述

出生在南韓的TA。

1997年開始從事電腦圖形視覺化工作後，在這個行業已經有21年經驗了。

在多個網絡遊戲公司引領過美術團隊，之前在allegorithmic擔任TA負責人，在中國網易盤古工作室擔任TA總監，現在是巨人網絡TA部門的總負責人。

懶惰的人才有創意”是他堅信并執行的哲學道理。