CUDA學習（四十六）

Warp 比對函數：

__match_any_sync和__match_all_sync對warp内的線程之間的變量執行廣播和比較操作。

由計算能力7.x或更高的裝置支援。

Synopsys：

T可以是int，unsigned int，long，unsigned long，long long，unsigned long long，float或double。

描述：

<code>__match_sync（）</code>内部函數允許在同步線程中命名的線程之後，對warp中的線程之間的值進行廣播和比較。

<code>__match_any_sync</code>：

傳回掩碼中具有相同值的線程的掩碼。

<code>__match_all_sync</code>:

如果掩碼中的所有線程都具有相同的值，則傳回掩碼; 否則傳回0。 如果掩碼中的所有線程都具有相同的值值，則Predicate pred将設定為true; 否則謂詞設定為false。

新的* _sync比對内部函數接收訓示參與調用的線程的掩碼。 表示線程的通道ID的位必須為每個參與線程設定，以確定它們在内部由硬體執行之前正确收斂。 在掩碼中命名的所有非退出線程都必須使用相同的掩碼執行相同的内部屬性，否則結果未定義。

Warp Shuffle函數：

<code>__shfl_sync</code>，<code>__shfl_up_sync</code>，<code>__shfl_down_sync</code>和<code>__shfl_xor_sync</code>在變形内的線程之間交換變量。

由計算能力3.x或更高的裝置支援。

棄用聲明：自CUDA 9.0起，<code>__shfl</code>，<code>__shfl_up</code>，<code>__shfl_down</code>和<code>__shfl_xor</code>已被棄用。

概要：

T可以是int，unsigned int，long，unsigned long，long long，unsigned long long，float或double。 通過包含cuda_fp16.h頭檔案，T也可以是<code>__half</code>或<code>__half2</code>。

<code>__shfl_sync（）</code>内部函數允許在不使用共享記憶體的情況下交換變形内的線程之間的變量。 對于warp中的所有活動線程（并在mask中命名），交換同時發生，根據類型，每個線程移動4或8個位元組的資料。

變形中的線程稱為通道，可能具有介于0和warpSize-1（含）之間的索引。 支援四種源通道尋址模式：

<code>__shfl_sync()</code>：

從索引lane直接複制。

<code>__shfl_up_sync()</code>:

從ID較低的lane複制相對于呼叫者。

<code>__shfl_down_sync()</code>:

從ID較高的lane複制相對于呼叫者。

<code>__shfl_xor_sync()：</code>

基于自己lane ID的按位XOR從車道複制

線程隻能從另一個主動參與<code>__shfl_sync（）</code>指令的線程讀取資料。 如果目标線程處于非活動狀态，則檢索到的值不确定。

所有的<code>__shfl_sync（）</code>内部函數都有一個可選的寬度參數，它可以改變内部的行為。 寬度必須是2的幂的值; 如果寬度不是2的幂，或者是大于warpSize的數字，結果是不确定的。

<code>__shfl_sync（）</code>傳回由srcLane給出ID的線程持有的var的值。 如果寬度小于warpSize，那麼warp的每個子部分都會作為一個開始邏輯通道ID為0的獨立實體運作。如果srcLane超出範圍[0：width-1]，則傳回的值對應于var的值 由srcLane模寬度（即在同一小節内）。

<code>__shfl_up_sync（）</code>通過從呼叫者的laneID減去delta來計算源lane ID。 傳回結果lane辨別所儲存的var的值：實際上，var通過三角形lane向上移動變形。 如果寬度小于warpSize，那麼warp的每個子部分都表現為一個開始邏輯通道ID為0的獨立實體。源通道索引不會環繞寬度值，是以實際上較低delta通道将保持不變。

<code>__shfl_down_sync（）</code>通過将delta添加到呼叫者的lane ID來計算源lane ID。 傳回結果lane 辨別所儲存的var的值将被傳回：這具有通過三角形lane 将變形向下移動var的效果。 如果寬度小于warpSize，那麼warp的每個子部分都表現為一個獨立的實體，其起始邏輯通道ID為0.對于<code>__shfl_up_sync（）</code>，源通道的ID号不會環繞width的值，是以 上三角洲lane 将保持不變。

<code>__shfl_xor_sync（）</code>通過調用laneMask執行呼叫者lane ID的按位XOR來計算源線路ID：傳回由生成的lane ID儲存的var值。 如果width小于warpSize，那麼每組寬度連續的線程都可以通路先前線程組中的元素，但是如果它們嘗試從後面的線程組通路元素，則會傳回它們自己的var值。 該模式實作了一種蝴蝶尋址模式，例如用于樹狀縮減和廣播。

新的* _sync shfl内部函數接收一個訓示參與調用的線程的掩碼。 表示線程的通道ID的位必須為每個參與線程設定，以確定它們在内部由硬體執行之前正确收斂。 在掩碼中命名的所有非退出線程都必須使用相同的掩碼執行相同的内部屬性，否則結果未定義。

傳回值：

所有的<code>__shfl_sync（）</code>内在函數都會将源通道ID中由var引用的4位元組字作為無符号整數傳回。 如果源通道辨別超出範圍或源線程退出，則傳回調用線程自己的變量。

注意：

寬度必須是2的幂（即2，4，8，16或32）。 其他值的結果未指定。